DPCT1069

メッセージ

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

説明

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

修正方法の提案#

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

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


1struct AAA { 
2    int *a; 
3}; 
4 
5__global__ void k(AAA obj) { 
6    obj.a[2] = 123; 
7} 
8 
9void 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> 
4struct AAA { 
5    int *a; 
6}; 
7 
8void k(AAA obj) { 
9    obj.a[2] = 123; 
10} 
11 
12void 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> 
3struct AAA { 
4    int *a; 
5}; 
6 
7void k(AAA obj) { 
8    obj.a[2] = 123; 
9} 
10 
11void 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}