Developer Guide and Reference

  • 2022.1
  • 04/11/2022
  • Public Content

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:
// 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:
// 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:
// 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.

Product and Performance Information

1

Performance varies by use, configuration and other factors. Learn more at www.Intel.com/PerformanceIndex.