DPCT1121#

Message#

Make sure that the <argument> which is used in the SYCL group function/algorithm is initialized.

Detailed Help#

For SYCL group function/algorithm in SYCL* spec 4.17.2 and 4.17.3, the input scalar value must be initialized by all work-items in the group. This is necessary for compatibility with standard C++ rules and to ensure that compilers do not optimize too aggressively.

Suggestions to Fix#

For example, this original CUDA* code:

 1__device__ void d(int var) {
 2  if (threadIdx.x == 0)
 3    var = 123;
 4  var = __shfl_sync(0xffffffff, var, 0);
 5}
 6
 7__global__ void kernel() {
 8  int var;
 9  d(var);
10}

results in the following migrated SYCL* code:

 1void d(int var, const sycl::nd_item<3> &item_ct1) {
 2  if (item_ct1.get_local_id(2) == 0)
 3    var = 123;
 4  /*
 5  DPCT1121:0: Make sure that the "var" which is used in the SYCL group
 6  function/algorithm is initialized.
 7  */
 8  var = dpct::select_from_sub_group(item_ct1.get_sub_group(), var, 0);
 9}
10void kernel(const sycl::nd_item<3> &item_ct1) {
11  int var;
12  d(var, item_ct1);
13}

which needs to be rewritten to:

1void d(int var, const sycl::nd_item<3> &item_ct1) {
2  if (item_ct1.get_local_id(2) == 0)
3    var = 123;
4  var = dpct::select_from_sub_group(item_ct1.get_sub_group(), var, 0);
5}
6void kernel(const sycl::nd_item<3> &item_ct1) {
7  int var = 0;
8  d(var, item_ct1);
9}