DPCT1117

メッセージ

<argument 2> には atomicInc/Dec 操作があり、<argument 2> の値はパフォーマンス向上のため移行によってスケーリングされました。DPCT1116 を参照してください。<argument 2> から <argument 1> の値を使用する場合も、<expression> を除算または乗算することでスケーリングする必要があります。ソースコードの調整が必要な場合があります。

詳細な説明

atomicInc/atomicDec の移行には、最適化された手法が適用されます。この最適化された手法を採用するには、移行後に値の一貫性と正当性を保証するため、スケーリングの調整が必要になります。詳細については、DPCT1116 を参照してください。

コードでアトミック変数のアドレスを使用する場合、アトミック変数のアドレスを介したすべての値の割り当てや演算後の参照も、ステップを乗算/除算することでスケーリングする必要があります。アトミック変数は、すべてのホスト側のメモリー操作に対し適切にスケーリングする必要があります。

修正方法の提案

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

 __device__ unsigned int a1; 
 __global__ void kernel(){ 
    ...
  unsigned int old_val = atomicInc(&a1, 0x7fffffff); 
     ... 
 } 
 int main(){ 
  unsigned int result; 
  unsigned int *d_addr; 
 cudaGetSymbolAddress((void **)&d_addr, a1); 
 cudaMemset(d_addr, 2, sizeof(unsigned int)); 
 kernel<<<1, 100>>>(); 
    ... 
 cudaMemcpy(&result, d_addr, sizeof(unsigned int), cudaMemcpyDeviceToHost); 
    ... 
}

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

  dpct::global_memory<unsigned int, 0> a1; 
  void kernel(unsigned int &a1){ 
      ... 
   /* 
   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 value of ‘a1’ outside atomic function to be scaled by dividing 2.
   */ 
   unsigned int old_val = dpct::atomic_fetch_add<sycl::access::address_space::generic_space>(&a1, 2) / 2; 
      ... 
  } 
 int main(){ 
  dpct::device_ext &dev_ct1 = dpct::get_current_device(); 
  sycl::queue &q_ct1 = dev_ct1.default_queue(); 
  unsigned int result; 
  unsigned int *d_addr; 
  /* 
  DPCT1117:3: There is atomicInc/Dec operation on 'a1' and value of 'a1' was scaled by migration for performance, refer to DPCT1116.Using value of 'a1' through 'd_addr' should also be scaled by dividing or multiplying 2.You may need to adjust the code.  */ 
  *((void **)&d_addr) = a1.get_ptr(); 
  q_ct1.memset(d_addr, 2, sizeof(unsigned int)).wait(); 
  q_ct1.submit( 
  [&](sycl::handler &cgh) { 
  a1.init(); 
  auto a1_ptr_ct1 = a1.get_ptr(); 
  cgh.parallel_for( 
    sycl::nd_range<3>(sycl::range<3>(1, 1, 100), sycl::range<3>(1, 1, 100)), 
    [=](sycl::nd_item<3> item_ct1) { 
    kernel1(*a1_ptr_ct1); 
    }); 
  }); 
  ... 
  q_ct1.memcpy(&result, d_addr, sizeof(unsigned int)).wait(); 
  ... 
}

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

 dpct::global_memory<unsigned int, 0> a1; 
 void kernel(unsigned int &a1){ 
     ... 
  unsigned int old_val = dpct::atomic_fetch_add<sycl::access::address_space::generic_space>(&a1, 2) / 2; 
     ... 
 } 
 int main(){ 
  dpct::device_ext &dev_ct1 = dpct::get_current_device(); 
  sycl::queue &q_ct1 = dev_ct1.default_queue(); 
 unsigned int result; 
 unsigned int *d_addr; 
 *((void **)&d_addr) = a1.get_ptr(); 
 q_ct1.memset(d_addr, 2 * 2, sizeof(unsigned int)).wait(); 
 q_ct1.submit( 
   [&](sycl::handler &cgh) { 
   a1.init(); 
   auto a1_ptr_ct1 = a1.get_ptr(); 
     cgh.parallel_for( 
     sycl::nd_range<3>(sycl::range<3>(1, 1, 100), sycl::range<3>(1, 1, 100)), 
     [=](sycl::nd_item<3> item_ct1) { 
     kernel1(*a1_ptr_ct1); 
     }); 
  }); 
   ... 
 q_ct1.memcpy(&result, d_addr, sizeof(unsigned int)).wait(); 
 result = result / 2; 
   ... 
}