Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
Original file line number Diff line number Diff line change
Expand Up @@ -33,11 +33,12 @@ double SubmitBufferKernel(queue& q, std::vector<T>& in, std::vector<T>& 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);
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -79,8 +79,20 @@ 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<Type>(size, q.get_context());
Type* out_zero_copy = malloc_host<Type>(size, q.get_context());
#else
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Add a comment here similar to the one in this commit: https://github.com/oneapi-src/oneAPI-samples/pull/1890/files#diff-5bf4c46164e0fe924b5cd1ed04885a6c86ebe3710e8f3a89bead31e3e9cf3c78 that motivates why buffer location must be specified for IPA flow.

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

And add the same to buffer_kernel.hpp in the !IS_BSP clause.

// 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<Type>(
size, q,
sycl::ext::intel::experimental::property::usm::buffer_location(0));
Type *out_zero_copy = sycl::malloc_host<Type>(
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) {
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -60,10 +60,24 @@ using ConsumePipe = pipe<class ConsumePipeClass, T>;
//
template <typename T>
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>});
#endif

return q.single_task<Producer>([=]() [[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<T> h_in_data(in_data);
#endif

for (size_t i = 0; i < size; i++) {
T data_from_host_memory = *(h_in_data + i);
Expand Down Expand Up @@ -95,10 +109,24 @@ event SubmitWorker(queue& q, size_t size) {
//
template <typename T>
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>});
#endif

return q.single_task<Consumer>([=]() [[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<T> h_out_data(out_data);
#endif

for (size_t i = 0; i < size; i++) {
T data_to_host_memory = ConsumePipe<T>::read();
Expand Down