DPCT1108
Contents
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, Intel® DPC++ Compatibility Tool 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 target = 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 // CHECK: /*
10 // CHECK: DPCT1108:{{[0-9]+}}: ‘__shfl_sync’ was migrated with the experimental feature masked sub_group function which may not be supported by all compilers or runtimes. You may need to adjust the code.
11 // CHECK: */
12 foo = dpct::select_from_sub_group(mask, item_ct1.get_sub_group(), 10, 0);
13 }
14}
which is rewritten to:
1void kernel(const sycl::nd_item<3> &item_ct1){
2 int target = 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 // Don’t 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}