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<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);
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..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
@@ -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
+    // 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) {
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..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
@@ -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);
@@ -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();