0

环境：Ubuntu 18.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.

有点解释，以防你想知道为什么......

在开发数据并行代码时，我经常开发英特尔曾经称之为“基本函数”的东西。这些被编写用于对应用程序的单个元素（SYCL 将其称为工作项）进行操作。我一直发现使用基本的软件开发环境更容易做到这一点，易于测试，并且更普遍地可重用（标量、SIMD、CUDA 等）。

在对单个元素进行测试之后，通过扩展调用代码而无需重写/重新测试函数，移动到数据并行非常容易：

    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); 将是我所需要的，而不必重写函数来理解 id 和/或访问器。

上述代码的 SYCL 问题是，在内核 C++ 中，我似乎无法将访问器指针重铸为 2D 指针。这在应用程序 C++ 中是允许的（参见上面的代码）。

也许这是一种纠正代码的坏方法，但它使开发/测试适用于标量和数据并行代码的代码变得容易，并使库在某种程度上具有可移植性。它还提供了一种绕过缓冲区/访问器的 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()”会产生相同的结果。

奇怪的是，解决错误的“初始化”部分：

    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;
                       ^~~~~~~~~~~~~~~~~~~

所以如果/当有人有时间时，我仍然很好奇......

4

1 回答 1

0

简短回答：不是 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;
}

使用常规 clang 或 gcc 编译时重现您的问题-pedantic。然而，默认情况下，gcc 支持可变长度数组作为 C++ 中的扩展，因此即使它不是有效的 C++，代码也可以编译。

根据 C++ 的要求，您的问题可以通过nelem转换为编译时常量来解决。可变长度数组是较新版本的 C 的一部分，但不是 C++ 的一部分。

于 2020-05-19T15:12:04.597 回答