DPCT1042

メッセージ

SYCL* カーネルに渡される引数のサイズが、カスタムではない SYCL* デバイスの最小サイズ制限 (1024) を超えています。ハードウェア引数のサイズの上限は、info::device::max_parameter_size で照会できます。引数のサイズがハードウェアの制限を超える場合、このコードを書き換える必要があります。

詳細な説明

カスタムでない SYCL* デバイスの SYCL* カーネルに渡される引数のサイズには制限があります (SYCL* 2020 仕様の「4.6.4.2 デバイス情報記述子」を参照)。

この警告が出力された場合は、SYCL* カーネルのラムダでキャプチャーされるアクセサーやその他の引数の数を減らすように、コードを手動で調整する必要があります。

次のセクションの例では、同じタイプの 2 つのバッファーをマージすることで、1 つのアクセサーを削除しています。

修正方法の提案

コードを確認して、調整してください。

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

  #define ARRAY_SIZE 2 

  __constant__ int device0[ARRAY_SIZE]; 
  __constant__ int device1[ARRAY_SIZE]; 
  __constant__ int device2[ARRAY_SIZE]; 
  ...
  __constant__ int device30[ARRAY_SIZE]; 
  __constant__ int device31[ARRAY_SIZE]; 

 // kernel function declaration 
 __global__ void kernel(int *out) { 
   int i = blockDim.x * blockIdx.x + threadIdx.x; 
   out[i] = device0[i] + device1[i] + device2[i] + ... + 
     device30[i] + device31[i]; 
 } 

 void test_function(int *out) { 
   kernel<<<1, 1>>>(out); 
 }

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

  #define ARRAY_SIZE 2 

  static dpct::constant_memory<int, 1> device0(ARRAY_SIZE); 
  static dpct::constant_memory<int, 1> device1(ARRAY_SIZE); 
  static dpct::constant_memory<int, 1> device2(ARRAY_SIZE); 
  ... 
  static dpct::constant_memory<int, 1> device30(ARRAY_SIZE); 
  static dpct::constant_memory<int, 1> device31(ARRAY_SIZE); 

 // kernel function declaration 
 void kernel(int *out, const sycl::nd_item<3> &item_ct1, int const *device0, 
 int const *device1, int const *device2, ..., int const *device30, 
 int const *device31) { 
   int i = item_ct1.get_local_range(2) * item_ct1.get_group(2) + 
   item_ct1.get_local_id(2); 
   out[i] = device0[i] + device1[i] + device2[i] + 
     ... + 
   device30[i] + device31[i]; 
 } 

 void test_function(int *out) { 
   device0.init(); 
   device1.init(); 
   device2.init(); 
   ... 
   device30.init(); 
   device31.init(); 

 /* 
 DPCT1042:0: The size of the arguments passed to the SYCL kernel exceeds the 
 minimum size limit (1024) for a non-custom SYCL device.You can get the 
 hardware argument size limit by querying info::device::max_parameter_size.You 
 may need to rewrite this code if the size of the arguments exceeds the 
 hardware limit.
 */ 
   dpct::get_out_of_order_queue().submit([&](sycl::handler &cgh) { 
     auto device0_acc_ct1 = device0.get_access(cgh); 
     auto device1_acc_ct1 = device1.get_access(cgh); 
     auto device2_acc_ct1 = device2.get_access(cgh); 
     ...   auto device30_acc_ct1 = device30.get_access(cgh); 
     auto device31_acc_ct1 = device31.get_access(cgh); 
     dpct::access_wrapper<int *> out_acc_ct0(out, cgh); 

     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(out_acc_ct0.get_raw_pointer(), item_ct1, 
         device0_acc_ct1.get_pointer(), device1_acc_ct1.get_pointer(), 
         device2_acc_ct1.get_pointer(), ..., 
         device30_acc_ct1.get_pointer(), device31_acc_ct1.get_pointer()); 
     }); 
   }); 
 }

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

  #define ARRAY_SIZE 2 

  static dpct::constant_memory<int, 1> device0(ARRAY_SIZE * 32); 

  // kernel function declaration 
  void kernel(int *out, const sycl::nd_item<3> &item_ct1, int const *device0) { 
  int i = item_ct1.get_local_range(2) * item_ct1.get_group(2) + 
  item_ct1.get_local_id(2); 
    for (int j = 0; j < 32; j++) { 
     out[i] += device0[ARRAY_SIZE * j + i]; 
   } 
 } 

 void test_function(int *out) { 
   device0.init(); 

   dpct::get_out_of_order_queue().submit([&](sycl::handler &cgh) { 
     auto device0_acc_ct1 = device0.get_access(cgh); 
     dpct::access_wrapper<int *> out_acc_ct0(out, cgh); 

     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(out_acc_ct0.get_raw_pointer(), item_ct1, 
       device0_acc_ct1.get_pointer()); 
     }); 
   }); 
 }