DPCT1085

Message

The function %0 requires sub-group size to be %1, while other sub-group functions in same DPC++ 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 code:

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
// original code:

_global_ void kernel() {
  int Input, Output1, Output2, Lane;
  ...

  // original code logic
  Output1 = __shfl(Input, Lane, 32);
  Output2 = __shfl_xor(Input, Lane, 16);
}

results in the following migrated DPC++ code:

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
// migrated DPC++ code:

void kernel(int WarpSize, sycl::nd_item<3> item_ct1) {
  int Input, Output1, Output2, Lane;
  ...

  // original code logic
  Output1 = Item_ct1.get_sub_group().shuffle(Input, Lane, 32);
  /* DPCT1085 */
  Output2 = Item_ct1.get_sub_group().shuffle_xor(Input, Lane, 16);
}

which is manually adjusted to:

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
// adjusted DPC++ code:

void kernel(int WarpSize, sycl::nd_item<3> item_ct1) {
  int Input, Output, SrcLane;
  ...

  // redesigned code logic
  Output1 = Item_ct1.get_sub_group().shuffle(Input, Lane, 32);
  /* DPCT1085 */
  // redesign the code logic to unify sub-group size in the same kernel
  Output2 = Item_ct1.get_sub_group().shuffle_xor(Input, Lane, 32);
}

Suggestions to Fix

Code requires manual fix. Rewrite the code manually.