DPCT1042

メッセージ

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

説明

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

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

以下の例では、同じ型の 2 つのバッファーをマージすることで、1 つのアクセサーを削除しています。

この例では、SYCL* カーネルのラムダでキャプチャーされる引数がサイズ制限を超えています。

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
 ... // declarations for device0 to device30
 dpct::constant_memory<int, 1> device31(ARRAY_SIZE);
 dpct::constant_memory<int, 1> device32(ARRAY_SIZE);

 // kernel function declaration
 void kernel(sycl::nd_item<3> item_ct1, int *device0, ..., int *device32);

 void test_function() {
   ... // memcpy operations for device0 to device30
   dpct::dpct_memcpy(device31.get_ptr(), host, ARRAY_SIZE * sizeof(int));
   dpct::dpct_memcpy(device32.get_ptr(), host, ARRAY_SIZE * sizeof(int));
   {
     dpct::get_default_queue().submit([&](sycl::handler &cgh) {
       ... // accessors for device0 to device30, captured by SYCL kernel lambda
           auto device31_acc_ct1 = device31.get_access(cgh);
       auto device32_acc_ct1 = device32.get_access(cgh);
       cgh.parallel_for(sycl::nd_range<3>(sycl::range<3>(1), sycl::range<3>(1)),
                        [=](sycl::nd_item<3> item_ct1) {
                          kernel(item_ct1, ... /*arguments device0 to device30*/
                                           device31_acc_ct1.get_pointer(),
                                 device32_acc_ct1.get_pointer());
                        });
     });
   }
 }

次のようにコードを更新することで、同じ型の 2 つのバッファー (device31 と device32) を 1 つ (device31) に統合し、アクセサーの数を 1 つ減らします。

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
 ... // declarations for device0 to device30
 dpct::constant_memory<int, 1> device32(ARRAY_SIZE + ARRAY_SIZE); // one buffer instead of two

 // kernel function declaration
 void kernel(sycl::nd_item<3> item_ct1, int *device0, ..., int *device32);

 void test_function() {
   ... // memcpy operations for device0 to device30
   dpct::dpct_memcpy(device31.get_ptr(), host, ARRAY_SIZE * sizeof(int));
   // memory copy with offset:
   dpct::dpct_memcpy(device31.get_ptr() + ARRAY_SIZE, host,
                     ARRAY_SIZE * sizeof(int));
   {
     dpct::get_default_queue().submit([&](sycl::handler &cgh) {
       ... // accessors for device0 to device30, captured by SYCL kernel lambda
           auto device31_acc_ct1 =
               device31.get_access(cgh); // only one accessor instead of 2
       cgh.parallel_for(sycl::nd_range<3>(sycl::range<3>(1), sycl::range<3>(1)),
                        [=](sycl::nd_item<3> item_ct1) {
                          // last parameter is modified to use offset and the
                          // same accessor as previous parameter
                          kernel(item_ct1, ... /*arguments device0 to device30*/
                                           device31_acc_ct1.get_pointer(),
                                 device31_acc_ct1.get_pointer() + ARRAY_SIZE);
                        });
     });
   }
 }

修正方法の提案

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