将访问器强制转换为内核代码中的 C++ 指针(尤其是 (int (*)[Nelem])
casting accessors to C++ pointers in kernel code (esp. (int (*)[Nelem])
环境:Ubuntu18.04,OneAPI beta 6
下面是完整代码,但这里是有问题的错误:
#dpcpp -O2 -g -o so2 so2.cpp -lOpenCL -lsycl
so2.cpp:64:38: error: cannot cast from type 'global_ptr<int>' (aka 'multi_ptr<int, access::address_space::global_space>') to pointer type 'int (*)[nelem]'
int (*xptr)[nelem] = (int (*)[nelem])xaccessor.get_pointer();
^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
1 error generated.
一些解释,以防您想知道为什么....
在开发数据并行代码时,我经常开发英特尔过去称之为"elemental functions"的东西。这些被编写为对应用程序的单个元素(SYCL 称之为工作项)进行操作。我一直发现使用基本的 SW 开发环境更容易做到这一点,易于测试,并且更普遍地可重用(标量、SIMD、CUDA 等)。
在对单个元素进行测试后,通过扩展调用代码而无需 rewrite/retest 函数:
即可轻松实现数据并行化
int x[NELEM]
fn1(x, NELEM)
变成
int x[NPROC][NELEM]
for (int p=0; p<NPROC; p++) fn1(x[p], NELEM);
在 SYCL 内核中,fn1(x[item.get_linear_id()], NELEM);这就是我所需要的,而不必重写函数来理解 ids and/or 访问器。
上述代码的 SYCL 问题是在内核 C++ 中我似乎无法将访问器指针重铸为 2D 指针。这在应用程序 C++ 中是允许的(参见上面的代码)。
也许这是正确代码的糟糕方法,但它可以很容易地 develop/test 适用于标量和数据并行代码的代码,并使库具有一定的可移植性。它还提供了一种绕过 buffers/accessors.
上的 SYCL 3 维度限制的方法
无论如何,我很好奇真正的 SYCL 程序员会怎么想。
玩具示例的完整代码:
#include <CL/sycl.hpp>
#include <cstdio>
namespace sycl = cl::sycl;
const int Nproc=3;
const int Nelem=4;
/** elemental function **/
void
fn1(int *h, int n)
{
for (int i=0; i<n; i++) h[i] = 10*h[i]+2*i;
}
int
main(int argc, char *argv[])
{
/** Make some memory **/
int x1d[Nproc * Nelem];
for (int j=0; j<Nproc; j++) {
for (int i=0; i<Nelem; i++) x1d[j*Nelem+i] = 10*j+i;
}
printf("1D\n");
for (int i=0; i<Nelem; i++) {
printf("%d : ", i);
for (int j=0; j<Nproc; j++) printf("%d ", x1d[j*Nelem+i]);
printf("\n");
}
/** Reshape it into 2D **/
int (*x2d)[Nelem] = (int (*)[Nelem])x1d;
for (int j=0; j<Nproc; j++) fn1(x2d[j], Nelem);
printf("2D\n");
for (int i=0; i<Nelem; i++) {
printf("%d : ", i);
for (int j=0; j<Nproc; j++) printf("%d ", x2d[j][i]);
printf("\n");
}
/** SYCL setup **/
sycl::device dev = sycl::default_selector().select_device();
std::cout << "Device: "
<< "name: " << dev.get_info<sycl::info::device::name>() << std::endl
<< "vendor: " << dev.get_info<sycl::info::device::vendor>() << std::endl;
sycl::queue q(dev);
{
sycl::buffer<int, 1> xbuffer(x1d, sycl::range<1> {Nproc*Nelem});
q.submit([&](sycl::handler& cgh) {
int nelem = Nelem;
auto xaccessor = xbuffer.get_access<sycl::access::mode::read_write, sycl::access::target::global_buffer>(cgh);
cgh.parallel_for<class k0>(
sycl::range<1> {Nproc},
[=] (sycl::item<1> item) {
int idx = item.get_linear_id();
#if 0
int *xptr = (int *)xaccessor.get_pointer(); // doing this does work so we _can_ get a real pointer
fn1(xptr + nelem*idx, nelem);
#else
int (*xptr)[nelem] = (int (*)[nelem])xaccessor.get_pointer();
//int *ptr = (int *)xaccessor.get_pointer(); // splitting it into two doesn't work either
//int (*xptr)[nelem] = (int (*)[nelem])ptr;
fn1(xptr[idx], nelem);
#endif
}
);
}
);
}
printf("2D SYCL\n");
for (int i=0; i<Nelem; i++) {
printf("%d : ", i);
for (int j=0; j<Nproc; j++) printf("%d ", x1d[j*Nelem+i]);
printf("\n");
}
}
编辑 1:
根据 illuhad 的评论,我试图充实一些替代方案。
首先,这两行注释似乎应该按照他的建议进行:
int *ptr = (int *)xaccessor.get_pointer();
int (*xptr)[nelem] = (int (*)[nelem])ptr;
但实际上它会产生此错误:
error: cannot initialize a variable of type 'int (*)[nelem]' with an rvalue of type 'int (*)[nelem]'
int (*xptr)[nelem] = (int (*)[nelem])ptr;
^ ~~~~~~~~~~~~~~~~~~~
在 get_pointer 的末尾添加一个 "get()" 得到相同的结果。
奇怪的是,解决错误的 "initialize" 部分:
int *ptr = (int *)xaccessor.get_pointer().get();
int (*xptr)[nelem];
xptr = (int (*)[nelem])ptr;
产生有趣的错误:
error: incompatible pointer types assigning to 'int (*)[nelem]' from 'int (*)[nelem]'
xptr = (int (*)[nelem])ptr;
^~~~~~~~~~~~~~~~~~~
所以if/when有人有时间,我还是很好奇...
简短回答:不是 SYCL 问题 ;)
根据您的编辑 1,很明显如果行
int *ptr = (int *)xaccessor.get_pointer();
int (*xptr)[nelem] = (int (*)[nelem])ptr;
在第二行导致转换错误,它不可能真的是 DPC++/SYCL 问题,因为这里只涉及 int 指针的变体,与 SYCL 无关。
事实上,问题在于 nelem
不是编译时常量。所以,下面的非SYCL测试程序
int main(){
int nelem = 10;
int* ptr = nullptr;
int (*xptr)[nelem] = (int (*)[nelem])ptr;
}
使用 -pedantic
使用常规 clang 或 gcc 编译时重现您的问题。然而,默认情况下,gcc 支持可变长度数组作为 C++ 中的扩展,因此即使代码不是有效的 C++,它也会编译。
根据 C++ 的要求,通过将 nelem
转换为编译时常量,您的问题已得到解决。可变长度数组是较新版本的 C 的一部分,但不是 C++ 的一部分。
环境:Ubuntu18.04,OneAPI beta 6
下面是完整代码,但这里是有问题的错误:
#dpcpp -O2 -g -o so2 so2.cpp -lOpenCL -lsycl
so2.cpp:64:38: error: cannot cast from type 'global_ptr<int>' (aka 'multi_ptr<int, access::address_space::global_space>') to pointer type 'int (*)[nelem]'
int (*xptr)[nelem] = (int (*)[nelem])xaccessor.get_pointer();
^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
1 error generated.
一些解释,以防您想知道为什么....
在开发数据并行代码时,我经常开发英特尔过去称之为"elemental functions"的东西。这些被编写为对应用程序的单个元素(SYCL 称之为工作项)进行操作。我一直发现使用基本的 SW 开发环境更容易做到这一点,易于测试,并且更普遍地可重用(标量、SIMD、CUDA 等)。
在对单个元素进行测试后,通过扩展调用代码而无需 rewrite/retest 函数:
即可轻松实现数据并行化 int x[NELEM]
fn1(x, NELEM)
变成
int x[NPROC][NELEM]
for (int p=0; p<NPROC; p++) fn1(x[p], NELEM);
在 SYCL 内核中,fn1(x[item.get_linear_id()], NELEM);这就是我所需要的,而不必重写函数来理解 ids and/or 访问器。
上述代码的 SYCL 问题是在内核 C++ 中我似乎无法将访问器指针重铸为 2D 指针。这在应用程序 C++ 中是允许的(参见上面的代码)。
也许这是正确代码的糟糕方法,但它可以很容易地 develop/test 适用于标量和数据并行代码的代码,并使库具有一定的可移植性。它还提供了一种绕过 buffers/accessors.
上的 SYCL 3 维度限制的方法无论如何,我很好奇真正的 SYCL 程序员会怎么想。
玩具示例的完整代码:
#include <CL/sycl.hpp>
#include <cstdio>
namespace sycl = cl::sycl;
const int Nproc=3;
const int Nelem=4;
/** elemental function **/
void
fn1(int *h, int n)
{
for (int i=0; i<n; i++) h[i] = 10*h[i]+2*i;
}
int
main(int argc, char *argv[])
{
/** Make some memory **/
int x1d[Nproc * Nelem];
for (int j=0; j<Nproc; j++) {
for (int i=0; i<Nelem; i++) x1d[j*Nelem+i] = 10*j+i;
}
printf("1D\n");
for (int i=0; i<Nelem; i++) {
printf("%d : ", i);
for (int j=0; j<Nproc; j++) printf("%d ", x1d[j*Nelem+i]);
printf("\n");
}
/** Reshape it into 2D **/
int (*x2d)[Nelem] = (int (*)[Nelem])x1d;
for (int j=0; j<Nproc; j++) fn1(x2d[j], Nelem);
printf("2D\n");
for (int i=0; i<Nelem; i++) {
printf("%d : ", i);
for (int j=0; j<Nproc; j++) printf("%d ", x2d[j][i]);
printf("\n");
}
/** SYCL setup **/
sycl::device dev = sycl::default_selector().select_device();
std::cout << "Device: "
<< "name: " << dev.get_info<sycl::info::device::name>() << std::endl
<< "vendor: " << dev.get_info<sycl::info::device::vendor>() << std::endl;
sycl::queue q(dev);
{
sycl::buffer<int, 1> xbuffer(x1d, sycl::range<1> {Nproc*Nelem});
q.submit([&](sycl::handler& cgh) {
int nelem = Nelem;
auto xaccessor = xbuffer.get_access<sycl::access::mode::read_write, sycl::access::target::global_buffer>(cgh);
cgh.parallel_for<class k0>(
sycl::range<1> {Nproc},
[=] (sycl::item<1> item) {
int idx = item.get_linear_id();
#if 0
int *xptr = (int *)xaccessor.get_pointer(); // doing this does work so we _can_ get a real pointer
fn1(xptr + nelem*idx, nelem);
#else
int (*xptr)[nelem] = (int (*)[nelem])xaccessor.get_pointer();
//int *ptr = (int *)xaccessor.get_pointer(); // splitting it into two doesn't work either
//int (*xptr)[nelem] = (int (*)[nelem])ptr;
fn1(xptr[idx], nelem);
#endif
}
);
}
);
}
printf("2D SYCL\n");
for (int i=0; i<Nelem; i++) {
printf("%d : ", i);
for (int j=0; j<Nproc; j++) printf("%d ", x1d[j*Nelem+i]);
printf("\n");
}
}
编辑 1:
根据 illuhad 的评论,我试图充实一些替代方案。
首先,这两行注释似乎应该按照他的建议进行:
int *ptr = (int *)xaccessor.get_pointer();
int (*xptr)[nelem] = (int (*)[nelem])ptr;
但实际上它会产生此错误:
error: cannot initialize a variable of type 'int (*)[nelem]' with an rvalue of type 'int (*)[nelem]'
int (*xptr)[nelem] = (int (*)[nelem])ptr;
^ ~~~~~~~~~~~~~~~~~~~
在 get_pointer 的末尾添加一个 "get()" 得到相同的结果。
奇怪的是,解决错误的 "initialize" 部分:
int *ptr = (int *)xaccessor.get_pointer().get();
int (*xptr)[nelem];
xptr = (int (*)[nelem])ptr;
产生有趣的错误:
error: incompatible pointer types assigning to 'int (*)[nelem]' from 'int (*)[nelem]'
xptr = (int (*)[nelem])ptr;
^~~~~~~~~~~~~~~~~~~
所以if/when有人有时间,我还是很好奇...
简短回答:不是 SYCL 问题 ;)
根据您的编辑 1,很明显如果行
int *ptr = (int *)xaccessor.get_pointer();
int (*xptr)[nelem] = (int (*)[nelem])ptr;
在第二行导致转换错误,它不可能真的是 DPC++/SYCL 问题,因为这里只涉及 int 指针的变体,与 SYCL 无关。
事实上,问题在于 nelem
不是编译时常量。所以,下面的非SYCL测试程序
int main(){
int nelem = 10;
int* ptr = nullptr;
int (*xptr)[nelem] = (int (*)[nelem])ptr;
}
使用 -pedantic
使用常规 clang 或 gcc 编译时重现您的问题。然而,默认情况下,gcc 支持可变长度数组作为 C++ 中的扩展,因此即使代码不是有效的 C++,它也会编译。
根据 C++ 的要求,通过将 nelem
转换为编译时常量,您的问题已得到解决。可变长度数组是较新版本的 C 的一部分,但不是 C++ 的一部分。