DPCT1042

Intel® DPC++ Compatibility Tool Developer Guide and Reference

Download PDF

ID 768918

Date 3/22/2024

Version

Public

A newer version of this document is available. Customers should click here to go to the newest version.

Visible to Intel only — GUID: GUID-91DFA275-0F17-436F-928A-438DE05F50D0

View Details

DPCT1042

Message

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.

Detailed Help

The size of the arguments passed to the SYCL* kernel for non-custom SYCL device has a limit (see SYCL 2020 standard, 4.6.4.2 Device information descriptors).

In cases where this warning occurs, you need to adjust the code manually to decrease the number of accessors or other arguments that are captured by the SYCL kernel lambda.

The example in the next section shows how you can remove one accessor by merging two buffers with the same type.

Suggestions to Fix

Review the code and adjust it.

For example, this original CUDA* code:

 #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);
 }

results in the following migrated SYCL code:

 #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());
         });
   });
 }

which is rewritten to:

 #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());
         });
   });
 }

Select Your Language

Using Intel.com Search

Quick Links

Recent Searches

Advanced Search

Only search in

Intel® DPC++ Compatibility Tool Developer Guide and Reference

DPCT1042

Message

Detailed Help

Suggestions to Fix