DPCT1050

Intel® DPC++ Compatibility Tool Developer Guide and Reference

Download PDF

ID 768918

Date 10/31/2024

Version

Public

Visible to Intel only — GUID: GUID-15D09574-7A47-473A-B1C1-4D385239D604

View Details

DPCT1050

Message

The template argument of the <type> could not be deduced. You need to update this code.

Detailed Help

This warning is generated when the template argument could not be deduced by the Intel® DPC++ Compatibility Tool because the variable of this type was not used directly in the code. Intel® DPC++ Compatibility Tool inserts “dpct_placeholder”, instead of type, in such cases.

Suggestions to Fix

Replace the “dpct_placeholder” with the real argument.

For example, this original CUDA* code:

 __global__ void kernel(const cudaTextureObject_t texObj) {}

 void foo() {
   float4 *d_data42;
   cudaArray_t a42;
   cudaMalloc(&d_data42, sizeof(float4) * 32 * 32);
   cudaChannelFormatDesc desc42 = cudaCreateChannelDesc(32, 32, 32, 32, cudaChannelFormatKindFloat);
   cudaMallocArray(&a42, &desc42, 32, 32);
   cudaMemcpyToArray(a42, 0, 0, d_data42, 32 * 32 * sizeof(float4), cudaMemcpyDeviceToDevice);
   cudaTextureObject_t tex42;
   cudaResourceDesc res42;
   cudaTextureDesc texDesc42;
   res42.resType = cudaResourceTypeArray;
   res42.res.array.array = a42;
   cudaCreateTextureObject(&tex42, &res42, &texDesc42, NULL);
   kernel<<<1, 1>>>(tex42);
 }

results in the following migrated SYCL* code:

 /*
 DPCT1050:1: The template argument of the image_accessor_ext could not be
 deduced. You need to update this code.
 */
 void kernel(const dpct::image_accessor_ext<
             dpct_placeholder /*Fix the type manually*/, 1>
                 texObj) {}

 void foo() {
   dpct::device_ext &dev_ct1 = dpct::get_current_device();
   sycl::queue &q_ct1 = dev_ct1.default_queue();
   sycl::float4 *d_data42;
   dpct::image_matrix_p a42;
   d_data42 = (sycl::float4 *)sycl::malloc_device(sizeof(sycl::float4) * 32 * 32,
                                                  q_ct1);
   dpct::image_channel desc42 =
       dpct::image_channel(32, 32, 32, 32, dpct::image_channel_data_type::fp);
   a42 = new dpct::image_matrix(desc42, sycl::range<2>(32, 32));
   dpct::dpct_memcpy(a42->to_pitched_data(), sycl::id<3>(0, 0, 0),
                     dpct::pitched_data(d_data42, 32 * 32 * sizeof(sycl::float4),
                                        32 * 32 * sizeof(sycl::float4), 1),
                     sycl::id<3>(0, 0, 0),
                     sycl::range<3>(32 * 32 * sizeof(sycl::float4), 1, 1));
   dpct::image_wrapper_base_p tex42;
   dpct::image_data res42;
   dpct::sampling_info texDesc42;

   res42.set_data(a42);
   tex42 = dpct::create_image_wrapper(res42, texDesc42);
   /*
   DPCT1050:0: The template argument of the image_accessor_ext could not be
   deduced. You need to update this code.
   */
   q_ct1.submit([&](sycl::handler &cgh) {
     auto tex42_acc = static_cast<dpct::image_wrapper<
         dpct_placeholder /*Fix the type manually*/, 1> *>(tex42)
                          ->get_access(cgh);

     auto tex42_smpl = tex42->get_sampler();

     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(dpct::image_accessor_ext<
                  dpct_placeholder /*Fix the type manually*/, 1>(tex42_smpl,
                                                                 tex42_acc));
         });
   });
 }

which is rewritten to:

 void kernel(const dpct::image_accessor_ext<sycl::float4, 2> texObj) {}

 void foo() {
   dpct::device_ext &dev_ct1 = dpct::get_current_device();
   sycl::queue &q_ct1 = dev_ct1.default_queue();
   sycl::float4 *d_data42;
   dpct::image_matrix_p a42;
   d_data42 = (sycl::float4 *)sycl::malloc_device(sizeof(sycl::float4) * 32 * 32,
                                                  q_ct1);
   dpct::image_channel desc42 =
       dpct::image_channel(32, 32, 32, 32, dpct::image_channel_data_type::fp);
   a42 = new dpct::image_matrix(desc42, sycl::range<2>(32, 32));
   dpct::dpct_memcpy(a42->to_pitched_data(), sycl::id<3>(0, 0, 0),
                     dpct::pitched_data(d_data42, 32 * 32 * sizeof(sycl::float4),
                                        32 * 32 * sizeof(sycl::float4), 1),
                     sycl::id<3>(0, 0, 0),
                     sycl::range<3>(32 * 32 * sizeof(sycl::float4), 1, 1));
   dpct::image_wrapper_base_p tex42;
   dpct::image_data res42;
   dpct::sampling_info texDesc42;

   res42.set_data(a42);
   tex42 = dpct::create_image_wrapper(res42, texDesc42);

   q_ct1.submit([&](sycl::handler &cgh) {
     auto tex42_acc =
         static_cast<dpct::image_wrapper<sycl::float4, 2> *>(tex42)->get_access(
             cgh);

     auto tex42_smpl = tex42->get_sampler();

     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(
               dpct::image_accessor_ext<sycl::float4, 2>(tex42_smpl, tex42_acc));
         });
   });
 }

Select Your Language

Using Intel.com Search

Quick Links

Recent Searches

Advanced Search

Only search in

Intel® DPC++ Compatibility Tool Developer Guide and Reference

DPCT1050

Message

Detailed Help

Suggestions to Fix