diff --git a/SYCL/Plugin/level_zero_eager_init.cpp b/SYCL/Plugin/level_zero_eager_init.cpp new file mode 100644 index 0000000000..9392c7ce2e --- /dev/null +++ b/SYCL/Plugin/level_zero_eager_init.cpp @@ -0,0 +1,57 @@ +// REQUIRES: level_zero, level_zero_dev_kit + +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %level_zero_options %s -o %t.out +// RUN: env ZE_DEBUG=1 SYCL_EAGER_INIT=1 %GPU_RUN_PLACEHOLDER %t.out 2>&1 %GPU_CHECK_PLACEHOLDER +// +// The test is to check that under SYCL_EAGER_INIT=1 there is no calls to +// heavy L0 initialization in the hot reportable path. +// +// CHECK-LABEL: HOT HOT HOT +// CHECK-NOT: zeCommandQueueCreate +// CHECK-NOT: zeCommandListCreate +// CHECK-NOT: zeFenceCreate +// + +#include + +#include +#include + +constexpr cl::sycl::access::mode sycl_read = cl::sycl::access::mode::read; +constexpr cl::sycl::access::mode sycl_write = cl::sycl::access::mode::write; + +/* This is the class used to name the kernel for the runtime. + * This must be done when the kernel is expressed as a lambda. */ +template class SimpleVadd; + +template +void simple_vadd(cl::sycl::queue &Queue, const std::array &VA, + const std::array &VB, std::array &VC) { + cl::sycl::range<1> numOfItems{N}; + cl::sycl::buffer bufferA(VA.data(), numOfItems); + cl::sycl::buffer bufferB(VB.data(), numOfItems); + cl::sycl::buffer bufferC(VC.data(), numOfItems); + + Queue.submit([&](cl::sycl::handler &cgh) { + auto accessorA = bufferA.template get_access(cgh); + auto accessorB = bufferB.template get_access(cgh); + auto accessorC = bufferC.template get_access(cgh); + + cgh.parallel_for>( + numOfItems, [=](cl::sycl::id<1> wiID) { + accessorC[wiID] = accessorA[wiID] + accessorB[wiID]; + }); + }); +} + +int main() { + const size_t array_size = 4; + std::array A = {{1, 2, 3, 4}}, + B = {{1, 2, 3, 4}}, C; + cl::sycl::queue Q; + + // simple_vadd(Q, A, B, C); + std::cerr << "\n\n\nHOT HOT HOT\n\n\n"; + simple_vadd(Q, A, B, C); + return 0; +}