DPCT1069#

メッセージ#

カーネル関数の <argument name> に、逆参照できない仮想ポインターが含まれています。usm-level=restricted でコードの移行を試してください。

詳細な説明#

インテル® DPC++ 互換性ツールは --usm-level=none オプションが使用されている場合、内部仮想ポインターを処理できません。

修正方法の提案#

--usm-level=none を指定せずにコードを再度移行します。

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

1  struct AAA { 
2   int *a; 
3  }; 
4 
5  __global__ void k(AAA obj) { 
6   obj.a[2] = 123; 
7  } 
8 
9  void foo() { 
10  AAA obj; 
11  int *a; 
12  cudaMalloc(&a, sizeof(int) * 10); 
13  obj.a = a; 
14  k<<<1, 1>>>(obj); 
15 }

このコードは、--usm-level=none を使用して移行され、以下の SYCL* コードが生成されます。

1  #define DPCT_USM_LEVEL_NONE 
2  #include <sycl/sycl.hpp> 
3  #include <dpct/dpct.hpp> 
4  struct AAA { 
5   int *a; 
6  }; 
7 
8  void k(AAA obj) { 
9   obj.a[2] = 123; 
10 } 
11 
12 void foo() { 
13  AAA obj; 
14  int *a; 
15  a = (int *)dpct::dpct_malloc(sizeof(int) * 10); 
16  obj.a = a; 
17  /* 
18  DPCT1069:0: The argument 'obj' of the kernel function contains virtual 
19  pointer(s), which cannot be dereferenced.Try to migrate the code with 
20  "usm-level=restricted".
21  */ 
22  dpct::get_out_of_order_queue().parallel_for( 
23    sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), 
24    [=](sycl::nd_item<3> item_ct1) { 
25    k(obj); 
26  }); 
27 }

--usm-level=none オプションなしで再度移行されたコードを以下に示します。

1  #include <sycl/sycl.hpp> 
2  #include <dpct/dpct.hpp> 
3  struct AAA { 
4   int *a; 
5  }; 
6 
7  void k(AAA obj) { 
8   obj.a[2] = 123; 
9  } 
10 
11 void foo() { 
12  sycl::device dev_ct1; 
13  sycl::queue q_ct1(dev_ct1, 
14  sycl::property_list{sycl::property::queue::in_order()}); 
15  AAA obj; 
16  int *a; 
17  a = sycl::malloc_device<int>(10, q_ct1); 18  obj.a = a; 
19  q_ct1.parallel_for( 
20    sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), 
21    [=](sycl::nd_item<3> item_ct1) { 
22    k(obj); 
23  }); 
24 }