DPCT1111#

メッセージ#

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

詳細な説明#

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

修正方法の提案#

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

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

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

1  void kernel(const sycl::nd_item<3> &item_ct1, int *b) { 
2 
3   item_ct1.barrier(sycl::access::fence_space::local_space); 
4  } 
5 
6  void 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 }

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

1  void kernel(const sycl::nd_item<3> &item_ct1, int *b) { 
2   item_ct1.barrier(sycl::access::fence_space::local_space); 
3  } 
4 
5  void 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 }