DPCT1039

Intel® DPC++ Compatibility Tool Developer Guide and Reference

Download PDF

ID 768918

Date 6/24/2024

Version

Public

A newer version of this document is available. Customers should click here to go to the newest version.

Visible to Intel only — GUID: GUID-9D5FCBF3-D6DF-444D-A741-822434440325

View Details

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);
 }

Select Your Language

Using Intel.com Search

Quick Links

Recent Searches

Advanced Search

Only search in

Intel® DPC++ Compatibility Tool Developer Guide and Reference

DPCT1039

Message

Detailed Help

Suggestions to Fix