Skip to content

[FPGA] Update zero_copy_kernel sample to support IP-Authoring #11

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
wangdi4 wants to merge 4 commits into
base: development
Choose a base branch
from
Open

[FPGA] Update zero_copy_kernel sample to support IP-Authoring #11

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
4eca0b7
Update zero_copy_kernel sample
wangdi4 Sep 21, 2023
4097fe7
update comments
wangdi4 Sep 21, 2023
49d2e1f
tweak
wangdi4 Sep 21, 2023
fbe8dfc
small update on comment
wangdi4 Sep 22, 2023
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
9 changes: 9 additions & 0 deletions ...YCL_FPGA/Tutorials/DesignPatterns/zero_copy_data_transfer/src/zero_copy_data_transfer.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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<Type>(size, q.get_context());
Type* out_zero_copy = malloc_host<Type>(size, q.get_context());
#else
Copy link
Collaborator

@tiwaria1 tiwaria1 Sep 21, 2023

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

@tiwaria1 tiwaria1 Sep 21, 2023

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.

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
16 changes: 16 additions & 0 deletions ...ng/C++SYCL_FPGA/Tutorials/DesignPatterns/zero_copy_data_transfer/src/zero_copy_kernel.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -60,10 +60,18 @@ using ConsumePipe = pipe<class ConsumePipeClass, T>;
//
template <typename T>
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<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 +103,18 @@ event SubmitWorker(queue& q, size_t size) {
//
template <typename T>
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<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