Visible to Intel only — GUID: GUID-5F4AD815-4367-475E-811C-C2B8CA18E850
Visible to Intel only — GUID: GUID-5F4AD815-4367-475E-811C-C2B8CA18E850
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
Intel® DPC++ Compatibility Tool 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:
__device__ void bar1(double *out1, double *d) {
double i = 30;
double &d2 = *d;
sincos(i, out1, &d2);
}
__global__ void kernel(double *out1, double *out2) {
double d;
bar1(out1, &d);
*out2 = d;
}
results in the following migrated SYCL code:
void bar1(double *out1, double *d) {
double i = 30;
double &d2 = *d;
/*
DPCT1081:0: The generated code assumes that "&d2" 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".
*/
*out1 = sycl::sincos(
i, sycl::address_space_cast<sycl::access::address_space::global_space,
sycl::access::decorated::yes, double>(&d2));
}
void kernel(double *out1, double *out2) {
double d;
bar1(out1, &d);
*out2 = d;
}
which is rewritten to:
void bar1(double *out1, double *d) {
double i = 30;
double &d2 = *d;
*out1 = sycl::sincos(
i, sycl::address_space_cast<sycl::access::address_space::private_space,
sycl::access::decorated::yes, double>(&d2));
}
void kernel(double *out1, double *out2) {
double d;
bar1(out1, &d);
*out2 = d;
}