Visible to Intel only — GUID: GUID-B31D758C-5D0F-4979-AB98-5126490F2329
Visible to Intel only — GUID: GUID-B31D758C-5D0F-4979-AB98-5126490F2329
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:
__global__ void kernel() {
int lane_id = threadIdx.x % 32;
int foo = 0, result = 0;
int mask = 0xf;
if (lane_id == 0) {
result = 10;
}
if (lane_id & mask) {
foo = __shfl_sync(mask, result, 0);
}
}
results in the following migrated SYCL code:
void kernel(const sycl::nd_item<3> &item_ct1) {
int lane_id = item_ct1.get_local_id(2) % 32;
int foo = 0, result = 0;
int mask = 0xf;
if (lane_id == 0) {
result = 10;
}
if (lane_id & mask) {
/*
DPCT1108:0: '__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.
*/
foo = dpct::experimental::select_from_sub_group(
mask, item_ct1.get_sub_group(), result, 0);
}
}
which is rewritten to:
void kernel(const sycl::nd_item<3> &item_ct1) {
int lane_id = item_ct1.get_local_id(2) % 32;
int foo = 0, result = 0;
int mask = 0xf;
if (lane_id == 0) {
result = 10;
}
// If not use experimental feature masked sub-group function
int foo_tmp = dpct::select_from_sub_group(item_ct1.get_sub_group(), result, 0);
if (lane_id & mask) {
foo = foo_tmp;
}
}