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 }