Visible to Intel only — GUID: GUID-B2CCADE4-544F-41C8-A322-A48CF7B7EE28
Visible to Intel only — GUID: GUID-B2CCADE4-544F-41C8-A322-A48CF7B7EE28
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();
...
}