3

我正在浏览 CCL 代码示例以及 oneapi 工具包。在下面的 DPC++(SYCL) 代码中，最初 sendbuf 在 cpu 端创建了一个缓冲区，但未初始化，在卸载到目标设备的部分中，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);
4

2 回答 2

3

在“ auto dev_acc_sbuf = sendbuf.get_access<mode::write>(cgh);”行中，dev_acc_sbuf是访问的句柄，sendbuf而不是单独的缓冲区。在 dev_acc_sbuf 句柄中所做的更改会反映到原始缓冲区，即 sendbuffer 。这是 SYCL 的一个优势，因为在内核范围内所做的更改会自动复制回原始变量

于 2019-11-20T13:09:59.967 回答
2

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

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

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

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

以下是更多说明的示例：

#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;
}
于 2019-11-26T07:41:23.977 回答