DPCT1115

Intel® DPC++ Compatibility Tool Developer Guide and Reference

Download PDF

ID 768918

Date 11/07/2023

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-30250AED-EA5B-40E2-B158-0D736F406924

View Details

DPCT1115

Message

The sycl::ext::oneapi::group_local_memory is used to allocate group-local memory at the none kernel functor scope of a work-group data parallel kernel. You may need to adjust the code.

Detailed Help

The sycl::ext::oneapi::group_local_memory can be used to allocate group-local memory at the kernel functor scope of a work-group data parallel kernel. The restriction that group-local variables must be defined at kernel functor scope may be lifted in a future version of this extension.

Refer to sycl_ext_oneapi_local_memory.asciidoc for more details.

Suggestions to Fix

For example, this original CUDA* code:

template <int S> __device__ void devfun() {
  __shared__ int slm1[32 * S];
  ...
}

template <int S> __global__ void kernel() {
  __shared__ int slm2[S];
  devfun<S>();
}

void hostfun() { kernel<256><<<1, 1>>>(); }

results in the following migrated SYCL* code:

template <int S> inline void devfun(int *p, const sycl::nd_item<3> &item_ct1) {
  /*
  DPCT1115:0: The sycl::ext::oneapi::group_local_memory is used to allocate
  group-local memory at the none kernel functor scope of a work-group data
  parallel kernel. You may need to adjust the code.
  */
  auto &slm1 =
      *sycl::ext::oneapi::group_local_memory<int[32 * S]>(item_ct1.get_group());
  ...
}

template <int S> __dpct_inline__ void kernel(const sycl::nd_item<3> &item_ct1) {
  auto &slm2 =
      *sycl::ext::oneapi::group_local_memory<int[S]>(item_ct1.get_group());
  devfun<S>(item_ct1);
}

void hostfun() { dpct::get_default_queue().parallel_for(
  sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
  [=](sycl::nd_item<3> item_ct1) {
    kernel<256>(item_ct1);
  });
}

which is rewritten to:

template <int S> inline void devfun(int *slm1) {
  ...
}

template <int S> __dpct_inline__ void kernel(int *slm1, int *slm2) {

  devfun<S>(slm1);
}

void hostfun() { dpct::get_default_queue().submit(
  [&](sycl::handler &cgh) {
    sycl::local_accessor<int, 1> slm1_acc_ct1(sycl::range<1>(32 * 256), cgh);
    sycl::local_accessor<int, 1> slm2_acc_ct1(sycl::range<1>(256), 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) {
        kernel<256>(slm1_acc_ct1.get_pointer(), slm2_acc_ct1.get_pointer());
      });
  });

Select Your Language

Using Intel.com Search

Quick Links

Recent Searches

Advanced Search

Only search in

Intel® DPC++ Compatibility Tool Developer Guide and Reference

DPCT1115

Message

Detailed Help

Suggestions to Fix