From 95aee578d4ad2e086f74b64cc4c5ad7658ff7896 Mon Sep 17 00:00:00 2001 From: gregory Date: Wed, 10 Aug 2022 11:14:19 -0700 Subject: [PATCH 1/4] Introduce a test for mix of unnamed ESIMD and nonESIMD kernels compilation --- .../regression/sycl_esimd_mixed_unnamed.cpp | 125 ++++++++++++++++++ 1 file changed, 125 insertions(+) create mode 100644 SYCL/ESIMD/regression/sycl_esimd_mixed_unnamed.cpp diff --git a/SYCL/ESIMD/regression/sycl_esimd_mixed_unnamed.cpp b/SYCL/ESIMD/regression/sycl_esimd_mixed_unnamed.cpp new file mode 100644 index 0000000000..cde1d3f5d7 --- /dev/null +++ b/SYCL/ESIMD/regression/sycl_esimd_mixed_unnamed.cpp @@ -0,0 +1,125 @@ +//==--- sycl_esimd_mixed_unnamed.cpp - DPC++ ESIMD on-device test ---------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// This is basic test for mixing unnamed SYCL and ESIMD kernels in the same +// source and in the same program . + +// REQUIRES: gpu +// UNSUPPORTED: cuda || hip +// UNSUPPORTED: esimd_emulator +// RUN: %clangxx -fsycl %s -o %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out +// XFAIL: esimd_emulator + +#include "esimd_test_utils.hpp" + +#include +#include +#include + +using namespace cl::sycl; + +bool checkResult(const std::vector &A, int Inc) { + int err_cnt = 0; + unsigned Size = A.size(); + + for (unsigned i = 0; i < Size; ++i) { + if (A[i] != i + Inc) + if (++err_cnt < 10) + std::cerr << "failed at A[" << i << "]: " << A[i] << " != " << i + Inc + << "\n"; + } + + if (err_cnt > 0) { + std::cout << " pass rate: " + << ((float)(Size - err_cnt) / (float)Size) * 100.0f << "% (" + << (Size - err_cnt) << "/" << Size << ")\n"; + return false; + } + return true; +} + +int main(void) { + constexpr unsigned Size = 32; + constexpr unsigned VL = 16; + + std::vector A(Size); + + for (unsigned i = 0; i < Size; ++i) { + A[i] = i; + } + + try { + buffer bufa(A.data(), range<1>(Size)); + + // We need that many workgroups + cl::sycl::range<1> GlobalRange{Size}; + // We need that many threads in each group + cl::sycl::range<1> LocalRange{1}; + + queue q(gpu_selector{}, esimd_test::createExceptionHandler()); + + auto dev = q.get_device(); + std::cout << "Running on " << dev.get_info() << "\n"; + + auto e = q.submit([&](handler &cgh) { + auto PA = bufa.get_access(cgh); + cgh.parallel_for(GlobalRange * LocalRange, + [=](id<1> i) { PA[i] = PA[i] + 1; }); + }); + e.wait(); + } catch (cl::sycl::exception const &e) { + std::cout << "SYCL exception caught: " << e.what() << '\n'; + return 2; + } + + if (checkResult(A, 1)) { + std::cout << "SYCL kernel passed\n"; + } else { + std::cout << "SYCL kernel failed\n"; + return 1; + } + + try { + buffer bufa(A.data(), range<1>(Size)); + + // We need that many workgroups + cl::sycl::range<1> GlobalRange{Size / VL}; + // We need that many threads in each group + cl::sycl::range<1> LocalRange{1}; + + queue q(esimd_test::ESIMDSelector{}, esimd_test::createExceptionHandler()); + + auto dev = q.get_device(); + std::cout << "Running on " << dev.get_info() << "\n"; + + auto e = q.submit([&](handler &cgh) { + auto PA = bufa.get_access(cgh); + cgh.parallel_for(GlobalRange * LocalRange, + [=](id<1> i) SYCL_ESIMD_KERNEL { + using namespace sycl::ext::intel::esimd; + unsigned int offset = i * VL * sizeof(float); + simd va; + va.copy_from(PA, offset); + simd vc = va + 1; + vc.copy_to(PA, offset); + }); + }); + e.wait(); + } catch (cl::sycl::exception const &e) { + std::cout << "SYCL exception caught: " << e.what() << '\n'; + return 2; + } + + if (checkResult(A, 2)) { + std::cout << "ESIMD kernel passed\n"; + } else { + std::cout << "ESIMD kernel failed\n"; + return 1; + } + return 0; +} From a577c3562875592a4cbf83582c399c17e4048ade Mon Sep 17 00:00:00 2001 From: gregory Date: Wed, 10 Aug 2022 12:55:12 -0700 Subject: [PATCH 2/4] Fix a build issue --- SYCL/ESIMD/regression/sycl_esimd_mixed_unnamed.cpp | 6 ++---- 1 file changed, 2 insertions(+), 4 deletions(-) diff --git a/SYCL/ESIMD/regression/sycl_esimd_mixed_unnamed.cpp b/SYCL/ESIMD/regression/sycl_esimd_mixed_unnamed.cpp index cde1d3f5d7..e989856f62 100644 --- a/SYCL/ESIMD/regression/sycl_esimd_mixed_unnamed.cpp +++ b/SYCL/ESIMD/regression/sycl_esimd_mixed_unnamed.cpp @@ -15,8 +15,6 @@ // RUN: %GPU_RUN_PLACEHOLDER %t.out // XFAIL: esimd_emulator -#include "esimd_test_utils.hpp" - #include #include #include @@ -61,7 +59,7 @@ int main(void) { // We need that many threads in each group cl::sycl::range<1> LocalRange{1}; - queue q(gpu_selector{}, esimd_test::createExceptionHandler()); + queue q; auto dev = q.get_device(); std::cout << "Running on " << dev.get_info() << "\n"; @@ -92,7 +90,7 @@ int main(void) { // We need that many threads in each group cl::sycl::range<1> LocalRange{1}; - queue q(esimd_test::ESIMDSelector{}, esimd_test::createExceptionHandler()); + queue q; auto dev = q.get_device(); std::cout << "Running on " << dev.get_info() << "\n"; From 98b3916b708ec096be6639df7ff5a6a9bbca3ad2 Mon Sep 17 00:00:00 2001 From: gregory Date: Wed, 17 Aug 2022 23:34:57 -0700 Subject: [PATCH 3/4] Eliminate build warnings --- SYCL/ESIMD/regression/sycl_esimd_mixed_unnamed.cpp | 14 +++++++------- 1 file changed, 7 insertions(+), 7 deletions(-) diff --git a/SYCL/ESIMD/regression/sycl_esimd_mixed_unnamed.cpp b/SYCL/ESIMD/regression/sycl_esimd_mixed_unnamed.cpp index e989856f62..873a27c5f1 100644 --- a/SYCL/ESIMD/regression/sycl_esimd_mixed_unnamed.cpp +++ b/SYCL/ESIMD/regression/sycl_esimd_mixed_unnamed.cpp @@ -19,7 +19,7 @@ #include #include -using namespace cl::sycl; +using namespace ::sycl; bool checkResult(const std::vector &A, int Inc) { int err_cnt = 0; @@ -55,9 +55,9 @@ int main(void) { buffer bufa(A.data(), range<1>(Size)); // We need that many workgroups - cl::sycl::range<1> GlobalRange{Size}; + ::sycl::range<1> GlobalRange{Size}; // We need that many threads in each group - cl::sycl::range<1> LocalRange{1}; + ::sycl::range<1> LocalRange{1}; queue q; @@ -70,7 +70,7 @@ int main(void) { [=](id<1> i) { PA[i] = PA[i] + 1; }); }); e.wait(); - } catch (cl::sycl::exception const &e) { + } catch (::sycl::exception const &e) { std::cout << "SYCL exception caught: " << e.what() << '\n'; return 2; } @@ -86,9 +86,9 @@ int main(void) { buffer bufa(A.data(), range<1>(Size)); // We need that many workgroups - cl::sycl::range<1> GlobalRange{Size / VL}; + ::sycl::range<1> GlobalRange{Size / VL}; // We need that many threads in each group - cl::sycl::range<1> LocalRange{1}; + ::sycl::range<1> LocalRange{1}; queue q; @@ -108,7 +108,7 @@ int main(void) { }); }); e.wait(); - } catch (cl::sycl::exception const &e) { + } catch (::sycl::exception const &e) { std::cout << "SYCL exception caught: " << e.what() << '\n'; return 2; } From 31145e4d9586cb882cba936d879502e9b5f2f743 Mon Sep 17 00:00:00 2001 From: gregory Date: Wed, 17 Aug 2022 23:40:15 -0700 Subject: [PATCH 4/4] Use sycl/sycl.hpp instead of CL/sycl.hpp --- SYCL/ESIMD/regression/sycl_esimd_mixed_unnamed.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/SYCL/ESIMD/regression/sycl_esimd_mixed_unnamed.cpp b/SYCL/ESIMD/regression/sycl_esimd_mixed_unnamed.cpp index 873a27c5f1..046a767dd1 100644 --- a/SYCL/ESIMD/regression/sycl_esimd_mixed_unnamed.cpp +++ b/SYCL/ESIMD/regression/sycl_esimd_mixed_unnamed.cpp @@ -15,9 +15,9 @@ // RUN: %GPU_RUN_PLACEHOLDER %t.out // XFAIL: esimd_emulator -#include #include #include +#include using namespace ::sycl;