Visible to Intel only — GUID: GUID-CF1C51DA-9501-4575-981C-8B5D03355662
Visible to Intel only — GUID: GUID-CF1C51DA-9501-4575-981C-8B5D03355662
DPCT1039
Message
The generated code assumes that <parameter name> points to the global memory address space. If it points to a local memory address space, replace <function name> with <function name>.
Detailed Help
Intel® DPC++ Compatibility Tool deduces whether the first parameter of an atomic function points to a global memory address space or a local memory space, using the last assignment’s RHS value of the first parameter of the atomic function. If the last assignment is in an if/while/do/for statement or the first argument of the atomic function depends on a function or template argument, the deduction result may be incorrect. You need to verify the generated code to determine if the first parameter of the atomic function actually points to the local memory address space. If it does, then replace the atomic function name with an atomic function name that includes the template parameters, as pointed to in the warning message.
Suggestions to Fix
If the first parameter of an atomic function points to a local memory address space, replace the atomic function name with an atomic function name that includes the template parameters.
For example, this original CUDA* code:
__device__ void d(int* s) {
atomicAdd(s, 1);
}
__global__ void k() {
__shared__ int shared;
d(&shared);
}
results in the following migrated SYCL* code:
void d(int* s) {
/*
DPCT1039:0: The generated code assumes that "s" points to the global memory
address space. If it points to a local memory address space, replace
"atomic_fetch_add" with
"atomic_fetch_add<sycl::access::address_space::local_space>".
*/
dpct::atomic_fetch_add(s, 1);
}
void k(int &shared) {
d(&shared);
}
which is rewritten to:
template<sycl::access::address_space as>
void d(int* s) {
dpct::atomic_fetch_add<as>(s, 1);
}
void k(int &shared) {
d<sycl::access::address_space::local_space>(&shared);
}