DPCT1102#

Message#

Zero-length arrays are not permitted in SYCL device code.

Detailed Help#

Zero-length arrays are not supported in the C++ standard or SYCL* specification. It may be supported by some host compiler implementations, but it is not supported by the IntelĀ® oneAPI DPC++/C++ Compiler for device code.

Suggestions to Fix#

Use an array with length greater than zero.

For example, this original CUDA* code:

 1  struct Foo {
 2    ...
 3    float data[0];
 4  };
 5
 6  __global__ void kernel(Foo *foo) {
 7    Foo->data[DATA_NUM - 1] = 123.f;
 8  }
 9
10  int main() {
11    Foo* foo;
12    cudaMalloc(&foo, sizeof(Foo) + DATA_NUM * sizeof(float));
13    kernel<<<1, 1>>>(foo);
14    ...
15  }

results in the following migrated SYCL code:

 1  struct Foo {
 2    ...
 3    std::byte data[0];
 4  };
 5
 6  void kernel(Foo *foo) {
 7    foo->data[DATA_NUM - 1] = 123.f;
 8  }
 9
10  int main() {
11    dpct::device_ext &dev_ct1 = dpct::get_current_device();
12    sycl::queue &q_ct1 = dev_ct1.default_queue();
13    Foo* foo;
14    foo = (Foo *)sycl::malloc_device(sizeof(Foo) + DATA_NUM * sizeof(float), q_ct1);
15    q_ct1.parallel_for(
16        sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
17        [=](sycl::nd_item<3> item_ct1) {
18          kernel(foo);
19        });
20    ...
21  }

which is rewritten to:

 1  struct Foo {
 2    ...
 3    std::byte data[DATA_NUM];
 4  };
 5
 6  void kernel(Foo *foo) {
 7    foo->data[DATA_NUM - 1] = 123.f;
 8  }
 9
10  int main() {
11    dpct::device_ext &dev_ct1 = dpct::get_current_device();
12    sycl::queue &q_ct1 = dev_ct1.default_queue();
13    Foo* foo;
14    foo = (Foo *)sycl::malloc_device(sizeof(Foo), q_ct1);
15    q_ct1.parallel_for(
16        sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
17        [=](sycl::nd_item<3> item_ct1) {
18          kernel(foo);
19        });
20    ...
21  }