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); 
    }); 
}