DPCT1081#

Message#

The generated code assumes that <pointer variable> points to the global memory address space. If it points to a local or private memory address space, replace address_space::global with address_space::local or address_space::private.

Detailed Help#

SYCLomatic tries to deduce the address space of the memory that <pointer variable> points to. If the tool cannot deduce the address space (for example, the pointer variable is used by multiple functions), the tool will use address_space::global and the warning will be emitted.

Suggestions to Fix#

Check the address space of the memory that <pointer variable> points to and replace address_space::global with the correct address space.

For example, this original CUDA* code:

 1__device__ void bar1(double *out1, double *d) {
 2  double i = 30;
 3  double &d2 = *d;
 4  sincos(i, out1, &d2);
 5}
 6
 7__global__ void kernel(double *out1, double *out2) {
 8  double d;
 9  bar1(out1, &d);
10  *out2 = d;
11}

results in the following migrated SYCL code:

 1void bar1(double *out1, double *d) {
 2  double i = 30;
 3  double &d2 = *d;
 4  /*
 5  DPCT1081:0: The generated code assumes that "&d2" points to the global memory
 6  address space. If it points to a local or private memory address space,
 7  replace "address_space::global" with "address_space::local" or
 8  "address_space::private".
 9  */
10  *out1 = sycl::sincos(
11      i, sycl::address_space_cast<sycl::access::address_space::global_space,
12                                  sycl::access::decorated::yes, double>(&d2));
13}
14
15void kernel(double *out1, double *out2) {
16  double d;
17  bar1(out1, &d);
18  *out2 = d;
19}

which is rewritten to:

 1void bar1(double *out1, double *d) {
 2  double i = 30;
 3  double &d2 = *d;
 4  *out1 = sycl::sincos(
 5      i, sycl::address_space_cast<sycl::access::address_space::private_space,
 6                                  sycl::access::decorated::yes, double>(&d2));
 7}
 8
 9void kernel(double *out1, double *out2) {
10  double d;
11  bar1(out1, &d);
12  *out2 = d;
13}