diff --git a/sycl/test-e2e/Basic/enqueue_barrier.cpp b/sycl/test-e2e/Basic/enqueue_barrier.cpp deleted file mode 100644 index 9793c83e73e14..0000000000000 --- a/sycl/test-e2e/Basic/enqueue_barrier.cpp +++ /dev/null @@ -1,62 +0,0 @@ -// RUN: %{build} -o %t.out -// RUN: env SYCL_UR_TRACE=2 %{run} %t.out 2>&1 | FileCheck %s - -#include -#include - -int main() { - sycl::context Context; - sycl::queue Q1(Context, sycl::default_selector_v); - - Q1.submit( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); - Q1.submit( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); - - // call handler::barrier() - Q1.submit([&](sycl::handler &cgh) { cgh.ext_oneapi_barrier(); }); - - Q1.submit( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); - Q1.submit( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); - - // call queue::ext_oneapi_submit_barrier() - Q1.ext_oneapi_submit_barrier(); - - sycl::queue Q2(Context, sycl::default_selector_v); - sycl::queue Q3(Context, sycl::default_selector_v); - - auto Event1 = Q1.submit( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); - - auto Event2 = Q2.submit( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); - - // call handler::barrier(const std::vector &WaitList) - Q3.submit([&](sycl::handler &cgh) { - cgh.ext_oneapi_barrier({Event1, Event2}); - }); - - Q3.submit( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); - - auto Event3 = Q1.submit( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); - - auto Event4 = Q2.submit( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); - - // call queue::ext_oneapi_submit_barrier(const std::vector &WaitList) - Q3.ext_oneapi_submit_barrier({Event3, Event4}); - - Q3.submit( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); - - return 0; -} - -// CHECK: <--- urEnqueueEventsWaitWithBarrier -// CHECK: <--- urEnqueueEventsWaitWithBarrier -// CHECK: <--- urEnqueueEventsWaitWithBarrier -// CHECK: <--- urEnqueueEventsWaitWithBarrier diff --git a/sycl/test-e2e/ESIMD/esimd_check_vc_codegen.cpp b/sycl/test-e2e/ESIMD/esimd_check_vc_codegen.cpp deleted file mode 100644 index 9576916b0e8bc..0000000000000 --- a/sycl/test-e2e/ESIMD/esimd_check_vc_codegen.cpp +++ /dev/null @@ -1,97 +0,0 @@ -//==-------- esimd_check_vc_codegen.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 -// -//===----------------------------------------------------------------------===// -// RUN: %{build} -o %t.out -// RUN: env SYCL_UR_TRACE=2 %{run} %t.out 2>&1 | FileCheck %s - -#include "esimd_test_utils.hpp" - -using namespace sycl; - -int main(void) { - constexpr unsigned Size = 1024 * 128; - constexpr unsigned VL = 16; - - float *A = new float[Size]; - float *B = new float[Size]; - float *C = new float[Size]; - - for (unsigned i = 0; i < Size; ++i) { - A[i] = B[i] = i; - C[i] = 0.0f; - } - - try { - buffer bufa(A, range<1>(Size)); - buffer bufb(B, range<1>(Size)); - buffer bufc(C, range<1>(Size)); - - // We need that many workgroups - range<1> GlobalRange{Size / VL}; - - // We need that many threads in each group - 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); - auto PB = bufb.get_access(cgh); - auto PC = bufc.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 vb; - vb.copy_from(PB, offset); - simd vc = va + vb; - vc.copy_to(PC, offset); - }); - }); - e.wait(); - } catch (sycl::exception const &e) { - std::cout << "SYCL exception caught: " << e.what() << '\n'; - - delete[] A; - delete[] B; - delete[] C; - return 1; - } - - int err_cnt = 0; - - for (unsigned i = 0; i < Size; ++i) { - if (A[i] + B[i] != C[i]) { - if (++err_cnt < 10) { - std::cout << "failed at index " << i << ", " << C[i] << " != " << A[i] - << " + " << B[i] << "\n"; - } - } - } - if (err_cnt > 0) { - std::cout << " pass rate: " - << ((float)(Size - err_cnt) / (float)Size) * 100.0f << "% (" - << (Size - err_cnt) << "/" << Size << ")\n"; - } - - delete[] A; - delete[] B; - delete[] C; - - std::cout << (err_cnt > 0 ? "FAILED\n" : "Passed\n"); - return err_cnt > 0 ? 1 : 0; -} - -// Don't use -NEXT here to split the line because we need to allow for the -// possbility of a BuildExp( that fails with UNSUPPORTED followed by a Build( -// that results in SUCCESS -// CHECK: <--- urProgramBuild{{(Exp)?}}({{.*}}-vc-codegen{{.*}} -> UR_RESULT_SUCCESS diff --git a/sycl/test-e2e/ProgramManager/multi_device_bundle/compile_link.cpp b/sycl/test-e2e/ProgramManager/multi_device_bundle/compile_link.cpp index 5589985e96d0e..1458f2d32c8fc 100644 --- a/sycl/test-e2e/ProgramManager/multi_device_bundle/compile_link.cpp +++ b/sycl/test-e2e/ProgramManager/multi_device_bundle/compile_link.cpp @@ -1,7 +1,7 @@ // REQUIRES: gpu && linux && (opencl || level_zero) // RUN: %{build} -o %t.out -// RUN: env NEOReadDebugKeys=1 CreateMultipleRootDevices=3 SYCL_UR_TRACE=2 %{run} %t.out +// RUN: env NEOReadDebugKeys=1 CreateMultipleRootDevices=3 %{run} %t.out // Test to check that we can compile and link a kernel bundle for multiple // devices and run the kernel on each device. diff --git a/sycl/unittests/kernel-and-program/KernelBuildOptions.cpp b/sycl/unittests/kernel-and-program/KernelBuildOptions.cpp index f994c62578c69..fdf1ff1e32746 100644 --- a/sycl/unittests/kernel-and-program/KernelBuildOptions.cpp +++ b/sycl/unittests/kernel-and-program/KernelBuildOptions.cpp @@ -119,3 +119,20 @@ TEST(KernelBuildOptions, KernelBundleBasic) { auto LinkBundle = sycl::link(ObjBundle, ObjBundle.get_devices()); EXPECT_EQ(BuildOpts, "-link-img"); } + +TEST(KernelBuildOptions, ESIMDParallelForBasic) { + sycl::unittest::UrMock<> Mock; + sycl::platform Plt = sycl::platform(); + setupCommonMockAPIs(Mock); + + const sycl::device Dev = Plt.get_devices()[0]; + sycl::queue Queue{Dev}; + + Queue.submit([&](sycl::handler &cgh) { + cgh.parallel_for( + sycl::range{1024}, [=](sycl::id<1>) /* SYCL_ESIMD_KERNEL */ {}); + }); + + EXPECT_EQ(BuildOpts, + "-compile-img -vc-codegen -disable-finalizer-msg -link-img"); +} diff --git a/sycl/unittests/queue/Barrier.cpp b/sycl/unittests/queue/Barrier.cpp new file mode 100644 index 0000000000000..a1e2160ba606f --- /dev/null +++ b/sycl/unittests/queue/Barrier.cpp @@ -0,0 +1,100 @@ +//==------------------- Barrier.cpp --- queue unit tests -------------------==// +// +// 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 +// +//===----------------------------------------------------------------------===// + +#include +#include + +#include + +static unsigned NumOfEventsWaitWithBarrierCalls = 0; + +static ur_result_t redefined_urEnqueueEventsWaitWithBarrier(void *) { + NumOfEventsWaitWithBarrierCalls++; + + return UR_RESULT_SUCCESS; +} + +TEST(Queue, HandlerBarrier) { + sycl::unittest::UrMock<> Mock; + mock::getCallbacks().set_before_callback( + "urEnqueueEventsWaitWithBarrier", + &redefined_urEnqueueEventsWaitWithBarrier); + NumOfEventsWaitWithBarrierCalls = 0; + + sycl::queue Q; + + Q.submit( + [&](sycl::handler &cgh) { cgh.single_task>([=]() {}); }); + Q.submit( + [&](sycl::handler &cgh) { cgh.single_task>([=]() {}); }); + + Q.submit([&](sycl::handler &cgh) { cgh.ext_oneapi_barrier(); }); + + ASSERT_EQ(NumOfEventsWaitWithBarrierCalls, 1u); +} + +TEST(Queue, ExtOneAPISubmitBarrier) { + sycl::unittest::UrMock<> Mock; + mock::getCallbacks().set_before_callback( + "urEnqueueEventsWaitWithBarrier", + &redefined_urEnqueueEventsWaitWithBarrier); + NumOfEventsWaitWithBarrierCalls = 0; + + sycl::queue Q; + + Q.submit( + [&](sycl::handler &cgh) { cgh.single_task>([=]() {}); }); + Q.submit( + [&](sycl::handler &cgh) { cgh.single_task>([=]() {}); }); + + Q.ext_oneapi_submit_barrier(); + + ASSERT_EQ(NumOfEventsWaitWithBarrierCalls, 1u); +} + +TEST(Queue, HandlerBarrierWithWaitList) { + sycl::unittest::UrMock<> Mock; + mock::getCallbacks().set_before_callback( + "urEnqueueEventsWaitWithBarrier", + &redefined_urEnqueueEventsWaitWithBarrier); + NumOfEventsWaitWithBarrierCalls = 0; + + sycl::queue Q1; + sycl::queue Q2; + sycl::queue Q3; + + auto E1 = Q1.submit( + [&](sycl::handler &cgh) { cgh.single_task>([=]() {}); }); + auto E2 = Q2.submit( + [&](sycl::handler &cgh) { cgh.single_task>([=]() {}); }); + + Q3.submit([&](sycl::handler &cgh) { cgh.ext_oneapi_barrier({E1, E2}); }); + + ASSERT_EQ(NumOfEventsWaitWithBarrierCalls, 1u); +} + +TEST(Queue, ExtOneAPISubmitBarrierWithWaitList) { + sycl::unittest::UrMock<> Mock; + mock::getCallbacks().set_before_callback( + "urEnqueueEventsWaitWithBarrier", + &redefined_urEnqueueEventsWaitWithBarrier); + NumOfEventsWaitWithBarrierCalls = 0; + + sycl::queue Q1; + sycl::queue Q2; + sycl::queue Q3; + + auto E1 = Q1.submit( + [&](sycl::handler &cgh) { cgh.single_task>([=]() {}); }); + auto E2 = Q2.submit( + [&](sycl::handler &cgh) { cgh.single_task>([=]() {}); }); + + Q3.ext_oneapi_submit_barrier({E1, E2}); + + ASSERT_EQ(NumOfEventsWaitWithBarrierCalls, 1u); +} diff --git a/sycl/unittests/queue/CMakeLists.txt b/sycl/unittests/queue/CMakeLists.txt index 07a5601c359c6..844e2509337fe 100644 --- a/sycl/unittests/queue/CMakeLists.txt +++ b/sycl/unittests/queue/CMakeLists.txt @@ -9,4 +9,5 @@ add_sycl_unittest(QueueTests OBJECT InOrderQueue.cpp InteropRetain.cpp Properties.cpp + Barrier.cpp )