DPCT1111#

メッセージ#

<migrated API> の入力引数をターゲット関数 <kernel function name> に基づいて確認します。

説明#

インテル® DPC++ 互換性ツールは、関数ポインターまたは Cufunction 変数から一部の引数を推測できません。引数値を手動で確認する必要があります。

修正方法の提案#

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

1extern __shared__ float a[]; 
2__global__ void kernel() { 
3 __shared__ int b[10]; 
4 __syncthreads(); 
5} 
6 
7void foo() { 
8 int numBlocks; 
9 cudaOccupancyMaxActiveBlocksPerMultiprocessor(&numBlocks, kernel, 128, 
10 sizeof(float) * 20); 
11 kernel<<<1, 128>>>(); 
12}

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

1void kernel(const sycl::nd_item<3> &item_ct1, int *b) { 
2 
3 item_ct1.barrier(sycl::access::fence_space::local_space); 
4} 
5 
6void foo() { 
7 int numBlocks; 
8 /* 
9 DPCT1111:0: Please verify the input arguments of 
10 dpct::experimental::calculate_max_active_wg_per_xecore base on the target 
11 function "kernel".12 */ 
13 dpct::experimental::calculate_max_active_wg_per_xecore( 
14 &numBlocks, 128, 
15 sizeof(float) * 20 + 
16 dpct_placeholder /* total shared local memory size */); 
17 dpct::get_default_queue().submit([&](sycl::handler &cgh) { 
18 sycl::local_accessor<int, 1> b_acc_ct1(sycl::range<1>(10), cgh); 
19 
20 cgh.parallel_for( 
21 sycl::nd_range<3>(sycl::range<3>(1, 1, 128), sycl::range<3>(1, 1, 128)), 
22 [=](sycl::nd_item<3> item_ct1) { 
23 kernel(item_ct1, b_acc_ct1.get_pointer()); 
24 }); 
25 }); 
26}

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

1void kernel(const sycl::nd_item<3> &item_ct1, int *b) { 
2 item_ct1.barrier(sycl::access::fence_space::local_space); 
3} 
4 
5void foo() { 
6 int numBlocks; 
7 dpct::experimental::calculate_max_active_wg_per_xecore( 
8 &numBlocks, 128, sizeof(float) * 20 + sizeof(int) * 10, 32, true, false); 
9 dpct::get_default_queue().submit([&](sycl::handler &cgh) { 
10 sycl::local_accessor<int, 1> b_acc_ct1(sycl::range<1>(10), cgh); 
11 
12 cgh.parallel_for( 
13 sycl::nd_range<3>(sycl::range<3>(1, 1, 128), sycl::range<3>(1, 1, 128)), 
14 [=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(32)]] { 
15 kernel(item_ct1, b_acc_ct1.get_pointer()); 
16 }); 
17 }); 
18}