Intel® DPC++ Compatibility Tool Developer Guide and Reference

Developer Guide and Reference

  • 2022.0
  • 12/06/2021
  • Public Content
Contents

DPCT1089

Message

The value of the subgroup 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 subgroup size argument as an integral constant expression. Check if the subgroup 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.