DPCT1085#

Message#

The function <function name> requires sub-group size to be <size>, while other sub-group functions in the same SYCL kernel require a different sub-group size. You may need to adjust the code.

Detailed Help#

Each kernel can only be decorated with one sub-group size. This warning is emitted when a kernel requires different sub-group sizes. Check if the sub-group size can be unified into one value, and if it cannot be unified, redesign the code logic.

For example, this original CUDA* code:

1  _global_ void kernel() {
2    int Input, Output1, Output2, Lane;
3    ...
4
5    // original code logic
6    Output1 = __shfl(Input, Lane, 32);
7    Output2 = __shfl_xor(Input, Lane, 16);
8  }

results in the following migrated SYCL* code:

1  void kernel(int WarpSize, sycl::nd_item<3> item_ct1) {
2    int Input, Output1, Output2, Lane;
3    ...
4
5    // original code logic
6    Output1 = Item_ct1.get_sub_group().shuffle(Input, Lane, 32);
7    /* DPCT1085 */
8    Output2 = Item_ct1.get_sub_group().shuffle_xor(Input, Lane, 16);
9  }

which is manually adjusted to:

 1  void kernel(int WarpSize, sycl::nd_item<3> item_ct1) {
 2    int Input, Output, SrcLane;
 3    ...
 4
 5    // redesigned code logic
 6    Output1 = Item_ct1.get_sub_group().shuffle(Input, Lane, 32);
 7    /* DPCT1085 */
 8    // redesign the code logic to unify sub-group size in the same kernel
 9    Output2 = Item_ct1.get_sub_group().shuffle_xor(Input, Lane, 32);
10  }

Suggestions to Fix#

Code requires manual fix. Rewrite the code manually.