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);
});
});
}
}
|
修正方法の提案¶
コードを確認して、調整してください。