DPCT1116#

メッセージ#

<original API> はパフォーマンス向上にため <migrated API> に移行され、<expression> は (UINT_MAX + 1) / ('<argument>' + 1) で計算されます。この移行では <value> の初期値を <expression> を乗算してスケーリングする必要があり、アトミック関数の外部で <value> の値を使用するには <expression> を除算してスケーリングする必要があります。

詳細な説明#

デフォルトでは、移行ツールは比較交換アトミック操作を使用して atomicInc と atomicDec 関数を移行します。アトミック API がクリティカル・パス内にある場合、パフォーマンスに影響します。

パフォーマンスを向上するため、移行ツールは atomicInc/atomicDec 関数の 2 番目の引数の値をチェックします。2n – 1 (1 < n <= 32) の場合、ツールはアトミック add API を使用して atomicInc/atomicDec を移行します。

例えば、atomicInc(addr, val) の場合、最適化では次のことが行われます。

  • atomicInc(addr, val) 関数を dpct::atomic_fetch_add(addr, step) に移行します。ここで、step は (UINT_MAX + 1) / (val + 1) で計算されます。この方法は、符号なし整数のオーバーフローを利用して、atomicInc/Dec() のラッピング動作を模倣します。

  • *addr の初期値を更新して step を乗算し、ラッピング動作の一貫性を保証します。

  • アトミック変数の外部で *addr 参照の値を更新し、step ごとに分割して必要とする値を取得します。

修正方法の提案

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

1  __device__ unsigned int a1 = 5; 
2  __global__ void kernel(unsigned int *result){ 
3     ...
4   unsigned int old_val = atomicInc(&a1, 0x7fffffff); 
5     ... 
6   *result = a1; 
7  }

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

1  dpct::global_memory<unsigned int, 0> a1(5 * 2); 
2  void kernel(unsigned int*result, unsigned int &a1){ 
3      ... 
4   /* 
5   DPCT1116:0: The atomicInc was migrated to dpct::atomic_fetch_add(&a1, 2) / 2 for performance, and 2 is computed by (UINT_MAX + 1) / ('0x7fffffff' + 1).This migration requires the initial value of 'a1' to be scaled by multiplying 2, and any usage of the value of 'a1' outside atomic function to be scaled by dividing 2.
6   */ 
7   unsigned int old_val = dpct::atomic_fetch_add<sycl::access::address_space::generic_space>(&a1, 2) / 2; 
8      ... 
9   *result = a1 / 2; // Dividing a1 by step 2 to get the intended value. 
10 }

このコードは次のように書き換えられます。

1  dpct::global_memory<unsigned int, 0> a1(5 * 2); 
2  void kernel(unsigned int*result, unsigned int &a1){ 
3      ... 
4   unsigned int old_val = dpct::atomic_fetch_add<sycl::access::address_space::generic_space>(&a1, 2) / 2; 
5    ... 
6   *result = a1 / 2; // Dividing a1 by step 2 to get the intended value. 
7  }