From 4eca0b7669c6e1d43d487df82b67a0abdcb1f57d Mon Sep 17 00:00:00 2001 From: "Wang, Di5" Date: Thu, 21 Sep 2023 06:05:06 -0700 Subject: [PATCH 1/4] Update zero_copy_kernel sample --- .../src/zero_copy_data_transfer.cpp | 9 +++++++++ .../src/zero_copy_kernel.hpp | 16 ++++++++++++++++ 2 files changed, 25 insertions(+) diff --git a/DirectProgramming/C++SYCL_FPGA/Tutorials/DesignPatterns/zero_copy_data_transfer/src/zero_copy_data_transfer.cpp b/DirectProgramming/C++SYCL_FPGA/Tutorials/DesignPatterns/zero_copy_data_transfer/src/zero_copy_data_transfer.cpp index 2da35473d8..034caca310 100644 --- a/DirectProgramming/C++SYCL_FPGA/Tutorials/DesignPatterns/zero_copy_data_transfer/src/zero_copy_data_transfer.cpp +++ b/DirectProgramming/C++SYCL_FPGA/Tutorials/DesignPatterns/zero_copy_data_transfer/src/zero_copy_data_transfer.cpp @@ -79,8 +79,17 @@ int main(int argc, char* argv[]) { // input and output data for the zero-copy version // malloc_host allocates memory specifically in the host's address space +#if defined(IS_BSP) Type* in_zero_copy = malloc_host(size, q.get_context()); Type* out_zero_copy = malloc_host(size, q.get_context()); +#else + Type *in_zero_copy = sycl::malloc_host( + size, q, + sycl::ext::intel::experimental::property::usm::buffer_location(0)); + Type *out_zero_copy = sycl::malloc_host( + size, q, + sycl::ext::intel::experimental::property::usm::buffer_location(0)); +#endif // ensure that we could allocate space for both the input and output if (in_zero_copy == NULL) { diff --git a/DirectProgramming/C++SYCL_FPGA/Tutorials/DesignPatterns/zero_copy_data_transfer/src/zero_copy_kernel.hpp b/DirectProgramming/C++SYCL_FPGA/Tutorials/DesignPatterns/zero_copy_data_transfer/src/zero_copy_kernel.hpp index f26fa879e5..17309c9278 100644 --- a/DirectProgramming/C++SYCL_FPGA/Tutorials/DesignPatterns/zero_copy_data_transfer/src/zero_copy_kernel.hpp +++ b/DirectProgramming/C++SYCL_FPGA/Tutorials/DesignPatterns/zero_copy_data_transfer/src/zero_copy_kernel.hpp @@ -60,10 +60,18 @@ using ConsumePipe = pipe; // template event SubmitProducer(queue& q, T* in_data, size_t size) { +#if !defined(IS_BSP) + sycl::ext::oneapi::experimental::annotated_arg h_in_data( + in_data, sycl::ext::oneapi::experimental::properties{ + sycl::ext::intel::experimental::buffer_location<0>}); +#endif + return q.single_task([=]() [[intel::kernel_args_restrict]] { +#if defined(IS_BSP) // using a host_ptr tells the compiler that this pointer lives in the // hosts address space host_ptr h_in_data(in_data); +#endif for (size_t i = 0; i < size; i++) { T data_from_host_memory = *(h_in_data + i); @@ -95,10 +103,18 @@ event SubmitWorker(queue& q, size_t size) { // template event SubmitConsumer(queue& q, T* out_data, size_t size) { +#if !defined(IS_BSP) + sycl::ext::oneapi::experimental::annotated_arg h_out_data( + out_data, sycl::ext::oneapi::experimental::properties{ + sycl::ext::intel::experimental::buffer_location<0>}); +#endif + return q.single_task([=]() [[intel::kernel_args_restrict]] { +#if defined(IS_BSP) // using a host_ptr tells the compiler that this pointer lives in the // hosts address space host_ptr h_out_data(out_data); +#endif for (size_t i = 0; i < size; i++) { T data_to_host_memory = ConsumePipe::read(); From 4097fe7bf2c77d3a36e0bab97bea1530e43062c0 Mon Sep 17 00:00:00 2001 From: "Wang, Di5" Date: Thu, 21 Sep 2023 13:10:36 -0700 Subject: [PATCH 2/4] update comments --- .../zero_copy_data_transfer/src/buffer_kernel.hpp | 9 +++++---- .../src/zero_copy_data_transfer.cpp | 6 ++++++ .../zero_copy_data_transfer/src/zero_copy_kernel.hpp | 6 ++++++ 3 files changed, 17 insertions(+), 4 deletions(-) diff --git a/DirectProgramming/C++SYCL_FPGA/Tutorials/DesignPatterns/zero_copy_data_transfer/src/buffer_kernel.hpp b/DirectProgramming/C++SYCL_FPGA/Tutorials/DesignPatterns/zero_copy_data_transfer/src/buffer_kernel.hpp index b06c2b0886..ee5428976f 100644 --- a/DirectProgramming/C++SYCL_FPGA/Tutorials/DesignPatterns/zero_copy_data_transfer/src/buffer_kernel.hpp +++ b/DirectProgramming/C++SYCL_FPGA/Tutorials/DesignPatterns/zero_copy_data_transfer/src/buffer_kernel.hpp @@ -33,11 +33,12 @@ double SubmitBufferKernel(queue& q, std::vector& in, std::vector& out, accessor in_a(in_buf, h, read_only); accessor out_a(out_buf, h, write_only, no_init); #else - // When targeting an FPGA family/part, the compiler does not know - // if the two kernels accesses the same memory location + // When targeting an FPGA family/part, the compiler infers memory + // interfaces based on the unique buffer_location property specified + // on kernel arguments // With this property, we tell the compiler that these buffers - // are in a location "1" whereas the pointers from ExplicitKernel - // are in the default location "0" + // are in a location "1" whereas the pointers from ZeroCopyKernel + // are in the location "0" sycl::ext::oneapi::accessor_property_list location_of_buffer{ ext::intel::buffer_location<1>}; accessor in_a(in_buf, h, read_only, location_of_buffer); diff --git a/DirectProgramming/C++SYCL_FPGA/Tutorials/DesignPatterns/zero_copy_data_transfer/src/zero_copy_data_transfer.cpp b/DirectProgramming/C++SYCL_FPGA/Tutorials/DesignPatterns/zero_copy_data_transfer/src/zero_copy_data_transfer.cpp index 034caca310..66ad416740 100644 --- a/DirectProgramming/C++SYCL_FPGA/Tutorials/DesignPatterns/zero_copy_data_transfer/src/zero_copy_data_transfer.cpp +++ b/DirectProgramming/C++SYCL_FPGA/Tutorials/DesignPatterns/zero_copy_data_transfer/src/zero_copy_data_transfer.cpp @@ -83,6 +83,12 @@ int main(int argc, char* argv[]) { Type* in_zero_copy = malloc_host(size, q.get_context()); Type* out_zero_copy = malloc_host(size, q.get_context()); #else + // When targeting an FPGA family/part, the compiler infers memory + // interfaces based on the unique buffer_location property specified + // on kernel arguments + // The USM pointers passed into the kernel must match with the buffer location + // of each kernel argument, i.e. the allocated host memory should be in + // the location "0", as requested in ZeroCopyKernel Type *in_zero_copy = sycl::malloc_host( size, q, sycl::ext::intel::experimental::property::usm::buffer_location(0)); diff --git a/DirectProgramming/C++SYCL_FPGA/Tutorials/DesignPatterns/zero_copy_data_transfer/src/zero_copy_kernel.hpp b/DirectProgramming/C++SYCL_FPGA/Tutorials/DesignPatterns/zero_copy_data_transfer/src/zero_copy_kernel.hpp index 17309c9278..29615d8ab6 100644 --- a/DirectProgramming/C++SYCL_FPGA/Tutorials/DesignPatterns/zero_copy_data_transfer/src/zero_copy_kernel.hpp +++ b/DirectProgramming/C++SYCL_FPGA/Tutorials/DesignPatterns/zero_copy_data_transfer/src/zero_copy_kernel.hpp @@ -61,6 +61,12 @@ using ConsumePipe = pipe; template event SubmitProducer(queue& q, T* in_data, size_t size) { #if !defined(IS_BSP) + // When targeting an FPGA family/part, the compiler infers memory + // interfaces based on the unique buffer_location property specified + // on kernel arguments + // With this property, we tell the compiler that these buffers + // are in a location "0" whereas the pointers from BufferKernel + // are in the location "1" sycl::ext::oneapi::experimental::annotated_arg h_in_data( in_data, sycl::ext::oneapi::experimental::properties{ sycl::ext::intel::experimental::buffer_location<0>}); From 49d2e1f8a4d154ecc93f3b72cb45c61cdbe5cffe Mon Sep 17 00:00:00 2001 From: "Wang, Di5" Date: Thu, 21 Sep 2023 13:37:40 -0700 Subject: [PATCH 3/4] tweak --- .../src/zero_copy_data_transfer.cpp | 9 +++------ 1 file changed, 3 insertions(+), 6 deletions(-) diff --git a/DirectProgramming/C++SYCL_FPGA/Tutorials/DesignPatterns/zero_copy_data_transfer/src/zero_copy_data_transfer.cpp b/DirectProgramming/C++SYCL_FPGA/Tutorials/DesignPatterns/zero_copy_data_transfer/src/zero_copy_data_transfer.cpp index 66ad416740..6ad2c67391 100644 --- a/DirectProgramming/C++SYCL_FPGA/Tutorials/DesignPatterns/zero_copy_data_transfer/src/zero_copy_data_transfer.cpp +++ b/DirectProgramming/C++SYCL_FPGA/Tutorials/DesignPatterns/zero_copy_data_transfer/src/zero_copy_data_transfer.cpp @@ -83,12 +83,9 @@ int main(int argc, char* argv[]) { Type* in_zero_copy = malloc_host(size, q.get_context()); Type* out_zero_copy = malloc_host(size, q.get_context()); #else - // When targeting an FPGA family/part, the compiler infers memory - // interfaces based on the unique buffer_location property specified - // on kernel arguments - // The USM pointers passed into the kernel must match with the buffer location - // of each kernel argument, i.e. the allocated host memory should be in - // the location "0", as requested in ZeroCopyKernel + // The USM pointers passed into the kernel must be allocated with the same + // buffer_location as the one specified on the kernel argument with the + // annotated_arg class. Type *in_zero_copy = sycl::malloc_host( size, q, sycl::ext::intel::experimental::property::usm::buffer_location(0)); From fbe8dfc2f40c950f90c8adf5fde40b85eaf99eaf Mon Sep 17 00:00:00 2001 From: "Wang, Di5" Date: Fri, 22 Sep 2023 08:18:35 -0700 Subject: [PATCH 4/4] small update on comment --- .../zero_copy_data_transfer/src/zero_copy_kernel.hpp | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/DirectProgramming/C++SYCL_FPGA/Tutorials/DesignPatterns/zero_copy_data_transfer/src/zero_copy_kernel.hpp b/DirectProgramming/C++SYCL_FPGA/Tutorials/DesignPatterns/zero_copy_data_transfer/src/zero_copy_kernel.hpp index 29615d8ab6..a2c9f1d83c 100644 --- a/DirectProgramming/C++SYCL_FPGA/Tutorials/DesignPatterns/zero_copy_data_transfer/src/zero_copy_kernel.hpp +++ b/DirectProgramming/C++SYCL_FPGA/Tutorials/DesignPatterns/zero_copy_data_transfer/src/zero_copy_kernel.hpp @@ -110,6 +110,12 @@ event SubmitWorker(queue& q, size_t size) { template event SubmitConsumer(queue& q, T* out_data, size_t size) { #if !defined(IS_BSP) + // When targeting an FPGA family/part, the compiler infers memory + // interfaces based on the unique buffer_location property specified + // on kernel arguments + // With this property, we tell the compiler that these buffers + // are in a location "0" whereas the pointers from BufferKernel + // are in the location "1" sycl::ext::oneapi::experimental::annotated_arg h_out_data( out_data, sycl::ext::oneapi::experimental::properties{ sycl::ext::intel::experimental::buffer_location<0>});