diff --git a/sycl/test-e2e/AOT/accelerator.cpp b/sycl/test-e2e/AOT/accelerator.cpp new file mode 100644 index 0000000000000..602ba415f4f50 --- /dev/null +++ b/sycl/test-e2e/AOT/accelerator.cpp @@ -0,0 +1,12 @@ +//=-- accelerator.cpp - compilation for fpga emulator dev using opencl-aot --=// +// +// 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 +// +//===---------------------------------------------------------------------===// + +// REQUIRES: opencl-aot, accelerator + +// RUN: %clangxx -fsycl -fsycl-targets=spir64_fpga %S/Inputs/aot.cpp -o %t.out +// RUN: %ACC_RUN_PLACEHOLDER %t.out diff --git a/sycl/test-e2e/Assert/Inputs/kernels_in_file2.cpp b/sycl/test-e2e/Assert/Inputs/kernels_in_file2.cpp new file mode 100644 index 0000000000000..e07fdbb3c5560 --- /dev/null +++ b/sycl/test-e2e/Assert/Inputs/kernels_in_file2.cpp @@ -0,0 +1,45 @@ +#include "kernels_in_file2.hpp" + +#ifdef DEFINE_NDEBUG_INFILE2 +#define NDEBUG +#else +#undef NDEBUG +#endif + +#include + +using namespace sycl; +using namespace sycl::access; + +int calculus(int X) { + assert(X && "this message from calculus"); + return X * 2; +} + +void check_nil(int value) { assert(value && "this message from file2"); } + +static constexpr size_t BUFFER_SIZE = 4; + +void enqueueKernel_1_fromFile2(queue *Q) { + sycl::range<1> numOfItems{BUFFER_SIZE}; + sycl::buffer Buf(numOfItems); + + Q->submit([&](handler &CGH) { + auto Acc = Buf.template get_access(CGH); + + CGH.parallel_for( + numOfItems, [=](sycl::id<1> wiID) { check_nil(Acc[wiID]); }); + }); +} + +void enqueueKernel_2_fromFile2(queue *Q) { + sycl::range<1> numOfItems{BUFFER_SIZE}; + sycl::buffer Buf(numOfItems); + + Q->submit([&](handler &CGH) { + auto Acc = Buf.template get_access(CGH); + + CGH.parallel_for( + numOfItems, [=](sycl::id<1> wiID) { check_nil(Acc[wiID]); }); + }); +} diff --git a/sycl/test-e2e/Assert/Inputs/kernels_in_file2.hpp b/sycl/test-e2e/Assert/Inputs/kernels_in_file2.hpp new file mode 100644 index 0000000000000..851207ac66026 --- /dev/null +++ b/sycl/test-e2e/Assert/Inputs/kernels_in_file2.hpp @@ -0,0 +1,7 @@ +#include + +SYCL_EXTERNAL int calculus(int X); + +void enqueueKernel_1_fromFile2(sycl::queue *Q); + +void enqueueKernel_2_fromFile2(sycl::queue *Q); diff --git a/sycl/test-e2e/Assert/assert_in_kernels.cpp b/sycl/test-e2e/Assert/assert_in_kernels.cpp new file mode 100644 index 0000000000000..eaf67c0e42cb9 --- /dev/null +++ b/sycl/test-e2e/Assert/assert_in_kernels.cpp @@ -0,0 +1,24 @@ +// REQUIRES: linux + +// https://github.com/intel/llvm/issues/7634 +// UNSUPPORTED: hip + +// RUN: %clangxx -DSYCL_FALLBACK_ASSERT=1 -fsycl -fsycl-targets=%sycl_triple %s -o %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out &> %t.cpu.txt || true +// RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.cpu.txt +// RUN: %GPU_RUN_PLACEHOLDER %t.out &> %t.gpu.txt || true +// RUN: %GPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.gpu.txt +// Shouldn't fail on ACC as fallback assert isn't enqueued there +// RUN: %ACC_RUN_PLACEHOLDER %t.out &> %t.acc.txt +// RUN: %ACC_RUN_PLACEHOLDER FileCheck %s --check-prefix=CHECK-ACC --input-file %t.acc.txt +// +// CHECK-NOT: One shouldn't see this message +// CHECK: {{.*}}assert_in_kernels.hpp:25: void kernelFunc2(int *, int): {{.*}} [{{[0,2]}},0,0], {{.*}} [0,0,0] +// CHECK-SAME: Assertion `Buf[wiID] == 0 && "from assert statement"` failed. +// CHECK-NOT: test aborts earlier, one shouldn't see this message +// CHECK-NOT: The test ended. +// +// CHECK-ACC-NOT: {{.*}}assert_in_kernels.hpp:25: void kernelFunc2(int *, int): {{.*}} [{{[0,2]}},0,0], {{.*}} [0,0,0] +// CHECK-ACC: The test ended. + +#include "assert_in_kernels.hpp" diff --git a/sycl/test-e2e/Assert/assert_in_kernels.hpp b/sycl/test-e2e/Assert/assert_in_kernels.hpp new file mode 100644 index 0000000000000..834ddb22da00b --- /dev/null +++ b/sycl/test-e2e/Assert/assert_in_kernels.hpp @@ -0,0 +1,69 @@ +#include +#include +#include + +using namespace sycl; +using namespace sycl::access; + +void kernelFunc1(int *Buf, int wiID) { + Buf[wiID] = 9; + assert(Buf[wiID] != 0 && "One shouldn't see this message"); +} + +void assertTest1(queue &Q, buffer &Buf) { + Q.submit([&](handler &CGH) { + auto Acc = Buf.template get_access(CGH); + + CGH.parallel_for( + Buf.get_range(), [=](sycl::id<1> wiID) { kernelFunc1(&Acc[0], wiID); }); + }); +} + +void kernelFunc2(int *Buf, int wiID) { + if (wiID % 2 != 0) + Buf[wiID] = 0; + assert(Buf[wiID] == 0 && "from assert statement"); +} + +void assertTest2(queue &Q, buffer &Buf) { + Q.submit([&](handler &CGH) { + auto Acc = Buf.template get_access(CGH); + + CGH.parallel_for( + Buf.get_range(), [=](sycl::id<1> wiID) { kernelFunc2(&Acc[0], wiID); }); + }); +} + +void kernelFunc3(int *Buf, int wiID) { + if (wiID == 0) + assert(false && "test aborts earlier, one shouldn't see this message"); + Buf[wiID] = 9; +} + +void assertTest3(queue &Q, buffer &Buf) { + Q.submit([&](handler &CGH) { + auto Acc = Buf.template get_access(CGH); + + CGH.parallel_for( + Buf.get_range(), [=](sycl::id<1> wiID) { kernelFunc3(&Acc[0], wiID); }); + }); +} + +int main(int Argc, const char *Argv[]) { + std::array Vec = {1, 2, 3, 4}; + sycl::range<1> numOfItems{Vec.size()}; + sycl::buffer Buf(Vec.data(), numOfItems); + + queue Q; + assertTest1(Q, Buf); + Q.wait(); + + assertTest2(Q, Buf); + Q.wait(); + + assertTest3(Q, Buf); + Q.wait(); + + std::cout << "The test ended." << std::endl; + return 0; +} diff --git a/sycl/test-e2e/Assert/assert_in_kernels_ndebug.cpp b/sycl/test-e2e/Assert/assert_in_kernels_ndebug.cpp new file mode 100644 index 0000000000000..4475083ffeb0c --- /dev/null +++ b/sycl/test-e2e/Assert/assert_in_kernels_ndebug.cpp @@ -0,0 +1,9 @@ +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DNDEBUG %S/assert_in_kernels.cpp -o %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER +// RUN: %GPU_RUN_PLACEHOLDER %t.out %GPU_CHECK_PLACEHOLDER +// RUN: %ACC_RUN_PLACEHOLDER %t.out %ACC_CHECK_PLACEHOLDER +// +// CHECK-NOT: One shouldn't see this message +// CHECK-NOT: from assert statement +// CHECK-NOT: test aborts earlier, one shouldn't see this message +// CHECK: The test ended. diff --git a/sycl/test-e2e/Assert/assert_in_kernels_win.cpp b/sycl/test-e2e/Assert/assert_in_kernels_win.cpp new file mode 100644 index 0000000000000..8b05ce7a06afe --- /dev/null +++ b/sycl/test-e2e/Assert/assert_in_kernels_win.cpp @@ -0,0 +1,22 @@ +// REQUIRES: windows +// RUN: %clangxx -DSYCL_FALLBACK_ASSERT=1 -fsycl -fsycl-targets=%sycl_triple %s -o %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out &> %t.txt || true +// RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt +// RUN: %GPU_RUN_PLACEHOLDER %t.out &> %t.txt || true +// RUN: %GPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt +// Shouldn't fail on ACC as fallback assert isn't enqueued there +// RUN: %ACC_RUN_PLACEHOLDER %t.out &> %t.txt +// RUN: %ACC_RUN_PLACEHOLDER FileCheck %s --check-prefix=CHECK-ACC --input-file %t.txt +// +// CHECK-NOT: One shouldn't see this message +// FIXME Windows version prints '(null)' instead of '' once in a +// while for some insane reason. +// CHECK: {{.*}}assert_in_kernels.hpp:25: {{|(null)}}: {{.*}} [{{[0,2]}},0,0], {{.*}} [0,0,0] +// CHECK-SAME: Assertion `Buf[wiID] == 0 && "from assert statement"` failed. +// CHECK-NOT: test aborts earlier, one shouldn't see this message +// CHECK-NOT: The test ended. +// +// CHECK-ACC-NOT: {{.*}}assert_in_kernels.hpp:25: {{|(null)}}: {{.*}} [{{[0,2]}},0,0], {{.*}} [0,0,0] +// CHECK-ACC: The test ended. + +#include "assert_in_kernels.hpp" diff --git a/sycl/test-e2e/Assert/assert_in_multiple_tus.cpp b/sycl/test-e2e/Assert/assert_in_multiple_tus.cpp new file mode 100644 index 0000000000000..7de9915a505cf --- /dev/null +++ b/sycl/test-e2e/Assert/assert_in_multiple_tus.cpp @@ -0,0 +1,25 @@ +// REQUIRES: linux + +// https://github.com/intel/llvm/issues/7634 +// UNSUPPORTED: hip + +// RUN: %clangxx -DSYCL_FALLBACK_ASSERT=1 -fsycl -fsycl-targets=%sycl_triple -I %S/Inputs %s %S/Inputs/kernels_in_file2.cpp -o %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out &> %t.cpu.txt || true +// RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.cpu.txt +// RUN: %GPU_RUN_PLACEHOLDER %t.out &> %t.gpu.txt || true +// RUN: %GPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.gpu.txt +// Shouldn't fail on ACC as fallback assert isn't enqueued there +// RUN: %ACC_RUN_PLACEHOLDER %t.out &> %t.acc.txt +// RUN: %ACC_RUN_PLACEHOLDER FileCheck %s --check-prefix=CHECK-ACC --input-file %t.acc.txt +// +// CUDA uses block/thread vs global/local id for SYCL, also it shows the +// position of a thread within the block, not the absolute ID. +// CHECK: {{.*}}kernels_in_file2.cpp:15: int calculus(int): {{global id: \[5|block: \[1}},0,0], {{local id|thread}}: [1,0,0] +// CHECK-SAME: Assertion `X && "this message from calculus"` failed. +// CHECK-NOT: this message from file2 +// CHECK-NOT: The test ended. +// +// CHECK-ACC-NOT: {{.*}}kernels_in_file2.cpp:15: int calculus(int): global id: [5,0,0], local id: [1,0,0] +// CHECK-ACC: The test ended. + +#include "assert_in_multiple_tus.hpp" diff --git a/sycl/test-e2e/Assert/assert_in_multiple_tus.hpp b/sycl/test-e2e/Assert/assert_in_multiple_tus.hpp new file mode 100644 index 0000000000000..cf65d6a4616af --- /dev/null +++ b/sycl/test-e2e/Assert/assert_in_multiple_tus.hpp @@ -0,0 +1,51 @@ +#include "Inputs/kernels_in_file2.hpp" +#include +#include + +#ifdef DEFINE_NDEBUG_INFILE1 +#define NDEBUG +#else +#undef NDEBUG +#endif + +#include + +using namespace sycl; +using namespace sycl::access; + +static constexpr size_t BUFFER_SIZE = 16; + +int checkFunction() { + int X = calculus(0); + assert(X && "Nil in result"); + return X; +} + +void enqueueKernel_1_fromFile1(queue *Q) { + sycl::range<1> numOfItems{BUFFER_SIZE}; + sycl::buffer Buf(numOfItems); + + Q->submit([&](handler &CGH) { + auto Acc = Buf.template get_access(CGH); + + CGH.parallel_for( + sycl::nd_range(Buf.get_range(), sycl::range<1>(4)), + [=](sycl::id<1> wiID) { + int X = 0; + if (wiID == 5) + X = checkFunction(); + Acc[wiID] = X; + }); + }); +} + +int main(int Argc, const char *Argv[]) { + + queue Q; + enqueueKernel_1_fromFile1(&Q); + enqueueKernel_2_fromFile2(&Q); + Q.wait(); + + std::cout << "The test ended." << std::endl; + return 0; +} diff --git a/sycl/test-e2e/Assert/assert_in_multiple_tus_one_ndebug.cpp b/sycl/test-e2e/Assert/assert_in_multiple_tus_one_ndebug.cpp new file mode 100644 index 0000000000000..8acbce9277bbd --- /dev/null +++ b/sycl/test-e2e/Assert/assert_in_multiple_tus_one_ndebug.cpp @@ -0,0 +1,24 @@ +// REQUIRES: linux + +// https://github.com/intel/llvm/issues/7634 +// UNSUPPORTED: hip + +// RUN: %clangxx -DSYCL_FALLBACK_ASSERT=1 -fsycl -fsycl-targets=%sycl_triple -DDEFINE_NDEBUG_INFILE2 -I %S/Inputs %S/assert_in_multiple_tus.cpp %S/Inputs/kernels_in_file2.cpp -o %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out &> %t.cpu.txt || true +// RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.cpu.txt +// RUN: %GPU_RUN_PLACEHOLDER %t.out &> %t.gpu.txt || true +// RUN: %GPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.gpu.txt +// Shouldn't fail on ACC as fallback assert isn't enqueued there +// RUN: %ACC_RUN_PLACEHOLDER %t.out &> %t.acc.txt +// RUN: %ACC_RUN_PLACEHOLDER FileCheck %s --check-prefix=CHECK-ACC --input-file %t.acc.txt +// +// CHECK-NOT: this message from calculus +// CUDA uses block/thread vs global/local id for SYCL, also it shows the +// position of a thread within the block, not the absolute ID. +// CHECK: {{.*}}assert_in_multiple_tus.hpp:20: int checkFunction(): {{global id: \[5|block: \[1}},0,0], +// CHECK-SAME: {{.*}} [1,0,0] Assertion `X && "Nil in result"` failed. +// CHECK-NOT: this message from file2 +// CHECK-NOT: The test ended. +// +// CHECK-ACC-NOT: {{.*}}assert_in_multiple_tus.hpp:20: int checkFunction(): {{.*}} +// CHECK-ACC: The test ended. diff --git a/sycl/test-e2e/Assert/assert_in_multiple_tus_one_ndebug_win.cpp b/sycl/test-e2e/Assert/assert_in_multiple_tus_one_ndebug_win.cpp new file mode 100644 index 0000000000000..e2add269bd1d7 --- /dev/null +++ b/sycl/test-e2e/Assert/assert_in_multiple_tus_one_ndebug_win.cpp @@ -0,0 +1,20 @@ +// REQUIRES: windows +// RUN: %clangxx -DSYCL_FALLBACK_ASSERT=1 -fsycl -fsycl-targets=%sycl_triple -DDEFINE_NDEBUG_INFILE2 -I %S/Inputs %S/assert_in_multiple_tus.cpp %S/Inputs/kernels_in_file2.cpp -o %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out &> %t.txt || true +// RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt +// RUN: %GPU_RUN_PLACEHOLDER %t.out &> %t.txt || true +// RUN: %GPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt +// Shouldn't fail on ACC as fallback assert isn't enqueued there +// RUN: %ACC_RUN_PLACEHOLDER %t.out &> %t.txt || true +// RUN: %ACC_RUN_PLACEHOLDER FileCheck %s --check-prefix=CHECK-ACC --input-file %t.txt +// +// CHECK-NOT: this message from calculus +// FIXME Windows version prints '(null)' instead of '' once in a +// while for some insane reason. +// CHECK: {{.*}}assert_in_multiple_tus.hpp:20: {{|(null)}}: {{.*}} [5,0,0], +// CHECK-SAME: {{.*}} [1,0,0] Assertion `X && "Nil in result"` failed. +// CHECK-NOT: this message from file2 +// CHECK-NOT: The test ended. +// +// CHECK-ACC-NOT: {{.*}}assert_in_multiple_tus.hpp:20: {{|(null)}}: {{.*}} [5,0,0], +// CHECK-ACC: The test ended. diff --git a/sycl/test-e2e/Assert/assert_in_multiple_tus_win.cpp b/sycl/test-e2e/Assert/assert_in_multiple_tus_win.cpp new file mode 100644 index 0000000000000..1915bd8ed80f5 --- /dev/null +++ b/sycl/test-e2e/Assert/assert_in_multiple_tus_win.cpp @@ -0,0 +1,21 @@ +// REQUIRES: windows +// RUN: %clangxx -DSYCL_FALLBACK_ASSERT=1 -fsycl -fsycl-targets=%sycl_triple -I %S/Inputs %s %S/Inputs/kernels_in_file2.cpp -o %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out &> %t.txt || true +// RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt +// RUN: %GPU_RUN_PLACEHOLDER %t.out &> %t.txt || true +// RUN: %GPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt +// Shouldn't fail on ACC as fallback assert isn't enqueued there +// RUN: %ACC_RUN_PLACEHOLDER %t.out &> %t.txt +// RUN: %ACC_RUN_PLACEHOLDER FileCheck %s --check-prefix=CHECK-ACC --input-file %t.txt +// +// FIXME Windows version prints '(null)' instead of '' once in a +// while for some insane reason. +// CHECK: {{.*}}kernels_in_file2.cpp:15: {{|(null)}}: {{.*}} [5,0,0], {{.*}} [1,0,0] +// CHECK-SAME: Assertion `X && "this message from calculus"` failed. +// CHECK-NOT: this message from file2 +// CHECK-NOT: The test ended. +// +// CHECK-ACC-NOT: {{.*}}kernels_in_file2.cpp:15: {{|(null)}}: {{.*}} [5,0,0], {{.*}} [1,0,0] +// CHECK-ACC: The test ended. + +#include "assert_in_multiple_tus.hpp" diff --git a/sycl/test-e2e/Assert/assert_in_one_kernel.cpp b/sycl/test-e2e/Assert/assert_in_one_kernel.cpp new file mode 100644 index 0000000000000..f253de34ac1b4 --- /dev/null +++ b/sycl/test-e2e/Assert/assert_in_one_kernel.cpp @@ -0,0 +1,22 @@ +// REQUIRES: linux + +// https://github.com/intel/llvm/issues/7634 +// UNSUPPORTED: hip + +// RUN: %clangxx -DSYCL_FALLBACK_ASSERT=1 -fsycl -fsycl-targets=%sycl_triple %s -o %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out &> %t.cpu.txt || true +// RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.cpu.txt +// RUN: %GPU_RUN_PLACEHOLDER %t.out &> %t.gpu.txt || true +// RUN: %GPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.gpu.txt +// Shouldn't fail on ACC as fallback assert isn't enqueued there +// RUN: %ACC_RUN_PLACEHOLDER %t.out &> %t.acc.txt +// RUN: %ACC_RUN_PLACEHOLDER FileCheck %s --check-prefix=CHECK-ACC --input-file %t.acc.txt +// +// CHECK: {{.*}}assert_in_one_kernel.hpp:10: void kernelFunc(int *, int): {{.*}} [{{[0-3]}},0,0], {{.*}} [0,0,0] +// CHECK-SAME: Assertion `Buf[wiID] != 0 && "from assert statement"` failed. +// CHECK-NOT: The test ended. +// +// CHECK-ACC-NOT: {{.*}}assert_in_one_kernel.hpp:10: void kernelFunc(int *, int): {{.*}} [{{[0-3]}},0,0], {{.*}} [0,0,0] +// CHECK-ACC: The test ended. + +#include "assert_in_one_kernel.hpp" diff --git a/sycl/test-e2e/Assert/assert_in_one_kernel.hpp b/sycl/test-e2e/Assert/assert_in_one_kernel.hpp new file mode 100644 index 0000000000000..5f37916fdfdef --- /dev/null +++ b/sycl/test-e2e/Assert/assert_in_one_kernel.hpp @@ -0,0 +1,34 @@ +#include +#include +#include + +using namespace sycl; +using namespace sycl::access; + +void kernelFunc(int *Buf, int wiID) { + Buf[wiID] = 0; + assert(Buf[wiID] != 0 && "from assert statement"); +} + +void assertTest() { + std::array Vec = {1, 2, 3, 4}; + sycl::range<1> numOfItems{Vec.size()}; + sycl::buffer Buf(Vec.data(), numOfItems); + + queue Q; + Q.submit([&](handler &CGH) { + auto Acc = Buf.template get_access(CGH); + + CGH.parallel_for( + numOfItems, [=](item<1> Item) { kernelFunc(&Acc[0], Item[0]); }); + }); + Q.wait(); +} + +int main(int Argc, const char *Argv[]) { + + assertTest(); + + std::cout << "The test ended." << std::endl; + return 0; +} diff --git a/sycl/test-e2e/Assert/assert_in_one_kernel_ndebug.cpp b/sycl/test-e2e/Assert/assert_in_one_kernel_ndebug.cpp new file mode 100644 index 0000000000000..63aeddee9be9d --- /dev/null +++ b/sycl/test-e2e/Assert/assert_in_one_kernel_ndebug.cpp @@ -0,0 +1,7 @@ +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DNDEBUG %S/assert_in_one_kernel.cpp -o %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER +// RUN: %GPU_RUN_PLACEHOLDER %t.out %GPU_CHECK_PLACEHOLDER +// RUN: %ACC_RUN_PLACEHOLDER %t.out %ACC_CHECK_PLACEHOLDER +// +// CHECK-NOT: from assert statement +// CHECK: The test ended. diff --git a/sycl/test-e2e/Assert/assert_in_one_kernel_win.cpp b/sycl/test-e2e/Assert/assert_in_one_kernel_win.cpp new file mode 100644 index 0000000000000..9eed7fe0653fd --- /dev/null +++ b/sycl/test-e2e/Assert/assert_in_one_kernel_win.cpp @@ -0,0 +1,20 @@ +// REQUIRES: windows +// RUN: %clangxx -DSYCL_FALLBACK_ASSERT=1 -fsycl -fsycl-targets=%sycl_triple %s -o %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out &> %t.txt || true +// RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt +// RUN: %GPU_RUN_PLACEHOLDER %t.out &> %t.txt || true +// RUN: %GPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt +// Shouldn't fail on ACC as fallback assert isn't enqueued there +// RUN: %ACC_RUN_PLACEHOLDER %t.out &> %t.txt +// RUN: %ACC_RUN_PLACEHOLDER FileCheck %s --check-prefix=CHECK-ACC --input-file %t.txt +// +// FIXME Windows version prints '(null)' instead of '' once in a +// while for some insane reason. +// CHECK: {{.*}}assert_in_one_kernel.hpp:10: {{|(null)}}: {{.*}} [{{[0-3]}},0,0], {{.*}} [0,0,0] +// CHECK-SAME: Assertion `Buf[wiID] != 0 && "from assert statement"` failed. +// CHECK-NOT: The test ended. +// +// CHECK-ACC-NOT: {{.*}}assert_in_one_kernel.hpp:10: {{|(null)}}: {{.*}} [{{[0-3]}},0,0], {{.*}} [0,0,0] +// CHECK-ACC: The test ended. + +#include "assert_in_one_kernel.hpp" diff --git a/sycl/test-e2e/Assert/assert_in_simultaneous_kernels.cpp b/sycl/test-e2e/Assert/assert_in_simultaneous_kernels.cpp new file mode 100644 index 0000000000000..a70b3c654483d --- /dev/null +++ b/sycl/test-e2e/Assert/assert_in_simultaneous_kernels.cpp @@ -0,0 +1,28 @@ +// REQUIRES: linux +// FIXME: Flaky on HIP and cuda +// UNSUPPORTED: hip || cuda +// RUN: %clangxx -DSYCL_FALLBACK_ASSERT=1 -fsycl -fsycl-targets=%sycl_triple %s -o %t.out %threads_lib +// RUN: %CPU_RUN_PLACEHOLDER %t.out &> %t.txt || true +// RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt +// +// Since this is a multi-threaded application enable memory tracking and +// deferred release feature in the Level Zero plugin to avoid releasing memory +// too early. This is necessary because currently SYCL RT sets indirect access +// flag for all kernels and the Level Zero runtime doesn't support deferred +// release yet. +// Suppress runtime from printing out error messages, so that the test can +// match on assert message generated by the toolchains. +// RUN: env SYCL_PI_LEVEL_ZERO_TRACK_INDIRECT_ACCESS_MEMORY=1 SYCL_PI_SUPPRESS_ERROR_MESSAGE=1 %GPU_RUN_PLACEHOLDER %t.out &> %t.txt || true +// RUN: %GPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt +// Shouldn't fail on ACC as fallback assert isn't enqueued there +// RUN: %ACC_RUN_PLACEHOLDER %t.out &> %t.txt +// RUN: %ACC_RUN_PLACEHOLDER FileCheck %s --check-prefix=CHECK-ACC --input-file %t.txt +// +// CHECK: {{.*}}assert_in_simultaneous_kernels.hpp:13: void assertFunc(): {{.*}}[9,7,0], {{.*}}[0,0,0] +// CHECK-SAME: Assertion `false && "from assert statement"` failed. +// CHECK-NOT: The test ended. +// +// CHECK-ACC-NOT: {{.*}}assert_in_simultaneous_kernels.hpp:13: void assertFunc(): {{.*}} [9,7,0], {{.*}} [0,0,0] +// CHECK-ACC: The test ended. + +#include "assert_in_simultaneous_kernels.hpp" diff --git a/sycl/test-e2e/Assert/assert_in_simultaneous_kernels.hpp b/sycl/test-e2e/Assert/assert_in_simultaneous_kernels.hpp new file mode 100644 index 0000000000000..8f6bb1e999d24 --- /dev/null +++ b/sycl/test-e2e/Assert/assert_in_simultaneous_kernels.hpp @@ -0,0 +1,75 @@ +#include +#include +#include +#include +#include + +using namespace sycl; +using namespace sycl::access; + +static constexpr size_t NUM_THREADS = 4; +static constexpr size_t RANGE_SIZE = 1024; + +void assertFunc() { assert(false && "from assert statement"); } + +template void assertTest(queue *Q) { + Q->submit([&](handler &CGH) { + CGH.parallel_for( + nd_range<2>{{RANGE_SIZE, RANGE_SIZE}, {1, 1}}, [=](nd_item<2> it) { + if (it.get_global_id(0) == 7 && it.get_global_id(1) == 9) + assertFunc(); + }); + }); + Q->wait(); +} + +void runTestForTid(queue *Q, size_t Tid) { + switch (Tid % 4) { + case 0: { + assertTest(Q); + break; + } + case 1: { + assertTest(Q); + break; + } + case 2: { + assertTest(Q); + break; + } + case 3: { + assertTest(Q); + break; + } + } +} + +int main(int Argc, const char *Argv[]) { + // On windows stderr output becomes messed up if several thread + // output simultaneously. Hence, setting explicit line buffering here. +#ifndef __SYCL_DEVICE_ONLY__ + if (setvbuf(stderr, nullptr, _IOLBF, BUFSIZ)) { + std::cerr << "Can't set line-buffering mode fo stderr\n"; + return 1; + } +#endif + + std::vector threadPool; + threadPool.reserve(NUM_THREADS); + + std::vector> Queues; + for (size_t i = 0; i < NUM_THREADS; ++i) { + Queues.push_back(std::make_unique()); + } + + for (size_t tid = 0; tid < NUM_THREADS; ++tid) { + threadPool.push_back(std::thread(runTestForTid, Queues[tid].get(), tid)); + } + + for (auto ¤tThread : threadPool) { + currentThread.join(); + } + + std::cout << "The test ended." << std::endl; + return 0; +} diff --git a/sycl/test-e2e/Assert/assert_in_simultaneous_kernels_win.cpp b/sycl/test-e2e/Assert/assert_in_simultaneous_kernels_win.cpp new file mode 100644 index 0000000000000..647dc8ff3d244 --- /dev/null +++ b/sycl/test-e2e/Assert/assert_in_simultaneous_kernels_win.cpp @@ -0,0 +1,26 @@ +// REQUIRES: windows +// RUN: %clangxx -DSYCL_FALLBACK_ASSERT=1 -fsycl -fsycl-targets=%sycl_triple %s -o %t.out %threads_lib +// RUN: %CPU_RUN_PLACEHOLDER %t.out &> %t.txt || true +// RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt +// +// Since this is a multi-threaded application enable memory tracking and +// deferred release feature in the Level Zero plugin to avoid releasing memory +// too early. This is necessary because currently SYCL RT sets indirect access +// flag for all kernels and the Level Zero runtime doesn't support deferred +// release yet. +// RUN: env SYCL_PI_LEVEL_ZERO_TRACK_INDIRECT_ACCESS_MEMORY=1 %GPU_RUN_PLACEHOLDER %t.out &> %t.txt || true +// RUN: %GPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt +// Shouldn't fail on ACC as fallback assert isn't enqueued there +// RUN: %ACC_RUN_PLACEHOLDER %t.out &> %t.txt +// RUN: %ACC_RUN_PLACEHOLDER FileCheck %s --check-prefix=CHECK-ACC --input-file %t.txt +// +// FIXME Windows version prints '(null)' instead of '' once in a +// while for some insane reason. +// CHECK: {{.*}}assert_in_simultaneous_kernels.hpp:13: {{|(null)}}: global id: [9,7,0], local id: [0,0,0] +// CHECK-SAME: Assertion `false && "from assert statement"` failed. +// CHECK-NOT: The test ended. +// +// CHECK-ACC-NOT: {{.*}}assert_in_simultaneous_kernels.hpp:13: {{|(null)}}: global id: [9,7,0], local id: [0,0,0] +// CHECK-ACC: The test ended. + +#include "assert_in_simultaneous_kernels.hpp" diff --git a/sycl/test-e2e/Assert/assert_in_simultaneously_multiple_tus.cpp b/sycl/test-e2e/Assert/assert_in_simultaneously_multiple_tus.cpp new file mode 100644 index 0000000000000..75fb682044343 --- /dev/null +++ b/sycl/test-e2e/Assert/assert_in_simultaneously_multiple_tus.cpp @@ -0,0 +1,115 @@ +// FIXME flaky fail on CUDA and HIP +// UNSUPPORTED: cuda || hip +// +// RUN: %clangxx -DSYCL_FALLBACK_ASSERT=1 -fsycl -fsycl-targets=%sycl_triple -I %S/Inputs %s %S/Inputs/kernels_in_file2.cpp -o %t.out %threads_lib +// RUN: %CPU_RUN_PLACEHOLDER %t.out &> %t.txt || true +// RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt +// +// Since this is a multi-threaded application enable memory tracking and +// deferred release feature in the Level Zero plugin to avoid releasing memory +// too early. This is necessary because currently SYCL RT sets indirect access +// flag for all kernels and the Level Zero runtime doesn't support deferred +// release yet. +// Suppress runtime from printing out error messages, so that the test can +// match on assert message generated by the toolchains. +// RUN: env SYCL_PI_LEVEL_ZERO_TRACK_INDIRECT_ACCESS_MEMORY=1 SYCL_PI_SUPPRESS_ERROR_MESSAGE=1 %GPU_RUN_PLACEHOLDER %t.out &> %t.txt || true +// RUN: %GPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt +// Shouldn't fail on ACC as fallback assert isn't enqueued there +// RUN: %ACC_RUN_PLACEHOLDER %t.out &> %t.txt +// RUN: %ACC_RUN_PLACEHOLDER FileCheck %s --check-prefix=CHECK-ACC --input-file %t.txt +// +// CHECK: {{this message from file1|this message from file2}} +// CHECK-NOT: The test ended. +// +// CHECK-ACC-NOT: {{this message from file1|this message from file2}} +// CHECK-ACC: The test ended. + +#include "Inputs/kernels_in_file2.hpp" +#include +#include +#include +#include + +#ifdef DEFINE_NDEBUG_INFILE1 +#define NDEBUG +#else +#undef NDEBUG +#endif + +#include + +using namespace sycl; +using namespace sycl::access; + +static constexpr size_t NUM_THREADS = 4; +static constexpr size_t BUFFER_SIZE = 10; + +template void enqueueKernel(queue *Q) { + sycl::range<1> numOfItems{BUFFER_SIZE}; + sycl::buffer Buf(numOfItems); + + Q->submit([&](handler &CGH) { + auto Acc = Buf.template get_access(CGH); + + CGH.parallel_for(numOfItems, [=](sycl::id<1> wiID) { + Acc[wiID] = 0; + if (wiID == 5) + assert(false && "this message from file1"); + }); + }); +} + +void runTestForTid(queue *Q, size_t Tid) { + switch (Tid % 4) { + case 0: { + enqueueKernel(Q); + Q->wait(); + break; + } + case 1: { + enqueueKernel(Q); + Q->wait(); + break; + } + case 2: { + enqueueKernel_1_fromFile2(Q); + Q->wait(); + break; + } + case 3: { + enqueueKernel_2_fromFile2(Q); + Q->wait(); + break; + } + } +} + +int main(int Argc, const char *Argv[]) { +#ifndef __SYCL_DEVICE_ONLY__ + // On windows stderr output becomes messed up if several thread + // output simultaneously. Hence, setting explicit line buffering here. + if (setvbuf(stderr, nullptr, _IOLBF, BUFSIZ)) { + std::cerr << "Can't set line-buffering mode fo stderr\n"; + return 1; + } +#endif + + std::vector threadPool; + threadPool.reserve(NUM_THREADS); + + std::vector> Queues; + for (size_t i = 0; i < NUM_THREADS; ++i) { + Queues.push_back(std::make_unique()); + } + + for (size_t tid = 0; tid < NUM_THREADS; ++tid) { + threadPool.push_back(std::thread(runTestForTid, Queues[tid].get(), tid)); + } + + for (auto ¤tThread : threadPool) { + currentThread.join(); + } + + std::cout << "The test ended." << std::endl; + return 0; +} diff --git a/sycl/test-e2e/Assert/assert_in_simultaneously_multiple_tus_one_ndebug.cpp b/sycl/test-e2e/Assert/assert_in_simultaneously_multiple_tus_one_ndebug.cpp new file mode 100644 index 0000000000000..266896108f702 --- /dev/null +++ b/sycl/test-e2e/Assert/assert_in_simultaneously_multiple_tus_one_ndebug.cpp @@ -0,0 +1,26 @@ +// FIXME flaky fail on CUDA +// FIXME HIP: https://github.com/intel/llvm/issues/7634 +// UNSUPPORTED: cuda, hip + +// RUN: %clangxx -DSYCL_FALLBACK_ASSERT=1 -fsycl -fsycl-targets=%sycl_triple -DDEFINE_NDEBUG_INFILE2 -I %S/Inputs %S/assert_in_simultaneously_multiple_tus.cpp %S/Inputs/kernels_in_file2.cpp -o %t.out %threads_lib +// RUN: %CPU_RUN_PLACEHOLDER %t.out &> %t.cpu.txt || true +// RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.cpu.txt +// +// Since this is a multi-threaded application enable memory tracking and +// deferred release feature in the Level Zero plugin to avoid releasing memory +// too early. This is necessary because currently SYCL RT sets indirect access +// flag for all kernels and the Level Zero runtime doesn't support deferred +// release yet. +// Suppress runtime from printing out error messages, so that the test can +// match on assert message generated by the toolchains. +// RUN: env SYCL_PI_LEVEL_ZERO_TRACK_INDIRECT_ACCESS_MEMORY=1 SYCL_PI_SUPPRESS_ERROR_MESSAGE=1 %GPU_RUN_PLACEHOLDER %t.out &> %t.gpu.txt || true +// RUN: %GPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.gpu.txt +// Shouldn't fail on ACC as fallback assert isn't enqueued there +// RUN: %ACC_RUN_PLACEHOLDER %t.out &> %t.acc.txt +// RUN: %ACC_RUN_PLACEHOLDER FileCheck %s --check-prefix=CHECK-ACC --input-file %t.acc.txt +// +// CHECK: this message from file1 +// CHECK-NOT: this message from file2 +// CHECK-NOT: The test ended. +// +// CHECK-ACC: The test ended. diff --git a/sycl/test-e2e/AsyncHandler/default_async_handler.cpp b/sycl/test-e2e/AsyncHandler/default_async_handler.cpp new file mode 100644 index 0000000000000..15f5a56b56727 --- /dev/null +++ b/sycl/test-e2e/AsyncHandler/default_async_handler.cpp @@ -0,0 +1,24 @@ +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out &> %t.txt || true +// RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt +// RUN: %GPU_RUN_PLACEHOLDER %t.out &> %t.txt || true +// RUN: %GPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt +// RUN: %ACC_RUN_PLACEHOLDER %t.out &> %t.txt || true +// RUN: %ACC_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt + +#include + +using namespace sycl; + +int main() { + queue Q; + Q.submit([&](handler &CGH) { + CGH.host_task([=]() { + throw std::runtime_error("Exception thrown from host_task."); + }); + }).wait_and_throw(); + return 0; +} + +// CHECK: Default async_handler caught exceptions: +// CHECK-NEXT: Exception thrown from host_task. diff --git a/sycl/test-e2e/AtomicRef/accessor.cpp b/sycl/test-e2e/AtomicRef/accessor.cpp new file mode 100644 index 0000000000000..058d33e0d9078 --- /dev/null +++ b/sycl/test-e2e/AtomicRef/accessor.cpp @@ -0,0 +1,105 @@ +// RUN: %clangxx -fsycl -fsycl-unnamed-lambda -fsycl-targets=%sycl_triple %s -o %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out + +#include +#include +#include +#include +#include +#include +using namespace sycl; +using namespace sycl::ext::oneapi; + +// Equivalent to add_test from add.cpp +// Uses atomic_accessor instead of atomic_ref +template void accessor_test(queue q, size_t N) { + T sum = 0; + std::vector output(N, 0); + { + buffer sum_buf(&sum, 1); + buffer output_buf(output.data(), output.size()); + + q.submit([&](handler &cgh) { +#if __cplusplus > 201402L + static_assert( + std::is_same>::value, + "atomic_accessor type incorrectly deduced"); +#endif + auto sum = + atomic_accessor( + sum_buf, cgh); + auto out = + output_buf.template get_access(cgh); + cgh.parallel_for(range<1>(N), [=](item<1> it) { + int gid = it.get_id(0); + static_assert( + std::is_same>::value, + "atomic_accessor returns incorrect atomic_ref"); + out[gid] = sum[0].fetch_add(T(1)); + }); + }); + } + + // All work-items increment by 1, so final value should be equal to N + assert(sum == N); + + // Intermediate values should be unique + std::sort(output.begin(), output.end()); + assert(std::unique(output.begin(), output.end()) == output.end()); + + // Fetch returns original value: will be in [0, N-1] + auto min_e = output[0]; + auto max_e = output[output.size() - 1]; + assert(min_e == 0 && max_e == N - 1); +} + +// Simplified form of accessor_test for local memory +template +void local_accessor_test(queue q, size_t N, size_t L = 8) { + assert(N % L == 0); + std::vector output(N / L, 0); + { + buffer output_buf(output.data(), output.size()); + q.submit([&](handler &cgh) { + auto sum = + atomic_accessor(1, cgh); + auto out = output_buf.template get_access(cgh); + cgh.parallel_for(nd_range<1>(N, L), [=](nd_item<1> it) { + int grp = it.get_group(0); + sum[0].store(0); + it.barrier(); + static_assert( + std::is_same>::value, + "local atomic_accessor returns incorrect atomic_ref"); + T result = sum[0].fetch_add(T(1)); + if (result == it.get_local_range(0) - 1) { + out[grp] = result; + } + }); + }); + } + + // All work-items increment by 1, and last in the group writes out old value + // All values should be L-1 + assert(std::all_of(output.begin(), output.end(), + [=](T x) { return x == L - 1; })); +} + +int main() { + queue q; + constexpr int N = 32; + accessor_test(q, N); + local_accessor_test(q, N); + std::cout << "Test passed." << std::endl; +} diff --git a/sycl/test-e2e/AtomicRef/add.cpp b/sycl/test-e2e/AtomicRef/add.cpp new file mode 100644 index 0000000000000..909daea98595c --- /dev/null +++ b/sycl/test-e2e/AtomicRef/add.cpp @@ -0,0 +1,8 @@ +// RUN: %clangxx -fsycl -fsycl-unnamed-lambda -fsycl-device-code-split=per_kernel -fsycl-targets=%sycl_triple %s -o %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %ACC_RUN_PLACEHOLDER %t.out + +#include "add.h" + +int main() { add_test_all(); } diff --git a/sycl/test-e2e/AtomicRef/add.h b/sycl/test-e2e/AtomicRef/add.h new file mode 100644 index 0000000000000..48f24342202d2 --- /dev/null +++ b/sycl/test-e2e/AtomicRef/add.h @@ -0,0 +1,350 @@ +#pragma once + +#ifndef TEST_GENERIC_IN_LOCAL +#define TEST_GENERIC_IN_LOCAL 0 +#endif + +#include +#include +#include +#include +#include +#include +#include + +using namespace sycl; + +template