DPCT1114#

メッセージ#

元のコードではソースのホストメモリーがページング可能なメモリーであることを想定して、cudaMemcpy は非同期 memcpy に移行されます。メモリーがページング可能でない場合、memcpy API がイベントを返すと wait() を呼び出すことで同期動作を保証します。

詳細な説明#

cudaMemcpy 関数は通常、同期を行います。ただし、ページング可能なホストメモリーを使用してホストからデバイスにコピーする場合、動作は非同期となります。移行時に --optimize-migration オプションを使用すると、移行ツールはホストメモリーがページング可能であると想定して、cudaMemcpy を非同期 memcpy に移行します。これにより、他のタスクとの同時メモリー転送が可能となり、パフォーマンスが向上します。ソースメモリーがホストメモリーに固定されている場合、ユーザーが同期動作を保証するため memcpy API が返すイベントで wait() する必要があります。

修正方法の提案

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

1  Int N = 100; 
2  float *src, *dst; 
3  cudaMalloc(&dst, sizeof(float) * N); 
4  cudaMallocHost(&src, sizeof(float) * N); 
5  for(int i = 0; i < N; i++){ 
6   Src[i] = i; 
7  } 
8  cudaMemcpy(dst, src, sizeof(float) * N, cudaMemcpyHostToDevice);

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

1  sycl::device dev_ct1; 
2  sycl::queue q_ct1(dev_ct1, sycl::property_list{sycl::property::queue::in_order()}); 
3  float *src, *dst; 
4  dst = sycl::malloc_device<float>(N, q_ct1); 
5  src = sycl::malloc_host<float>(N, q_ct1); 
6  for(int i = 0; i < N; i++){ 
7   src[i] = i; 
8  } 
9  /* 
10 DPCT1114:1: cudaMemcpy is migrated to asynchronization memcpy, assuming in the original code the source host memory is pageable memory.If the memory is not pageable, call wait() on event return by memcpy API to ensure synchronization behavior.
11 */ 
12 q_ct1.memcpy(dst, src, sizeof(float) * N);

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

1  sycl::device dev_ct1; 
2  sycl::queue q_ct1(dev_ct1, sycl::property_list{sycl::property::queue::in_order()}); 
3  float *src, *dst; 
4  dst = sycl::malloc_device<float>(N, q_ct1); 
5  src = sycl::malloc_host<float>(N, q_ct1); 
6  for(int i = 0; i < N; i++){ 
7   src[i] = i; 
8  } 
9  q_ct1.memcpy(dst, src, sizeof(float) * N).wait(); // src is allocated by cudaMallocHost with page-locked memory on host, so call wait().