sycl/dpc++ 访问器与内核函数对象中的 global_ptr

Question

使用 Intel OneAPI beta6 的以下玩具代码。

#include <CL/sycl.hpp>
#include <iostream>

namespace sycl = cl::sycl;

const int SIZE=1;

class Increment_accessor {
  public:
    Increment_accessor(sycl::accessor<int, 1, sycl::access::mode::read_write, sycl::access::target::global_buffer> ptr_) : ptr {ptr_} {}
    void operator()(sycl::item<1> item) {
      ptr[item.get_linear_id()]++;
    }
  private:
    sycl::accessor<int, 1, sycl::access::mode::read_write, sycl::access::target::global_buffer> ptr;
};

class Increment_pointer {
  public:
    Increment_pointer(sycl::global_ptr<int> ptr_) : ptr {ptr_} {} 
    void operator()(sycl::item<1> item) {
      ptr[item.get_linear_id()]++;
    }
  private:
    sycl::global_ptr<int> ptr;
};

int 
main(int argc, char *argv[])
{
  sycl::device dev = sycl::default_selector().select_device();
  sycl::queue q(dev);
  int hbuffer[SIZE] = {};

  {
    sycl::buffer<int, 1> hbuf(hbuffer, sycl::range<1> {SIZE});
    q.submit([&](sycl::handler& cgh) {
        auto harray = hbuf.get_access<sycl::access::mode::read_write, sycl::access::target::global_buffer>(cgh);
        // !!! Uncomment _one_ of the following lines to compile !!!
        //Increment_accessor increment {harray};
        //Increment_pointer increment {harray};
        //Increment_pointer increment {harray.get_pointer()};
        cgh.parallel_for<class kernel1>(
            sycl::range<1> {SIZE}, 
            increment
        );
      }
      ); 
  }

  for (int i=0; i<SIZE; i++) std::cout << "hbuffer[" << i << "]= " << hbuffer[i] << std::endl;
}

问题：为什么 Increment_pointer 版本是 "wrong"？没有 compile/runtime 错误。您只是没有在最后获得递增的 hbuffer。（我玩过一些类似的版本，其中 operator() 中的 ptr 最终为 0x0）。

我还在学习在"SYCL"中思考，所以欢迎详细解释。

Answer 1

如果我没理解错的话，您是在问为什么您的代码在使用 Increment_accessor 时有效，但在使用 Increment_pointer 时却中断了。或者，更一般地说，是否可以构建一个接受指针而不是访问器作为参数的内核函数？

SYCL 规范对此不是很清楚，但 4.7.6.3 节给出了提示：

A SYCL accessor can be a device accessor in which case it provides access to data within a SYCL kernel function, or a host accessor in which case it provides immediate access on the host. If an accessor has the access target access::target::global_buffer, access::target::constant_buffer, access::target::local,access::target::image or access::target::image_array then it is considered a device accessor, and therefore can only be used within a SYCL kernel function

因此，设备访问器在主机上无效。现在，在基于指针的版本中，您正在调用 get_pointer() （或依赖从访问器到指针的隐式转换，这将是等效的）。但此时，你仍然在命令组范围内，不在内核内，即parallel_forincrement内核内的代码。命令组范围始终在 SYCL 中的主机上进行评估，因为在那里构造的访问器告诉 SYCL 运行时如何构建任务图以及 SYCL 任务图中的依赖节点有哪些。

因此我们可以将问题归结为调用 get_pointer() 是否已经符合 使用设备访问器 的条件。如果是这样，则意味着在 SYCL 内核之外使用了设备访问器，这违反了规范的引用部分并会使此代码非法。

我认为调用 get_pointer() 已经对应于“使用访问器”。例如，get_pointer() 要求访问器已经具有可以指向的有效设备内存分配。但是，在命令组范围内，这种分配甚至可能还不存在，因为高效的 SYCL 运行时可能只会在需要之前在后台懒惰地在设备内存中进行必要的分配。但是在命令组评估期间，命令组甚至还没有完全提交给 SYCL 运行时，因为它实际上仍处于正在提交的阶段。因此，SYCL 实现通常无法保证 get_pointer() 在此阶段已经正常工作。

因此，最好假设主机上的设备访问器仅仅是对访问数据的描述，并且只有作为一种在使用时实际允许数据访问的机制才具有更多意义里面的内核。

现在，对于是否有可能让内核接受指针而不是访问器作为参数的更普遍的问题：据我所知，在 SYCL 1.2.1 中不可能拥有访问器并以某种方式将它们转换为指针主机因为上述问题。但是，如果您根本不使用访问器，则可以使用 Intel unified shared memory (USM) 扩展。此扩展尚未出现在常规 SYCL 1.2.1 中，但在英特尔 oneAPI DPC++ 中可用。 USM 允许您显式创建使用指针管理的设备可访问分配。然后，您可以直接在内核中使用这些指针。

更新：USM 已合并到 2021 年 2 月发布的官方 SYCL 2020 标准中，现在可用于除 DPC++ 之外的多个 SYCL 实现。

sycl/dpc++ 访问器与内核函数对象中的 global_ptr

sycl/dpc++ accessor vs. global_ptr in kernel function object

sycl

intel-oneapi

dpc++