Intel® DPC++ Compatibility Tool Developer Guide and Reference
A newer version of this document is available. Customers should click here to go to the newest version.
DPCT1118
Message
SYCL group functions and algorithms must be encountered in converged control flow.
Detailed Help
SYCL* group functions and algorithms must be encountered in converged control flow for all work-items in a work-group. If a work-item completes the kernel and exits early without reaching a collective work-group operation like barrier, all the other work-items in the work-group reaching the collective operation will wait for the exited work-item.
Refer to How do I fix the issue of SYCL* code hanging due to work group level synchronization, such as a group barrier used in a conditional statement? for additional information.
Suggestions to Fix
For example, this original CUDA* code:
__global__ void kernel(float *data) {
int tid = threadIdx.x;
if (tid < 32) {
if (data[tid] < data[tid + 32]) {
data[tid] = data[tid + 32];
}
__syncthreads();
...
}
}results in the following migrated SYCL code:
void kernel(float *data, const sycl::nd_item<3> &item_ct1) {
int tid = item_ct1.get_local_id(2);
if (tid < 32) {
if (data[tid] < data[tid + 32]) {
data[tid] = data[tid + 32];
}
/*
DPCT1118:0: SYCL group functions and algorithms must be encountered in converged control flow. You should check this condition holds.
*/
/*
DPCT1065:1: Consider replacing sycl::nd_item::barrier() with sycl::nd_item::barrier(sycl::access::fence_space::local_space) for better performance if there is no access to global memory.
*/
item_ct1.barrier();
}
}which is rewritten to:
void kernel(float *data, const sycl::nd_item<3> &item_ct1) {
int tid = item_ct1.get_local_id(2);
if (tid < 32) {
if (data[tid] < data[tid + 32]) {
data[tid] = data[tid + 32];
}
}
item_ct1.barrier();
...
}