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

Posted 2023-02-22

【中文标题】sycl/dpc++ 访问器与内核函数对象中的 global_ptr

【英文标题】：sycl/dpc++ accessor vs. global_ptr in kernel function object

【发布时间】：2020-05-11 23:45:04

【问题描述】：

以下玩具代码使用英特尔 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”思考，所以欢迎详细解释。

【问题讨论】：

您提供的这段代码无法编译，增量未定义，恐怕我不确定您想用它做什么。您是否看过我们在 GitHub 存储库中提供的一些 SYCL 示例？ github.com/codeplaysoftware/computecpp-sdk/tree/master/samples

对不起，我不清楚。您需要取消注释源代码中的某一行以获得正确的函数对象。我已经编辑了代码以使其更清晰。

【参考方案1】：

如果我理解正确，您是在问为什么您的代码在使用 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_forincrement 内核内的代码。命令组范围始终在 SYCL 中的主机上进行评估，因为在那里构造的访问器告诉 SYCL 运行时如何构建任务图以及 SYCL 任务图中的依赖节点具有哪些依赖项。

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

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

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

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

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

【讨论】：

现在我明白了。感谢您的精彩解释。现在我知道我在错误的范围内使用了访问器，我有点想知道编译器和/或运行时是否应该将其标记为错误。无论如何，请继续关注。我很难理解 SYCL 中隐藏/隐含（以及描述不足）的行为，并希望将来我会有更多的 *** 问题。

太棒了！在基于编译器的 SYCL 实施（例如英特尔 DPC++）中，诊断此类情况并打印错误或警告应该是可能的，而且听起来是个好主意。我建议您在 github.com/intel/llvm 上打开一个问题，并向英特尔 DPC++ 开发人员提出这个建议。请注意，有一些 SYCL 实现依赖于 SYCL 作为纯 C++ 语法，因此将 SYCL 实现为纯 C++ 库。对于这些实现，打印此类诊断是不可能的，因为没有 SYCL 特定的编译器支持。

有趣的用例...我在gitlab.khronos.org/sycl/Specification/issues/346中为规范创建了一个内部问题