DPCT1054

Intel® DPC++ Compatibility Tool Developer Guide and Reference

Download PDF

ID 768918

Date 10/31/2024

Version

Public

Visible to Intel only — GUID: GUID-5BE6D83D-98D2-4AE2-B55C-1594A2014A00

View Details

DPCT1054

Message

The type of variable <variable name> is declared in device function with the name <type>. Adjust the code to make the <type> declaration visible at the accessor declaration point.

Detailed Help

This warning will be emitted when the type of a __shared__ variable is declared in a device function.

In the migrated code:

in the device function, the tool adds:

# uint8_t* input parameter # name to the type, if unnamed, for example type_ct1 # a type cast from uint8_t* to the original type, after the original type declaration.
in the call side, when defining local accessor, the first template argument will be set to uint8_t[sizeof(<original type, for example "type_ct1">)].

For example, this original CUDA* code:

// header file
template <typename T> __global__ void k() {
  __shared__ union {
    T t;
    ...
  } a;
  ...
}

// source file
void foo() { k<int><<<1, 1>>>(); }

results in the following migrated SYCL* code:

// header file
template <typename T> void k(uint8_t *a_ct1) {
  union type_ct1 {
    T t;
    ...
  };
  type_ct1 &a = *(type_ct1 *)a_ct1;
  ...
}

// source file
void foo() {
  dpct::get_default_queue().submit([&](sycl::handler &cgh) {
    /*
    DPCT1054:0: The type of variable a is declared in device function with the
    name type_ct1. Adjust the code to make the type_ct1 declaration visible at
    the accessor declaration point.
    */
    sycl::local_accessor<uint8_t[sizeof(type_ct1)], 0> a_ct1_acc_ct1(cgh);

    cgh.parallel_for(
        sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
        [=](sycl::nd_item<3> item_ct1) {
          k<int>(a_ct1_acc_ct1.get_pointer());
        });
  });
}

which is rewritten to:

// header file
template <typename T>
union type_ct1 {
  T t;
  ...
};
template <typename T> void k(uint8_t *a_ct1) {
  type_ct1<T> &a = *(type_ct1<T> *)a_ct1;
  ...
}

// source file
void foo() {
  dpct::get_default_queue().submit([&](sycl::handler &cgh) {
    sycl::local_accessor<uint8_t[sizeof(type_ct1<T>)], 0> a_ct1_acc_ct1(cgh);

    cgh.parallel_for(
        sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
        [=](sycl::nd_item<3> item_ct1) {
          k<int>(a_ct1_acc_ct1.get_pointer());
        });
  });
}

Suggestions to Fix

Move the type declaration so that it will be visible at the accessor declaration point or replace the sizeof(<original type> with size in bytes needed for the original type.

Select Your Language

Using Intel.com Search

Quick Links

Recent Searches

Advanced Search

Only search in

Intel® DPC++ Compatibility Tool Developer Guide and Reference

DPCT1054

Message

Detailed Help

Suggestions to Fix