DPCT1089#
Message#
The value of the sub-group size attribute argument <argument name> cannot be evaluated. 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 SYCLomatic 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 CUDA* code:
1 _global_ void kernel(int WarpSize) { int Input, Output, Lane; ... Output = __shfl(Input, Lane, WarpSize); }
2 ...
3 if(condition)
4 { kernel<<<GirdSize, BlockSize>>>(16); }
5 else
6 { kernel<<<GirdSize, BlockSize>>>(32); }
results in the following migrated SYCL* code:
1 void kernel(int WarpSize, sycl::nd_item<3> item_ct1) {
2 int Input, Output, Lane; ... Output = Item_ct1.get_sub_group().shuffle(Input, Lane);
3 }
4 ...
5 if(condition) {
6 /* DPCT1089 */
7 q_ct1.parallel_for(
8 sycl::nd_range<3>(GridSize * BlockSize, BlockSize),
9 [=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(dpct_placeholder)]] { kernel(16, item_ct1); });
10 } else {
11 /* DPCT1089 */
12 q_ct1.parallel_for(
13 sycl::nd_range<3>(GridSize * BlockSize, BlockSize),
14 [=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(dpct_placeholder)]] { kernel(32, item_ct1); }
15 }
which is manually adjusted to:
1 void kernel(int WarpSize, sycl::nd_item<3> item_ct1) { int Input, Output, Lane; ... Output = Item_ct1.get_sub_group().shuffle(Input, Lane); }
2 ...
3 if(condition) {
4 /* DPCT1089 */
5 q_ct1.parallel_for(
6 sycl::nd_range<3>(GridSize * BlockSize, BlockSize),
7 [=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(16)]]
8 { kernel(16, item_ct1); }
9 );
10 } else {
11 /* DPCT1089 */
12 q_ct1.parallel_for(
13 sycl::nd_range<3>(GridSize * BlockSize, BlockSize),
14 [=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(32)]]
15 { kernel(32, item_ct1); }
16 }
Suggestions to Fix#
Code requires manual fix. Replace “dpct_placeholder” with an integral constant expression.