DPCT1108#

Message#

<original API> was migrated with the experimental feature <feature name> which may not be supported by all compilers or runtimes. You may need to adjust the code.

Detailed Help#

To support better migration, SYCLomatic provides the -use-experimental-features option to apply experimental features during migration. <feature name> may not be supported by all SYCL compilers or runtimes. If the target SYCL compiler and runtime does not support <feature name>, you need to adjust the code and not use <feature name>.

Suggestions to Fix#

For example, this original CUDA* code:

 1__global__ void kernel() {
 2  int lane_id = threadIdx.x % 32;
 3  int foo = 0, result = 0;
 4  int mask = 0xf;
 5  if (lane_id == 0) {
 6    result = 10;
 7  }
 8  if (lane_id & mask) {
 9    foo = __shfl_sync(mask, result, 0);
10  }
11}

results in the following migrated SYCL code:

 1void kernel(const sycl::nd_item<3> &item_ct1) {
 2  int lane_id = item_ct1.get_local_id(2) % 32;
 3  int foo = 0, result = 0;
 4  int mask = 0xf;
 5  if (lane_id == 0) {
 6    result = 10;
 7  }
 8  if (lane_id & mask) {
 9    /*
10    DPCT1108:0: '__shfl_sync' was migrated with the experimental feature masked
11    sub_group function which may not be supported by all compilers or runtimes.
12    You may need to adjust the code.
13    */
14    foo = dpct::experimental::select_from_sub_group(
15        mask, item_ct1.get_sub_group(), result, 0);
16  }
17}

which is rewritten to:

 1void kernel(const sycl::nd_item<3> &item_ct1) {
 2  int lane_id = item_ct1.get_local_id(2) % 32;
 3  int foo = 0, result = 0;
 4  int mask = 0xf;
 5  if (lane_id == 0) {
 6    result = 10;
 7  }
 8  // If not use experimental feature masked sub-group function
 9  int foo_tmp = dpct::select_from_sub_group(item_ct1.get_sub_group(), result, 0);
10  if (lane_id & mask) {
11    foo = foo_tmp;
12  }
13}