2

以下玩具代码使用 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”进行思考，因此欢迎详细解释。

4

1 回答 1

5

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

SYCL 规范对此并不十分清楚，但第 4.7.6.3 节给出了提示：

SYCL 访问器可以是设备访问器，在这种情况下，它提供对 SYCL 内核函数内数据的访问，或者是主机访问器，在这种情况下，它提供对主机的即时访问。如果访问器具有访问目标 access::target::global_buffer, access::target::constant_buffer, access::target::local,access::target::image 或 access::target::image_array 则认为它设备访问器，因此只能在 SYCL 内核函数中使用

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

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

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

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

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

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

于 2020-05-12T13:46:31.950 回答