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 を使用できるか確認します。使用できない場合は、スレッドのロジックを再設計します。