CCL代码示例中的缓冲区以及oneapi工具包

Question

我浏览了CCL代码示例和oneapi工具包。在下面的DPC++(SYCL)代码中，最初sendbuf缓冲区是在CPU端创建的，并未初始化，并且在卸载到目标设备的部分修改了dev_acc_sbufid变量，该变量是内核范围中的一个变量。因此，这个变量(Dev_acc_sbuf)不会在程序中使用，它的值也不会复制回sendbuf.Then。在下一行中，sendbuf变量用于所有缩减。我不能理解改变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);“行中，dev_acc_sbuf是一个访问sendbuf而不是独立缓冲区的句柄。在dev_acc_sbuf句柄中所做的更改将反映到原始缓冲区，即发送缓冲区。这是SYCL的一个优点，因为在内核作用域中所做的更改会自动复制回原始变量

Answer 2

在大多数系统中，主机和设备不共享物理内存，CPU可能使用RAM，而GPU可能使用自己的全局内存。SYCL需要知道它将在主机和设备之间共享哪些数据。

为此，SYCL使用它的缓冲区，buffer类是元素类型和维度数量的泛型。当传递一个原始指针时，buffer(T* ptr，范围大小)构造函数将获得传递给它的内存的所有权。这意味着当缓冲区存在时，我们绝对不能自己使用该内存，这就是我们开始C++作用域的原因。在它们的作用域结束时，缓冲区将被销毁，内存将返回给用户。size参数是一个range对象，它必须具有与buffer相同的维数，并使用每个维度中的元素数目进行初始化。在这里，我们有一个维度和一个元素。

缓冲区不与特定的队列或上下文相关联，因此它们能够在多个设备之间透明地处理数据。

访问器用于从buffer对象访问对设备内存的请求控制。它们的模式将负责主机和设备之间的数据移动。因此，我们不必显式地将结果从一个设备复制到另一个主机。

以下是更多说明的示例：

#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;
}