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 版本“错误”？没有编译/运行时错误。您只是最后没有得到递增的 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_for increment 内核内的代码。命令组范围始终在 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 中，但可在 Intel oneAPI DPC++ 中使用。 USM 允许您显式创建使用指针管理的设备可访问分配。然后您可以直接在内核中使用这些指针。

更新:USM 已合并到 2021 年 2 月发布的官方 SYCL 2020 标准中，除了 DPC++ 之外，现在还可以在多种 SYCL 实现中使用。