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:

  #define DATA_NUM 16

  struct Foo {
    ...
    float data[0];
  };

  __global__ void kernel(Foo *foo) {
    foo->data[DATA_NUM - 1] = 123.f;
  }

  void test() {
    Foo* foo;
    cudaMalloc(&foo, sizeof(Foo) + DATA_NUM * sizeof(float));
    kernel<<<1, 1>>>(foo);
    ...
  }

results in the following migrated SYCL code:

  #define DATA_NUM 16

  struct Foo {
    ...
    /*
    DPCT1102:0: Zero-length arrays are not permitted in SYCL device code.
    */
    float data[0];
  };

  void kernel(Foo *foo) {
    foo->data[DATA_NUM - 1] = 123.f;
  }

  void test() {
    sycl::device dev_ct1;
    sycl::queue q_ct1(dev_ct1,
                      sycl::property_list{sycl::property::queue::in_order()});
    Foo* foo;
    foo =
        (Foo *)sycl::malloc_device(sizeof(Foo) + DATA_NUM * sizeof(float), q_ct1);
    q_ct1.parallel_for(
        sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
        [=](sycl::nd_item<3> item_ct1) {
          kernel(foo);
        });
    ...
  }

which is rewritten to:

  #define DATA_NUM 16

  struct Foo {
    ...
    std::byte data[DATA_NUM];
  };

  void kernel(Foo *foo) {
    foo->data[DATA_NUM - 1] = 123.f;
  }

  void test() {
    sycl::device dev_ct1;
    sycl::queue q_ct1(dev_ct1,
                      sycl::property_list{sycl::property::queue::in_order()});
    Foo* foo;
    foo = (Foo *)sycl::malloc_device(sizeof(Foo), q_ct1);
    q_ct1.parallel_for(
        sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
        [=](sycl::nd_item<3> item_ct1) {
          kernel(foo);
        });
    ...
  }