DPCT1113

メッセージ

<function name> が多次元カーネルで呼び出される場合、sycl::nd_item::barrier(sycl::access::fence_space::local_space) を sycl::nd_item::barrier() に置き換えることを検討してください。

詳細な説明

work-group 内の各ワーク項目間でグローバル・デバイス・メモリーへの重複するアクセスがない場合、__syncthreads() API は、sycl::nd_item::barrier(sycl::access::fence_space::local_space) に移行可能なカーネル実行の 1D インデックス空間を使用してカーネルを呼び出します。2D/3D インデックス空間を使用してカーネルを呼び出す場合、work-group 内の各ワーク項目からグローバルメモリーへのアクセスが重複する可能性があり、バリアを超えたワーク項目間でデータの依存関係が生じる可能性があります。この場合、sycl::nd_item::barrier() 呼び出しを sycl::access::fence_space::global_and_local に置き換える必要があります。

修正方法の提案

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

  __global__ void kernel(float *mem) { 
   unsigned int tid = threadIdx.x + blockDim.x * blockIdx.x; 
     ...
   mem[tid] = ...; 
   __syncthreads(); 
      ... 
      ... = mem[tid]; 
} 
 
 void foo(float *mem) { 
  kernel<<<16, 16>>>(mem); //1D index space of a SYCL kernel execution 
 }

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

 void kernel(float *mem, const sycl::nd_item<3> &item_ct1) { 
  unsigned int tid = item_ct1.get_local_id(2) + 
  item_ct1.get_local_range(2) * item_ct1.get_group(2); 
     ... 
  mem[tid] = ...; // global memory access without overlap among each work-item in a work-group 
  /* 
  DPCT1113:0: Consider replacing 
  sycl::nd_item::barrier(sycl::access::fence_space::local_space) with 
  sycl::nd_item::barrier() if function "kernel" is called in a multidimensional 
 kernel.
 */ 
 item_ct1.barrier(sycl::access::fence_space::local_space); 
    ... 
    ... = mem[tid]; // global memory access without overlap among each work-item in a work-group 
} 

void foo(float *mem) { 
 dpct::get_default_queue().parallel_for( 
   sycl::nd_range<3>(sycl::range<3>(1, 1, 16) * sycl::range<3>(1, 1, 16), 
   sycl::range<3>(1, 1, 16)), 
   [=](sycl::nd_item<3> item_ct1) { 
   kernel(mem, item_ct1); 
 }); 
}

移行後、移行された SYCL* カーネルコードを 2D カーネルに更新します。

 void kernel(float *mem, const sycl::nd_item<3> &item_ct1) { 
  unsigned int tidx = item_ct1.get_local_id(2) + 
  item_ct1.get_local_range(2) * item_ct1.get_group(2); 
  unsigned int tidy = item_ct1.get_local_id(1) + 
  item_ct1.get_local_range(1) * item_ct1.get_group(1); 

     ... 
  mem[tidx] = ...; // global memory access with overlap among each work-item in a work-group 
  mem[tidy] = ...; // global memory access with overlap among each work-item in a work-group 
 /* 
 DPCT1113:0: Consider replacing 
  sycl::nd_item::barrier(sycl::access::fence_space::local_space) with 
  sycl::nd_item::barrier() if function "kernel" is called in a multidimensional 
 kernel.
 */ 
 item_ct1.barrier(sycl::access::fence_space::local_space); 
    ... 
    ... = mem[tidx]; // global memory access with overlap among each work-item in a work-group 
    ... = mem[tidy]; // global memory access with overlap among each work-item in a work-group 
} 

void foo(float *mem) { 
 dpct::get_default_queue().parallel_for( 
   sycl::nd_range<3>(sycl::range<3>(1, 4, 4) * sycl::range<3>(1, 4, 4), 
   sycl::range<3>(1, 4, 4)), /*2D index space of a SYCL kernel execution */ 

   [=](sycl::nd_item<3> item_ct1) { 
   kernel(mem, item_ct1); 
 }); 
}

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

void kernel(float *mem, const sycl::nd_item<3> &item_ct1) { 
unsigned int tidx = item_ct1.get_local_id(2) + 
item_ct1.get_local_range(2) * item_ct1.get_group(2); 
unsigned int tidy = item_ct1.get_local_id(1) + 
item_ct1.get_local_range(1) * item_ct1.get_group(1); 

   ... 
  mem[tidx] = ...; // global memory access with overlap among each work-item in a work-group 
  mem[tidy] = ...; // global memory access with overlap among each work-item in a work-group 
 item_ct1.barrier(sycl::access::fence_space::global_and_local); 
    ... 
    ... = mem[tidx]; // global memory access with overlap among each work-item in a work-group 
    ... = mem[tidy]; // global memory access with overlap among each work-item in a work-group 
} 

void foo(float *mem) { 
 dpct::get_default_queue().parallel_for( 
   sycl::nd_range<3>(sycl::range<3>(1, 4, 4) * sycl::range<3>(1, 4, 4), 
   sycl::range<3>(1, 4, 4)), 
   [=](sycl::nd_item<3> item_ct1) { 
   kernel(mem, item_ct1); 
 }); 
}