DPCT1020#

メッセージ#

<api name> の移行が __global__ または __device__ 関数から呼び出される場合、サポートされません。代わりにホスト側の <api name> を使うようにコードを再設計する必要があります。この場合、SYCL* キューにこの呼び出しが自動的に送信されます。

説明#

この警告は、<api name> が SYCL* カーネルをコマンドキューに送信し、<api-name> の呼び出し元がコマンドキューに送信された SYCL* カーネルの場合に生成されます。これにより、デバイス側でカーネルをエンキューすることになりますが、これは SYCL* 1.2.1 ではサポートされていません。

修正方法の提案#

ホスト側の API を使うようにコードを再設計します。そうすることで、SYCL* キューにこの呼び出しが自動的に送信されます。

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

1__global__ void kernel(float *d_data) { 
2 int tid = threadIdx.x; 
3 d_data[tid + 1] = tid; 
4 
5 __syncthreads(); 
6 
7 if (tid == 0) { 
8 cublasHandle_t handle; 
9 cublasCreate(&handle); 
10 cublasSasum(handle, 128, d_data + 1, 1, d_data) 
11 cublasDestroy(handle); 
12 } 
13} 
14 
15void foo() { 
16 float *d_data; 
17 cudaMalloc((void **)&d_data, sizeof(float) * (1 + 128)); 
18 kernel<<<1, 128>>>(d_data); 
19 
20 float data; 
21 cudaMemcpy(data, d_data, sizeof(float), cudaMemcpyDeviceToHost); 
22 cudaFree(d_data); 
23}

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

1void kernel(float *d_data, sycl::nd_item<3> item_ct1) { 
2 int tid = item_ct1.get_local_id(2); 
3 d_data[tid + 1] = tid; 
4 
5 item_ct1.barrier(); 
6 
7 if (tid == 0) { 
8 /* 
9 DPCT1021:2: Migration of cublasHandle_t in __global__ or __device__ function 
10 is not supported.コードを再設計する必要があります11 */ 
12 cublasHandle_t handle; 
13 handle = &dpct::get_default_queue(); 
14 /* 
15 DPCT1020:1: Migration of cublasSasum, if it is called from __global__ or 
16 __device__ function, is not supported.You may need to redesign the code to 
17 use the host-side oneapi::mkl::blas::column_major::asum instead, which submits 
18 this call to the SYCL queue automatically.19 */ 
20 cublasSasum(handle, 128, d_data + 1, 1, d_data); 
21 handle = nullptr; 
22 } 
23} 
24 
25void foo() { 
26 dpct::device_ext &dev_ct1 = dpct::get_current_device(); 
27 sycl::queue &q_ct1 = dev_ct1.default_queue(); 
28 float *d_data; 
29 d_data = sycl::malloc_device<float>((1 + 128), q_ct1); 
30 q_ct1.parallel_for( 
31 sycl::nd_range<3>(sycl::range<3>(1, 1, 128), sycl::range<3>(1, 1, 128)), 
32 [=](sycl::nd_item<3> item_ct1) { 
33 kernel(d_data, item_ct1); 
34 }); 
35 
36 float data; 
37 q_ct1.memcpy(&data, d_data, sizeof(float)).wait(); 
38 sycl::free(d_data, q_ct1); 
39}

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

1void kernel(float *d_data, sycl::nd_item<3> item_ct1) { 
2 int tid = item_ct1.get_local_id(2); 
3 d_data[tid + 1] = tid; 
4} 
5 
6void foo() { 
7 dpct::device_ext &dev_ct1 = dpct::get_current_device(); 
8 sycl::queue &q_ct1 = dev_ct1.default_queue(); 
9 float *d_data; 
10 d_data = sycl::malloc_device<float>((1 + 128), q_ct1); 
11 q_ct1.parallel_for( 
12 sycl::nd_range<3>(sycl::range<3>(1, 1, 128), sycl::range<3>(1, 1, 128)), 
13 [=](sycl::nd_item<3> item_ct1) { 
14 kernel(d_data, item_ct1); 
15 }); 
16 oneapi::mkl::blas::column_major::asum(q_ct1, 128, d_data + 1, 1, d_data); 
17 
18 float data; 
19 q_ct1.memcpy(&data, d_data, sizeof(float)).wait(); 
20 sycl::free(d_data, q_ct1); 
21}