DPCT1085#

メッセージ#

関数 <function name> は sub-group サイズが <size> である必要がありますが、同じ SYCL* カーネル内のほかの sub-group 関数では異なる sub-group サイズが必要です。ソースコードの調整が必要な場合があります。

詳細な説明#

各カーネルは、1 つの sub-group サイズでしか修飾できません。この警告は、カーネルが異なる sub-group サイズを必要とする場合に出力されます。sub-group サイズを 1 つの値に統一できるかどうかを確認し、統一できない場合はコードのロジックを再設計します。

例えば、以下のオリジナル CUDA* コードについて考えてみます。

1   __global__ void kernel(int* data1, int* data2) { 
2     typedef cub::WarpScan<int> WarpScan; 
3     typedef cub::WarpScan<int, 16> WarpScan16; 
4 
5     typename WarpScan::TempStorage temp1; 
6     typename WarpScan16::TempStorage temp2; 
7 
8     int input = data1[threadIdx.x]; 
9     int output1 = 0; 
10    int output2 = 0; 
11    WarpScan(temp1).InclusiveSum(input, output1); 
12    data1[threadIdx.x] = output1; 
13    WarpScan16(temp2).InclusiveSum(input, output2); 
14    data2[threadIdx.x] = output1; 
15  } 
16 
17  void foo(int* data1, int* data2) { 
18    kernel<<<1, 32>>>(data1, data2); 
19  }

このコードは、以下の SYCL* コードに移行されます。

1   void kernel(int* data1, int* data2, const sycl::nd_item<3> &item_ct1) { 
2 
3     int input = data1[item_ct1.get_local_id(2)]; 
4     int output1 = 0; 
5     int output2 = 0; 
6     output1 = sycl::inclusive_scan_over_group(item_ct1.get_sub_group(), input, 
7     sycl::plus<>()); 
8     data1[item_ct1.get_local_id(2)] = output1; 
9   /* 
10  DPCT1085:0: The function inclusive_scan_over_group requires sub-group size to 
11  be 16, while other sub-group functions in the same SYCL kernel require a 
12  different sub-group size. You may need to adjust the code. 13  */ 
14    output2 = sycl::inclusive_scan_over_group(item_ct1.get_sub_group(), input, 
15    sycl::plus<>()); 
16    data2[item_ct1.get_local_id(2)] = output1; 
17  } 
18 
19  void foo(int* data1, int* data2) { 
20    dpct::get_in_order_queue().parallel_for( 
21      sycl::nd_range<3>(sycl::range<3>(1, 1, 32), sycl::range<3>(1, 1, 32)), 
22      [=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(32)]] { 
23      kernel(data1, data2, item_ct1); 
24    }); 
25  } 
26 
27  void foo(int* data) { 
28    dpct::get_in_order_queue().parallel_for( 
29      sycl::nd_range<3>(sycl::range<3>(1, 1, 32), sycl::range<3>(1, 1, 32)), 
30      [=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(32)]] { 
31      kernel(data, item_ct1);
32    }); 
33  }

このコードを以下のように手動で調整します。

1   void kernel(int* data1, int* data2, const sycl::nd_item<3> &item_ct1) { 
2 
3     int input = data1[item_ct1.get_local_id(2)]; 
4     int output1 = 0; 
5     int output2 = 0; 
6     output1 = sycl::inclusive_scan_over_group(item_ct1.get_sub_group(), input, 
7     sycl::plus<>()); 
8     data1[item_ct1.get_local_id(2)] = output1; 
9     output2 = sycl::inclusive_scan_over_group(item_ct1.get_sub_group(), input, 
10    sycl::plus<>()); 
11    data2[item_ct1.get_local_id(2)] = output1; 
12    item_ct1.barrier(); 
13    if (item_ct1.get_local_id(2) % 32 >= 16) { 
14      int warp_id = item_ct1.get_local_id(2) / 32; 
15      data2[item_ct1.get_local_id(2)] -= data2[warp_id * 32 + 15]; 
16    } 
17  } 
18 
19  void foo(int* data1, int* data2) { 
20    dpct::get_in_order_queue().parallel_for( 
21      sycl::nd_range<3>(sycl::range<3>(1, 1, 32), sycl::range<3>(1, 1, 32)), 
22      [=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(32)]] { 
23      kernel(data1, data2, item_ct1); 
24    }); 
25  }

修正方法の提案#

コードを手動で修正する必要があります。このコードを手動で書き換えてください。