Intel® DPC++ Compatibility Tool Developer Guide and Reference

Developer Guide and Reference

  • 2022.1
  • 04/11/2022
  • Public Content
Contents

DPCT1089

Message

The value of the sub-group size attribute argument
<argument name>
 cannot be evaluated by the Intel(R) DPC++ Compatibility Tool. Replace "dpct_placeholder" with an integral constant expression.

Detailed Help

The argument to the attribute must be an integral constant expression. This warning is emitted when the Intel® DPC++ Compatibility Tool cannot evaluate the value of the sub-group size argument as an integral constant expression. Check if the sub-group size attribute argument can be replaced by an integral constant expression, and if not, redesign the code logic.
For example, this original code:
// original code

_global_ void kernel(int WarpSize) {   int Input, Output, Lane;   ...   Output = __shfl(Input, Lane, WarpSize); }
...
if(condition)
{   kernel<<<GirdSize, BlockSize>>>(16); }
else
{   kernel<<<GirdSize, BlockSize>>>(32); }
results in the following migrated DPC++ code:
// migrated DPC++ code

void kernel(int WarpSize, sycl::nd_item<3> item_ct1)
{   int Input, Output, Lane;   ...   Output = Item_ct1.get_sub_group().shuffle(Input, Lane); }
...
if(condition) {
/* DPCT1089 */
q_ct1.parallel_for(
    sycl::nd_range<3>(GridSize * BlockSize, BlockSize),
    [=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(dpct_placeholder)]] {       kernel(16, item_ct1); });
} else {
/* DPCT1089 */
q_ct1.parallel_for(
    sycl::nd_range<3>(GridSize * BlockSize, BlockSize),
    [=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(dpct_placeholder)]] {       kernel(32, item_ct1); }
which is manually adjusted to:
// fixed DPC++ code

void kernel(int WarpSize, sycl::nd_item<3> item_ct1) {   int Input, Output, Lane;   ...   Output = Item_ct1.get_sub_group().shuffle(Input, Lane); }
...
if(condition) {
/* DPCT1089 */
q_ct1.parallel_for(
    sycl::nd_range<3>(GridSize * BlockSize, BlockSize),
    [=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(16)]]
{       kernel(16, item_ct1); }
);
} else {
/* DPCT1089 */
q_ct1.parallel_for(
    sycl::nd_range<3>(GridSize * BlockSize, BlockSize),
    [=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(32)]]
{       kernel(32, item_ct1); }

Suggestions to Fix

Code requires manual fix. Replace “dpct_placeholder” with an integral constant expression.

Product and Performance Information

1

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