From db42e369d8b12720231c0db4884ddd27f171b113 Mon Sep 17 00:00:00 2001 From: Mikhail Nikolskiy Date: Fri, 23 Apr 2021 07:03:25 -0700 Subject: [PATCH 01/10] [SYCL] E2E test fot interop_task with Level-Zero and OpenCL --- .../interop-level-zero-interop-task-mem.cpp | 65 +++++++++++++++++++ .../interop-opencl-interop-task-mem.cpp | 63 ++++++++++++++++++ 2 files changed, 128 insertions(+) create mode 100644 SYCL/Plugin/interop-level-zero-interop-task-mem.cpp create mode 100644 SYCL/Plugin/interop-opencl-interop-task-mem.cpp diff --git a/SYCL/Plugin/interop-level-zero-interop-task-mem.cpp b/SYCL/Plugin/interop-level-zero-interop-task-mem.cpp new file mode 100644 index 0000000000..f9caf93d87 --- /dev/null +++ b/SYCL/Plugin/interop-level-zero-interop-task-mem.cpp @@ -0,0 +1,65 @@ +// REQUIRES: level_zero, level_zero_dev_kit +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %level_zero_options %s -o %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out + +// Test for Level Zero interop_task. + +#include +// clang-format off +#include +#include +// clang-format on + +#define SIZE 16 + +class my_selector : public cl::sycl::device_selector { +public: + int operator()(const cl::sycl::device &dev) const override { + sycl::backend backend = dev.get_platform().get_backend(); + if (backend == cl::sycl::backend::level_zero && dev.is_gpu()) + return 1; + else + return 0; + } +}; + +int main() { + sycl::queue queue = sycl::queue(my_selector()); + + ze_context_handle_t ze_context = + queue.get_context().get_native(); + + try { + sycl::buffer buffer(SIZE); + sycl::image<2> image(sycl::image_channel_order::rgba, + sycl::image_channel_type::fp32, {SIZE, SIZE}); + + queue + .submit([&](cl::sycl::handler &cgh) { + auto buffer_acc = + buffer.get_access(cgh); + auto image_acc = + image.get_access(cgh); + cgh.interop_task([&](const cl::sycl::interop_handler &ih) { + void *device_ptr = + ih.get_mem(buffer_acc); + size_t size = 0; + zeMemGetAddressRange(ze_context, device_ptr, NULL, &size); + assert(size == SIZE); + + ze_image_handle_t ze_image = + ih.get_mem(image_acc); + assert(ze_image != nullptr); + }); + }) + .wait(); + } catch (cl::sycl::exception const &e) { + std::cout << "SYCL exception caught: " << e.what() << std::endl; + return e.get_cl_code(); + } catch (const char *msg) { + std::cout << "Exception caught: " << msg << std::endl; + return 1; + } + + return 0; +} diff --git a/SYCL/Plugin/interop-opencl-interop-task-mem.cpp b/SYCL/Plugin/interop-opencl-interop-task-mem.cpp new file mode 100644 index 0000000000..ba33f17442 --- /dev/null +++ b/SYCL/Plugin/interop-opencl-interop-task-mem.cpp @@ -0,0 +1,63 @@ +// REQUIRES: opencl + +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out + +// Test for OpenCL interop_task. + +#include +#include +#include + +#define SIZE 16 + +class my_selector : public cl::sycl::device_selector { +public: + int operator()(const cl::sycl::device &dev) const override { + sycl::backend backend = dev.get_platform().get_backend(); + if (backend == cl::sycl::backend::opencl && dev.is_gpu()) + return 1; + else + return 0; + } +}; + +int main() { + sycl::queue queue = sycl::queue(my_selector()); + + try { + sycl::buffer buffer(SIZE); + sycl::image<2> image(sycl::image_channel_order::rgba, + sycl::image_channel_type::fp32, {SIZE, SIZE}); + + queue + .submit([&](cl::sycl::handler &cgh) { + auto buffer_acc = + buffer.get_access(cgh); + auto image_acc = + image.get_access(cgh); + cgh.interop_task([&](const cl::sycl::interop_handler &ih) { + cl_mem buffer_mem = ih.get_mem(buffer_acc); + size_t size = 0; + clGetMemObjectInfo(buffer_mem, CL_MEM_SIZE, sizeof(size), + (void *)&size, nullptr); + assert(size == SIZE); + + cl_mem mem = ih.get_mem(image_acc); + size_t width = 0; + clGetImageInfo(mem, CL_IMAGE_WIDTH, sizeof(width), (void *)&width, + nullptr); + assert(width == SIZE); + }); + }) + .wait(); + } catch (cl::sycl::exception const &e) { + std::cout << "SYCL exception caught: " << e.what() << std::endl; + return e.get_cl_code(); + } catch (const char *msg) { + std::cout << "Exception caught: " << msg << std::endl; + return 1; + } + + return 0; +} From df07c9ebffd72c5de401fdd055ffebdf3f5591e8 Mon Sep 17 00:00:00 2001 From: Mikhail Nikolskii Date: Tue, 27 Apr 2021 00:29:44 +0300 Subject: [PATCH 02/10] Update interop-level-zero-interop-task-mem.cpp --- SYCL/Plugin/interop-level-zero-interop-task-mem.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/SYCL/Plugin/interop-level-zero-interop-task-mem.cpp b/SYCL/Plugin/interop-level-zero-interop-task-mem.cpp index f9caf93d87..372afaee50 100644 --- a/SYCL/Plugin/interop-level-zero-interop-task-mem.cpp +++ b/SYCL/Plugin/interop-level-zero-interop-task-mem.cpp @@ -40,7 +40,7 @@ int main() { buffer.get_access(cgh); auto image_acc = image.get_access(cgh); - cgh.interop_task([&](const cl::sycl::interop_handler &ih) { + cgh.interop_task([=](const cl::sycl::interop_handler &ih) { void *device_ptr = ih.get_mem(buffer_acc); size_t size = 0; From 07a90e8aea7ed8ec2c8623d331f28d4d863d69ab Mon Sep 17 00:00:00 2001 From: Mikhail Nikolskii Date: Tue, 27 Apr 2021 00:30:27 +0300 Subject: [PATCH 03/10] Update interop-opencl-interop-task-mem.cpp --- SYCL/Plugin/interop-opencl-interop-task-mem.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/SYCL/Plugin/interop-opencl-interop-task-mem.cpp b/SYCL/Plugin/interop-opencl-interop-task-mem.cpp index ba33f17442..2715a3f243 100644 --- a/SYCL/Plugin/interop-opencl-interop-task-mem.cpp +++ b/SYCL/Plugin/interop-opencl-interop-task-mem.cpp @@ -36,7 +36,7 @@ int main() { buffer.get_access(cgh); auto image_acc = image.get_access(cgh); - cgh.interop_task([&](const cl::sycl::interop_handler &ih) { + cgh.interop_task([=](const cl::sycl::interop_handler &ih) { cl_mem buffer_mem = ih.get_mem(buffer_acc); size_t size = 0; clGetMemObjectInfo(buffer_mem, CL_MEM_SIZE, sizeof(size), From 3e8fe24664121dedc3a196519d0bdc25e4c09768 Mon Sep 17 00:00:00 2001 From: Mikhail Nikolskiy Date: Wed, 28 Apr 2021 16:38:21 +0300 Subject: [PATCH 04/10] address review comments --- .../interop-level-zero-interop-task-mem.cpp | 46 +++++++------------ .../interop-opencl-interop-task-mem.cpp | 39 ++++++---------- 2 files changed, 31 insertions(+), 54 deletions(-) diff --git a/SYCL/Plugin/interop-level-zero-interop-task-mem.cpp b/SYCL/Plugin/interop-level-zero-interop-task-mem.cpp index f9caf93d87..ef58c023f5 100644 --- a/SYCL/Plugin/interop-level-zero-interop-task-mem.cpp +++ b/SYCL/Plugin/interop-level-zero-interop-task-mem.cpp @@ -1,6 +1,6 @@ // REQUIRES: level_zero, level_zero_dev_kit // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %level_zero_options %s -o %t.out -// RUN: %GPU_RUN_PLACEHOLDER %t.out +// RUN: env SYCL_BE=PI_LEVEL_ZERO %GPU_RUN_PLACEHOLDER %t.out // Test for Level Zero interop_task. @@ -10,50 +10,38 @@ #include // clang-format on -#define SIZE 16 +using namespace sycl; -class my_selector : public cl::sycl::device_selector { -public: - int operator()(const cl::sycl::device &dev) const override { - sycl::backend backend = dev.get_platform().get_backend(); - if (backend == cl::sycl::backend::level_zero && dev.is_gpu()) - return 1; - else - return 0; - } -}; +constexpr size_t SIZE = 16; int main() { - sycl::queue queue = sycl::queue(my_selector()); - - ze_context_handle_t ze_context = - queue.get_context().get_native(); + queue queue{}; try { - sycl::buffer buffer(SIZE); - sycl::image<2> image(sycl::image_channel_order::rgba, - sycl::image_channel_type::fp32, {SIZE, SIZE}); + buffer buffer(SIZE); + image<2> image(image_channel_order::rgba, image_channel_type::fp32, + {SIZE, SIZE}); + + ze_context_handle_t ze_context = + queue.get_context().get_native(); queue - .submit([&](cl::sycl::handler &cgh) { - auto buffer_acc = - buffer.get_access(cgh); - auto image_acc = - image.get_access(cgh); - cgh.interop_task([&](const cl::sycl::interop_handler &ih) { - void *device_ptr = - ih.get_mem(buffer_acc); + .submit([&](handler &cgh) { + auto buffer_acc = buffer.get_access(cgh); + auto image_acc = image.get_access(cgh); + cgh.interop_task([&](const interop_handler &ih) { + void *device_ptr = ih.get_mem(buffer_acc); size_t size = 0; zeMemGetAddressRange(ze_context, device_ptr, NULL, &size); assert(size == SIZE); ze_image_handle_t ze_image = - ih.get_mem(image_acc); + ih.get_mem(image_acc); assert(ze_image != nullptr); }); }) .wait(); - } catch (cl::sycl::exception const &e) { + } catch (exception const &e) { std::cout << "SYCL exception caught: " << e.what() << std::endl; return e.get_cl_code(); } catch (const char *msg) { diff --git a/SYCL/Plugin/interop-opencl-interop-task-mem.cpp b/SYCL/Plugin/interop-opencl-interop-task-mem.cpp index ba33f17442..f3d3886bb5 100644 --- a/SYCL/Plugin/interop-opencl-interop-task-mem.cpp +++ b/SYCL/Plugin/interop-opencl-interop-task-mem.cpp @@ -9,41 +9,30 @@ #include #include -#define SIZE 16 - -class my_selector : public cl::sycl::device_selector { -public: - int operator()(const cl::sycl::device &dev) const override { - sycl::backend backend = dev.get_platform().get_backend(); - if (backend == cl::sycl::backend::opencl && dev.is_gpu()) - return 1; - else - return 0; - } -}; +using namespace sycl; + +constexpr size_t SIZE = 16; int main() { - sycl::queue queue = sycl::queue(my_selector()); + queue queue{}; try { - sycl::buffer buffer(SIZE); - sycl::image<2> image(sycl::image_channel_order::rgba, - sycl::image_channel_type::fp32, {SIZE, SIZE}); + buffer buffer(SIZE); + image<2> image(image_channel_order::rgba, image_channel_type::fp32, + {SIZE, SIZE}); queue - .submit([&](cl::sycl::handler &cgh) { - auto buffer_acc = - buffer.get_access(cgh); - auto image_acc = - image.get_access(cgh); - cgh.interop_task([&](const cl::sycl::interop_handler &ih) { - cl_mem buffer_mem = ih.get_mem(buffer_acc); + .submit([&](handler &cgh) { + auto buffer_acc = buffer.get_access(cgh); + auto image_acc = image.get_access(cgh); + cgh.interop_task([=](const interop_handler &ih) { + cl_mem buffer_mem = ih.get_mem(buffer_acc); size_t size = 0; clGetMemObjectInfo(buffer_mem, CL_MEM_SIZE, sizeof(size), (void *)&size, nullptr); assert(size == SIZE); - cl_mem mem = ih.get_mem(image_acc); + cl_mem mem = ih.get_mem(image_acc); size_t width = 0; clGetImageInfo(mem, CL_IMAGE_WIDTH, sizeof(width), (void *)&width, nullptr); @@ -51,7 +40,7 @@ int main() { }); }) .wait(); - } catch (cl::sycl::exception const &e) { + } catch (exception const &e) { std::cout << "SYCL exception caught: " << e.what() << std::endl; return e.get_cl_code(); } catch (const char *msg) { From 7c083a51e9df6bad4c2f95c3092ac41104c491e4 Mon Sep 17 00:00:00 2001 From: Alexey Bader Date: Wed, 28 Apr 2021 21:08:20 +0300 Subject: [PATCH 05/10] Update SYCL/Plugin/interop-level-zero-interop-task-mem.cpp --- SYCL/Plugin/interop-level-zero-interop-task-mem.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/SYCL/Plugin/interop-level-zero-interop-task-mem.cpp b/SYCL/Plugin/interop-level-zero-interop-task-mem.cpp index ef58c023f5..cb4a28c254 100644 --- a/SYCL/Plugin/interop-level-zero-interop-task-mem.cpp +++ b/SYCL/Plugin/interop-level-zero-interop-task-mem.cpp @@ -1,6 +1,6 @@ // REQUIRES: level_zero, level_zero_dev_kit // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %level_zero_options %s -o %t.out -// RUN: env SYCL_BE=PI_LEVEL_ZERO %GPU_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out // Test for Level Zero interop_task. From b7a70839efb79d3d40811f5f192b3b9a52733cd2 Mon Sep 17 00:00:00 2001 From: Mikhail Nikolskii Date: Thu, 29 Apr 2021 16:57:38 +0300 Subject: [PATCH 06/10] OpenCL ICD loader is also needed to resolve direct OpenCL calls Co-authored-by: vladimirlaz --- SYCL/Plugin/interop-opencl-interop-task-mem.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/SYCL/Plugin/interop-opencl-interop-task-mem.cpp b/SYCL/Plugin/interop-opencl-interop-task-mem.cpp index f3d3886bb5..5ccf04443c 100644 --- a/SYCL/Plugin/interop-opencl-interop-task-mem.cpp +++ b/SYCL/Plugin/interop-opencl-interop-task-mem.cpp @@ -1,4 +1,4 @@ -// REQUIRES: opencl +// REQUIRES: opencl, opencl_icd // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out From b82c6495b505daa3fa205ab9fa3d3b4954b99da8 Mon Sep 17 00:00:00 2001 From: Mikhail Nikolskiy Date: Thu, 29 Apr 2021 19:41:57 +0300 Subject: [PATCH 07/10] add %opencl_lib and remove clang-format off/on --- SYCL/Plugin/interop-level-zero-interop-task-mem.cpp | 7 ++++--- SYCL/Plugin/interop-opencl-interop-task-mem.cpp | 2 +- 2 files changed, 5 insertions(+), 4 deletions(-) diff --git a/SYCL/Plugin/interop-level-zero-interop-task-mem.cpp b/SYCL/Plugin/interop-level-zero-interop-task-mem.cpp index cb4a28c254..f7f0e8bb72 100644 --- a/SYCL/Plugin/interop-level-zero-interop-task-mem.cpp +++ b/SYCL/Plugin/interop-level-zero-interop-task-mem.cpp @@ -4,11 +4,12 @@ // Test for Level Zero interop_task. -#include -// clang-format off +// Level-Zero #include + +// SYCL +#include #include -// clang-format on using namespace sycl; diff --git a/SYCL/Plugin/interop-opencl-interop-task-mem.cpp b/SYCL/Plugin/interop-opencl-interop-task-mem.cpp index 5ccf04443c..78390ac575 100644 --- a/SYCL/Plugin/interop-opencl-interop-task-mem.cpp +++ b/SYCL/Plugin/interop-opencl-interop-task-mem.cpp @@ -1,6 +1,6 @@ // REQUIRES: opencl, opencl_icd -// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out %opencl_lib // RUN: %GPU_RUN_PLACEHOLDER %t.out // Test for OpenCL interop_task. From 4239d5e87ef7400810440635b665133fc5262e49 Mon Sep 17 00:00:00 2001 From: Mikhail Nikolskiy Date: Fri, 9 Jul 2021 16:31:02 +0300 Subject: [PATCH 08/10] check ptr via zeMemGetAllocProperties --- SYCL/Plugin/interop-level-zero-interop-task-mem.cpp | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) diff --git a/SYCL/Plugin/interop-level-zero-interop-task-mem.cpp b/SYCL/Plugin/interop-level-zero-interop-task-mem.cpp index c927129836..8b8797b729 100644 --- a/SYCL/Plugin/interop-level-zero-interop-task-mem.cpp +++ b/SYCL/Plugin/interop-level-zero-interop-task-mem.cpp @@ -32,9 +32,10 @@ int main() { auto image_acc = image.get_access(cgh); cgh.interop_task([=](const interop_handler &ih) { void *device_ptr = ih.get_mem(buffer_acc); - size_t size = 0; - zeMemGetAddressRange(ze_context, device_ptr, NULL, &size); - assert(size == SIZE); + ze_memory_allocation_properties_t memAllocProperties{}; + ze_result_t res = zeMemGetAllocProperties(ze_context, device_ptr, &memAllocProperties, + nullptr); + assert(res == ZE_RESULT_SUCCESS); ze_image_handle_t ze_image = ih.get_mem(image_acc); From 4421eb1a6bf70676b73536afd2f391e6e5702e37 Mon Sep 17 00:00:00 2001 From: Mikhail Nikolskiy Date: Fri, 9 Jul 2021 16:32:16 +0300 Subject: [PATCH 09/10] clang-format --- SYCL/Plugin/interop-level-zero-interop-task-mem.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/SYCL/Plugin/interop-level-zero-interop-task-mem.cpp b/SYCL/Plugin/interop-level-zero-interop-task-mem.cpp index 8b8797b729..73c19b97d2 100644 --- a/SYCL/Plugin/interop-level-zero-interop-task-mem.cpp +++ b/SYCL/Plugin/interop-level-zero-interop-task-mem.cpp @@ -32,7 +32,7 @@ int main() { auto image_acc = image.get_access(cgh); cgh.interop_task([=](const interop_handler &ih) { void *device_ptr = ih.get_mem(buffer_acc); - ze_memory_allocation_properties_t memAllocProperties{}; + ze_memory_allocation_properties_t memAllocProperties{}; ze_result_t res = zeMemGetAllocProperties(ze_context, device_ptr, &memAllocProperties, nullptr); assert(res == ZE_RESULT_SUCCESS); From c1afbee61cf79ae1b52793f9c34d1442ad2c182c Mon Sep 17 00:00:00 2001 From: Mikhail Nikolskiy Date: Fri, 9 Jul 2021 20:31:44 +0300 Subject: [PATCH 10/10] clang-format --- SYCL/Plugin/interop-level-zero-interop-task-mem.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/SYCL/Plugin/interop-level-zero-interop-task-mem.cpp b/SYCL/Plugin/interop-level-zero-interop-task-mem.cpp index 73c19b97d2..5459740296 100644 --- a/SYCL/Plugin/interop-level-zero-interop-task-mem.cpp +++ b/SYCL/Plugin/interop-level-zero-interop-task-mem.cpp @@ -33,9 +33,9 @@ int main() { cgh.interop_task([=](const interop_handler &ih) { void *device_ptr = ih.get_mem(buffer_acc); ze_memory_allocation_properties_t memAllocProperties{}; - ze_result_t res = zeMemGetAllocProperties(ze_context, device_ptr, &memAllocProperties, - nullptr); - assert(res == ZE_RESULT_SUCCESS); + ze_result_t res = zeMemGetAllocProperties( + ze_context, device_ptr, &memAllocProperties, nullptr); + assert(res == ZE_RESULT_SUCCESS); ze_image_handle_t ze_image = ih.get_mem(image_acc);