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  // This call will wait for all work-items to arrive which will never happen since only work-items with tid < 8 will encounter this call. 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   // remove mask parameter, as it is not used 
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   // Use a temporary variable to save the result of sycl::shift_group_left() to make sure all work-items can encounter this call. 
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  // remove mask parameter, as it is not used 
14    array[tid] = SHFL_SYNC(array[tid], 4, item_ct1); 
15  }

修正方法の提案#

__activemask() の代わりに 0xffffffff を使用できるか確認します。使用できない場合は、スレッドのロジックを再設計します。