буферы в образцах кода CCL вместе с инструментарием oneapi

Question

Я просматривал образцы кода CCL вместе с инструментарием oneapi. В приведенном ниже коде DPC++(SYCL) изначально sendbuf буфер создается на стороне процессора и не инициализируется, а в той части, где происходит разгрузка на целевое устройство, изменяется переменная dev_acc_sbuf [id], которая является переменной в области ядра.. Следовательно, эта переменная (dev_acc_sbuf) не используется в программе и ее значение не копируется обратно в sendbuf. Затем в следующей строке переменная sendbuf используется для allreduce. Я не могу понять, как изменение dev_acc_sbuf влияет на sendbuf.

          cl::sycl::queue q;
cl::sycl::buffer<int, 1> sendbuf(COUNT);
          /* open sendbuf and modify it on the target device side */
q.submit([&](cl::sycl::handler& cgh) {
   auto dev_acc_sbuf = sendbuf.get_access<mode::write>(cgh);
   cgh.parallel_for<class allreduce_test_sbuf_modify>(range<1>{COUNT}, [=](item<1> id) {
       dev_acc_sbuf[id] += 1;
   });
});
/* invoke ccl_allreduce on the CPU side */
ccl_allreduce(&sendbuf,
              &recvbuf,
              COUNT,
              ccl_dtype_int,
              ccl_reduction_sum,
              NULL,
              NULL,
              stream,
              &request);

Answer 1

В соответствии "auto dev_acc_sbuf = sendbuf.get_access<mode::write>(cgh);"the dev_acc_sbuf это дескриптор, который обращается к sendbufа не отдельный буфер. Изменения, сделанные в дескрипторе dev_acc_sbuf, отражаются в исходном буфере, то есть в буфере отправки. Это преимущество SYCL, поскольку изменения, сделанные в области ядра, автоматически копируются обратно в исходную переменную.

Answer 2

В большинстве систем хост и устройство не совместно используют физическую память, ЦП может использовать ОЗУ, а графический процессор может использовать свою собственную глобальную память. SYCL необходимо знать, какие данные он будет передавать между хостом и устройствами.

Для этой цели SYCL использует свои буферы, класс буфера является универсальным по типу элемента и количеству измерений. Когда передается необработанный указатель, конструктор буфера (T* ptr, размер диапазона) становится владельцем переданной памяти. Это означает, что мы абсолютно не можем использовать эту память сами, пока существует буфер, поэтому мы начинаем область видимости C++. По окончании их действия буферы будут уничтожены, а память будет возвращена пользователю. Аргумент размера - это объект диапазона, который должен иметь такое же количество измерений, что и буфер, и инициализируется количеством элементов в каждом измерении. Здесь у нас есть одно измерение с одним элементом.

Буферы не связаны с конкретной очередью или контекстом, поэтому они могут прозрачно обрабатывать данные между несколькими устройствами.

Аксессоры используются для управления запросами доступа к памяти устройства из буферных объектов. Их режимы позаботятся о перемещении данных между хостом и устройством. Таким образом, нам не нужно явно копировать результат с устройства на хост.

Ниже приведен пример для пояснения:

#include <bits/stdc++.h>
#include <CL/sycl.hpp>

using namespace std;
class vector_addition;

int main(int, char**) {
   //creating host memory
   int *a=(int *)malloc(10*sizeof(int));
   int *b=(int *)malloc(10*sizeof(int));
   int *c=(int *)malloc(10*sizeof(int));

   for(int i=0;i<10;i++){
       a[i]=i;
       b[i]=10-i;
   }

   cl::sycl::default_selector device_selector;

   cl::sycl::queue queue(device_selector);
   std::cout << "Running on "<< queue.get_device().get_info<cl::sycl::info::device::name>()<< "\n";

 {
    //creating buffer from pointer of host memory
    cl::sycl::buffer<int, 1> a_sycl{a, cl::sycl::range<1>{10} };
    cl::sycl::buffer<int, 1> b_sycl{b, cl::sycl::range<1>{10} };
    cl::sycl::buffer<int, 1> c_sycl{c, cl::sycl::range<1>{10} };

    queue.submit([&] (cl::sycl::handler& cgh) {
       //creating accessor of buffer with proper mode
       auto a_acc = a_sycl.get_access<cl::sycl::access::mode::read>(cgh);
       auto b_acc = b_sycl.get_access<cl::sycl::access::mode::read>(cgh);
       auto c_acc = c_sycl.get_access<cl::sycl::access::mode::write>(cgh);//responsible for copying back to host memory 

       //kernel for execution
       cgh.parallel_for<class vector_addition>(cl::sycl::range<1>{ 10 }, [=](cl::sycl::id<1> idx) {
       c_acc[idx] = a_acc[idx] + b_acc[idx];
       });

    });
 }

 for(int i=0;i<10;i++){
     cout<<c[i]<<" ";    
 }
 cout<<"\n";
 return 0;
}