DPCT1086#
メッセージ#
__activemask()
は 0xffffffff
に置き換えられます。コードを調整する必要があります。
詳細な説明#
現在、SYCL* には、__activemask()
と同等の機能はありません。コードにスレッドを非アクティブにするスロー制御がある場合、スレッドのロジックを書き直す必要があります。
例えば、以下のオリジナル CUDA* コードについて考えてみます。
1 __device__ inline int SHFL_SYNC(unsigned mask, int val, unsigned offset,
2 unsigned w = warpSize) {
3 return __shfl_down_sync(mask, val, offset, w);
4 }
5
6 __global__ void kernel(int *array) {
7 unsigned int tid = threadIdx.x;
8 if (tid >= 8)
9 return;
10 unsigned mask = __activemask();
11 array[tid] = SHFL_SYNC(mask, array[tid], 4);
12 }
このコードは、以下の SYCL* コードに移行されます。
1 inline int SHFL_SYNC(unsigned mask, int val, unsigned offset,
2 const sycl::nd_item<3> &item_ct1, unsigned w = 0) {
3 /*
4 DPCT1023:0: The SYCL sub-group does not support mask options for
5 dpct::shift_sub_group_left.You can specify
6 "--use-experimental-features=masked-sub-group-operation" to use the
7 experimental helper function to migrate __shfl_down_sync.
8 */
9 if (!w) w = item_ct1.get_sub_group().get_local_range().get(0);
10 // この呼び出しは、すべてのワーク項目が到着するまで待機しますが、tid < 8 のワーク項目のみがこの呼び出しに遭遇するため、これは発生しません
11 return dpct::shift_sub_group_left(item_ct1.get_sub_group(), val, offset, w);
12 }
13
14 void kernel(int *array, const sycl::nd_item<3> &item_ct1) {
15 unsigned int tid = item_ct1.get_local_id(2);
16 if (tid >= 8)
17 return;
18 /*
19 DPCT1086:1: __activemask() is migrated to 0xffffffff.You may need to adjust
20 the code.
21 */
22 unsigned mask = 0xffffffff;
23 array[tid] = SHFL_SYNC(mask, array[tid], 4, item_ct1);
24 }
このコードは次のように書き換えられます。
1 // 使用されていないため、マスク・パラメーターを削除
2 inline int SHFL_SYNC(int val, unsigned offset,
3 const sycl::nd_item<3> &item_ct1, unsigned w = 0) {
4 if (!w) w = item_ct1.get_sub_group().get_local_range().get(0);
5 unsigned int tid = item_ct1.get_local_id(2);
6 // 一時変数を使用して sycl::shift_group_left() の結果を保存し、すべてのワーク項目がこの呼び出しを実行できるようにします
7 int v_tmp = sycl::shift_group_left(item_ct1.get_sub_group(), val, offset);
8 return (tid < 8) ? v_tmp : val;
9 }
10
11 void kernel(int *array, const sycl::nd_item<3> &item_ct1) {
12 unsigned int tid = item_ct1.get_local_id(2);
13 // 使用されていないため、マスク・パラメーターを削除
14 array[tid] = SHFL_SYNC(array[tid], 4, item_ct1);
15 }
修正方法の提案#
__activemask()
の代わりに 0xffffffff
を使用できるか確認します。使用できない場合は、スレッドのロジックを再設計します。