DPCT1120

メッセージ

<変数名> はデバイスメモリー内にあり、追加の引数としてデバイス関数に渡されます。ここで <変数名> はテンプレート変数であるため、デバイス関数内の追加パラメーターの宣言、コマンドグループで参照される変数、およびデバイス関数が正しくない可能性があります。ソースコードの調整が必要な場合があります。

詳細な説明

<変数名> はデバイスメモリー内にあり、追加の引数としてデバイス関数に渡されます。ここで <変数名> はテンプレート変数であるため、デバイス関数内の追加パラメーターの宣言、コマンドグループで参照される変数、およびデバイス関数もテンプレート変数である必要があり、移行結果が正しくない可能性があります。ソースコードの調整が必要な場合があります。

修正方法の提案

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

  struct Foo { 
   int x; 
   __host__ __device__ Foo() {} 
   __host__ __device__ ~Foo() {} 
  }; 
 
  template <typename T> __constant__ T cmem; 
 
  __global__ void kernel() { 
  ...
  int x = cmem<Foo>.x; 
  ... 
 } 
 
 void foo() { 
  ... 
  kernel<<<1, 1>>>(); 
  ... 
}

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

 struct Foo { 
  int x; 
  Foo() {} 
  ~Foo() {} 
 }; 

 /* 
 DPCT1120:0: cmem is in device memory and passed into device functions as an 
 extra arguments.As cmem is a template variable here, then the declaration of 
an extra parameter in device functions, the variable referenced in command 
group, and device functions may be incorrect as a result.You may need to 
adjust the code.
*/ 
template <typename T> static dpct::constant_memory<T, 0> cmem; 

void kernel(T cmem) { 
 ... 
 int x = cmem<Foo>.x; 
 ... 
} 

void foo() { 
 ... 
 cmem.init(); 

 dpct::get_in_order_queue().submit([&](sycl::handler &cgh) { 
 auto cmem_ptr_ct1 = cmem.get_ptr(); 

 cgh.parallel_for( 
   sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), 
   [=](sycl::nd_item<3> item_ct1) { 
   kernel(*cmem_ptr_ct1); 
   }); 
 }); 
 ... 
}

上記は次のように書き換える必要があります。

 struct Foo { 
  int x; 
  Foo() {} 
  ~Foo() {} 
 }; 

 template <typename T> static dpct::constant_memory<T, 0> cmem; 

 template <class T> 
void kernel(T cmem) { 
 ... 
 int x = cmem.x; 
 ... 
} 

void foo() { 
 ... 
 cmem<Foo>.init(); 

 dpct::get_in_order_queue().submit([&](sycl::handler &cgh) { 
 auto cmem_ptr_ct1 = cmem<Foo>.get_ptr(); 

 cgh.parallel_for( 
   sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), 
   [=](sycl::nd_item<3> item_ct1) { 
   kernel(*cmem_ptr_ct1); 
   }); 
 }); 
 ... 
}