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

Question

我正在浏览 CCL 代码示例以及 oneapi 工具包。在下面的 DPC++(SYCL) 代码中，最初 sendbuf 在 cpu 端创建了一个缓冲区并且未初始化，并且在卸载到目标设备的部分发生了 dev_acc_sbuf[id] 变量，它是修改了内核作用域中的变量。此变量 (dev_acc_sbuf) 因此未在程序中使用，其值也未在下一行中复制回 sendbuf.Then 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);”行中，dev_acc_sbuf 是访问 sendbuf 的句柄，而不是单独的缓冲区。在 dev_acc_sbuf 句柄中所做的更改会反映到原始缓冲区，即 sendbuffer 。这是 SYCL 的一个优势，因为在内核范围内所做的更改会自动复制回原始变量

Answer 2

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

为此，SYCL 使用其缓冲区，缓冲区 class 在元素类型和维数上是通用的。当传递一个原始指针时， buffer(T* ptr, range size) 构造函数获得它所传递的内存的所有权。这意味着我们绝对不能在缓冲区存在时自己使用该内存，这就是我们开始 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;
}

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

buffers in CCL code samples along with the oneapi toolkit

ccl

sycl

intel-ccl

intel-oneapi

dpc++