在 SYCL 中使用障碍

Posted

技术标签:

【中文标题】在 SYCL 中使用障碍【英文标题】:Using Barriers in SYCL 【发布时间】:2019-10-17 16:24:40 【问题描述】:

我正在 SYCL 中进行矩阵乘法,并且有一个工作代码,其中我在 parallel_for 中只使用了 range,而不是在 parallel_for 中使用 nd_range >。现在我想在其中使用障碍,据我所知,障碍只能与 nd_range 一起使用,对吧?我附上了我的代码的一部分,请告诉我是否可以在没有nd_range 的情况下完成,或者我应该使用nd_range 进行哪些更改。谢谢

queue.submit([&](cl::sycl::handler &cgh) 
    auto A = A_sycl.get_access<cl::sycl::access::mode::read>(cgh);
    auto B = B_sycl.get_access<cl::sycl::access::mode::read>(cgh);
    auto C = C_sycl.get_access<cl::sycl::access::mode::write>(cgh);

    cgh.parallel_for<class test>(
        cl::sycl::range<2>(4, 4), [=](cl::sycl::id<2> id) 
        c_access[id] = A[id] * Y[id.get(1)];
    );

);

【问题讨论】:

【参考方案1】:

使用 nd_range 可让您明确指定本地范围。为了能够在内核中设置工作组屏障,您还需要使用 nd_item 而不是 id 来访问更多 id 位置和大小,例如作为全局和本地 id、组范围和本地范围,以及屏障同步原语。

然后,您可以在完成读取/写入到设备本地内存(使用仅限设备的本地访问器)后放置一个屏障

而使用 rangeid 无法获得任何功能。它只是为了简化命令组设置和全局内存内核的编写,您希望运行时为您决定工作组大小,并有一种简单的方法来索引您的工作项,而不是传统的 OpenCL 方法无论您的内核多么简单或复杂,您都必须始终明确定义 NDRange(SYCL 中的 nd_range)。

这是一个简单的示例,假设您要启动 2D 内核。

myQueue.submit([&](cl::sycl::handler& cgh) 
    auto A_ptr = A_buf.get_access<cl::sycl::access::mode::read>(cgh);
    auto B_ptr = B_buf.get_access<cl::sycl::access::mode::read_write>(cgh);
    auto C_ptr = C_buf.get_access<cl::sycl::access::mode::write>(cgh);
    // scratch/local memory for faster memory access to compute the results
    cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write,
                       cl::sycl::access::target::local>
        C_scratch(range<1>size, cgh);

    cgh.parallel_for<example_kernel>(
        cl::sycl::nd_range<2>(range<2>size >> 3, size >> 3,   // 8, 8
                              range<2>size >> 4, size >> 4),  // 4, 4
        [=](cl::sycl::nd_item<2> item) 
          // get the 2D x and y indices
          const auto id_x = item.get_global_id(0);
          const auto id_y = item.get_global_id(1);
          // map the 2D x and y indices to a single linear,
          // 1D (kernel space) index
          const auto width =
              item.get_group_range(0) * item.get_local_range(0);
          // map the 2D x and y indices to a single linear,
          // 1D (work-group) index
          const auto index = id_x * width + id_y;
          // compute A_ptr * B_ptr into C_scratch
          C_scratch[index] = A_ptr[index] * B_ptr[index];
          // wait for result to be written (sync local memory read_write)
          item.barrier(cl::sycl::access::fence_space::local_space);
          // output result computed in local memory
          C_ptr[index] = C_scratch[index];
        );
  );

我使用主机数据和 SYCL 缓冲区的 1D 表示,它解释了从 2D 索引到单个线性 1D 索引的映射。

我希望这个解释有助于在您的案例中应用这些概念。

【讨论】:

感谢您的详细解释,这非常有帮助,我会尝试实施。再次感谢

以上是关于在 SYCL 中使用障碍的主要内容,如果未能解决你的问题,请参考以下文章

在 MacOS 中使用 SYCL 1.2

在 MacOS 中使用 SYCL 1.2

在 SYCL 中使用一个缓冲区还是多个缓冲区更有效?

在 SYCL 中使用一个缓冲区还是多个缓冲区更有效?

如何从 SYCL 内核中提取控制流?

如何从 SYCL 内核中提取控制流?