在 SYCL 中使用障碍
Using Barriers in SYCL
本文关键字:SYCL 更新时间:2023-10-16
我正在 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)];
});
});
使用nd_range可以显式指定本地范围。为了能够在内核中放置工作组屏障,您还需要使用nd_item而不是 id 来访问更多 id 位置和大小,例如全局和本地id、组范围和本地范围,以及屏障同步原语。
然后,您可以在完成对设备本地内存的读取/写入后放置屏障(使用仅限设备的本地访问器(。
而使用范围和id无法为您提供任何该功能。它只是为了简化命令组的设置和全局内存内核的编写,您希望运行时为您决定工作组大小,并有一种简单的方法来索引您的工作项,而不是传统的 OpenCL 方法,在传统的 OpenCL 方法中,您必须始终显式定义 NDRange(nd_range 在 SYCL 中(,无论您的内核多么简单或复杂。
下面是一个简单的示例,假设您要启动 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 缓冲区的一维表示,它解释了从二维索引到单个线性一维索引的映射。
我希望这个解释有助于将这些概念应用于您的案例。