[SYCL][CUDA] unnecessary memcpy for write buffers.

[Ryefalk]

When using write and discard_write accessors with the CUDA backend a Host to Device copy is made. Even though this is supposed to be write only which does not need such an action and this is wasted resources.

queue.submit([&] (cl::sycl::handler& cgh) {         
         auto input_acc = input.get_access<sycl::access::mode::read>(cgh);
         auto output_acc = output.get_access<sycl::access::mode::discard_write>(cgh);
         auto maxRange = sycl::nd_range<2>(sycl::range<2>{height, width / 4}, sycl::range<2>(1, 128));
         cgh.parallel_for<class test>(maxRange, [=](sycl::nd_item<2> item){
            output_acc[item.get_global_id()] = 0;
         });
});

I have never written an issue like this before so if you need additional information just ask.

[Ruyk]

is this only happening in the CUDA backend?

[Ryefalk]

We don't have the tooling to see if it happens in other backends as of right now. And after some consideration the copy should exist when using write, but not when using discard_write. As write seems to be designed for only changing some of the data.

[Ruyk]

Yes, you are correct. I am not sure if the decision to trigger that copy or not is up to the CUDA backend, or is made by the SYCL RT

[Ryefalk]

What additional information or help would you need to figure it out, or can I try to do so in some way?

[Ruyk]

Would be good to have your full code. How are you creating the sycl buffers? With an initial pointer or without it?

The problem I see is that, if you pass an initial pointer to the buffer, even if the first time is accessed is discard, the data will be copied into the device.
This is because the DPC++ runtime creates CUDA (or OpenCL) memory with initial data by default, irrespective on how it is used later on. It may not have an impact on OpenCL, but on CUDA we trigger a host to device copy on construction that is what I am afraid you are seeing.
If you are not passing an initial pointer to the SYCL buffer (i.e., you construct is as sycl::buffer foo(someRange)), then the first copy is triggered by the SYCL runtime, the CUDA backend simply "responds" to that request from the runtime to perform the copy. But if that is the case then the OpenCL implementation should also see this additional copy.
If that is the case, @s-kanaev or others may be better suited to figure out the problem.

[Ryefalk]

namespace sycl = cl::sycl;

using r_acc_t =
   sycl::accessor<unsigned char, 2, sycl::access::mode::read, sycl::access::target::global_buffer>;
using dw_acc_t =
   sycl::accessor<unsigned char, 2, sycl::access::mode::discard_write, sycl::access::target::global_buffer>;
using local_accessor_t = 
   sycl::accessor<unsigned char, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::local>;

int main(int, char**) { 
   size_t width;
   size_t height;
   char const* input_file_path = "/home/andry_admin/cpu-gpu/assets/lena.ppm";
   char const* output_file_path = "/home/andry_admin/cpu-gpu/assets/result.ppm";
   int* input_img_ptr = (int*)Image::load_ppm_gray(input_file_path, width, height);
   int* tmp_img_ptr = (int*) malloc(width * height);
   int* output_img_ptr = (int*) malloc(width * height);

   sycl::range<2> img_range(height, width);
   sycl::default_selector device_selector;
   sycl::property_list prop_list{sycl::property::queue::enable_profiling()};
   sycl::queue queue(device_selector, prop_list);
   std::cout << "Running on "
               << queue.get_device().get_info<sycl::info::device::name>()
               << "\n";

   size_t profiling_iters = 2;
   std::vector<sycl::event> event_list(profiling_iters);
   std::vector<time_point_t> start_time_list(profiling_iters);
   auto current_time = wall_clock_t::now();
   {
      sycl::buffer<int, 2> input(input_img_ptr, sycl::range<2>{height, width / 4}); //row, col
      sycl::buffer<int, 2> output(output_img_ptr, sycl::range<2>{height, width / 4});
      sycl::buffer<int, 2> temp(tmp_img_ptr, sycl::range<2>{height, width / 4});

      start_time_list.at(0) = wall_clock_t::now();
      event_list.at(0) = queue.submit([&] (cl::sycl::handler& cgh) {   
         auto local_input_acc = local_accessor_t(512, cgh);
         auto local_output_acc = local_accessor_t(512, cgh);
         auto input_acc = input.get_access<sycl::access::mode::read>(cgh);
         auto output_acc = output.get_access<sycl::access::mode::discard_write>(cgh);
         auto maxRange = sycl::nd_range<2>(sycl::range<2>{height, width / 4}, sycl::range<2>(1, 128));
         cgh.parallel_for(maxRange, box_kernel::row_local(input_acc, output_acc, local_input_acc, local_output_acc, 10));
      });

      start_time_list.at(1) = wall_clock_t::now();
      event_list.at(1) = queue.submit([&] (cl::sycl::handler& cgh) {   
         auto local_input_acc = local_accessor_t(512 * 4, cgh);
         auto local_output_acc = local_accessor_t(512 * 4, cgh);
         //auto temp_acc = temp.get_access<sycl::access::mode::read>(cgh);
         auto input_acc = output.get_access<sycl::access::mode::read>(cgh);
         auto output_acc = input.get_access<sycl::access::mode::discard_write>(cgh);
         auto maxRange = sycl::nd_range<2>(sycl::range<2>{height, width / 4}, sycl::range<2>(512, 1));
         cgh.parallel_for(maxRange, box_kernel::col_local(input_acc, output_acc, local_input_acc, local_output_acc, 10));
      });

      queue.wait_and_throw();
   }
   auto end_time = wall_clock_t::now();
   profiler<double> profiler(event_list, start_time_list);
   std::chrono::duration<double, std::milli> real_time = end_time - current_time;
   std::cout << real_time.count() << std::endl;
   Image::write_ppm_gray(output_file_path, width, height, (unsigned char*)input_img_ptr);
  
   return 0;
}

[Ryefalk]

We give it an initial pointer. This is not the entire code, but the timers and parallel_for can be exchange for dummies. We can't test out the openCL version as we don't have a non nvidia card on hand sadly.

[Ruyk]

      sycl::buffer<int, 2> input(input_img_ptr, sycl::range<2>{height, width / 4}); //row, col
      sycl::buffer<int, 2> temp(tmp_img_ptr, sycl::range<2>{height, width / 4});
      sycl::buffer<int, 2> output(output_img_ptr, sycl::range<2>{height, width / 4});

As suspected, those buffers pass an initial pointer and I think there is no way on the SYCL runtime to avoid the initial copy as it is currently implemented, Intel folks probably have a better answer

[Ryefalk]

This copy only happens when we use an accessor, not when the buffer is created but that is maybe what you mean. Additionally two device to host copy operations are made after the kernels have run instead of the expected one.

[Ruyk]

Can you try:

 sycl::buffer<int, 2> input(input_img_ptr, sycl::range<2>{height, width / 4}); //row, col
      sycl::buffer<int, 2> temp(sycl::range<2>{height, width / 4});
      sycl::buffer<int, 2> output(sycl::range<2>{height, width / 4});
      output.set_final_ptr(output_img_ptr);

Also you dont need an allocation of temp on the host

[Ryefalk]

[Ryefalk]

We realised this later so we are no longer using it.

[Ryefalk]

This removes both the extra copy before and after the kernel run. We just used output as temp and used input as write back as we are using a separable filter. The "output.set_final_ptr(output_img_ptr);" gives a compiler error for us. But it was not needed in this case.

[Ruyk]

set_final_data is the right name, sorry.
Does that produces the expected number of copies now?