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:

 1 __global__ void kernel(const cudaTextureObject_t texObj) {}
 2
 3 void foo() {
 4   float4 *d_data42;
 5   cudaArray_t a42;
 6   cudaMalloc(&d_data42, sizeof(float4) * 32 * 32);
 7   cudaChannelFormatDesc desc42 = cudaCreateChannelDesc(32, 32, 32, 32, cudaChannelFormatKindFloat);
 8   cudaMallocArray(&a42, &desc42, 32, 32);
 9   cudaMemcpyToArray(a42, 0, 0, d_data42, 32 * 32 * sizeof(float4), cudaMemcpyDeviceToDevice);
10   cudaTextureObject_t tex42;
11   cudaResourceDesc res42;
12   cudaTextureDesc texDesc42;
13   res42.resType = cudaResourceTypeArray;
14   res42.res.array.array = a42;
15   cudaCreateTextureObject(&tex42, &res42, &texDesc42, NULL);
16   kernel<<<1, 1>>>(tex42);
17 }

results in the following migrated SYCL* code:

 1 /*
 2 DPCT1050:1: The template argument of the image_accessor_ext could not be
 3 deduced. You need to update this code.
 4 */
 5 void kernel(const dpct::image_accessor_ext<
 6             dpct_placeholder /*Fix the type manually*/, 1>
 7                 texObj) {}
 8
 9 void foo() {
10   dpct::device_ext &dev_ct1 = dpct::get_current_device();
11   sycl::queue &q_ct1 = dev_ct1.default_queue();
12   sycl::float4 *d_data42;
13   dpct::image_matrix_p a42;
14   d_data42 = (sycl::float4 *)sycl::malloc_device(sizeof(sycl::float4) * 32 * 32,
15                                                  q_ct1);
16   dpct::image_channel desc42 =
17       dpct::image_channel(32, 32, 32, 32, dpct::image_channel_data_type::fp);
18   a42 = new dpct::image_matrix(desc42, sycl::range<2>(32, 32));
19   dpct::dpct_memcpy(a42->to_pitched_data(), sycl::id<3>(0, 0, 0),
20                     dpct::pitched_data(d_data42, 32 * 32 * sizeof(sycl::float4),
21                                        32 * 32 * sizeof(sycl::float4), 1),
22                     sycl::id<3>(0, 0, 0),
23                     sycl::range<3>(32 * 32 * sizeof(sycl::float4), 1, 1));
24   dpct::image_wrapper_base_p tex42;
25   dpct::image_data res42;
26   dpct::sampling_info texDesc42;
27
28   res42.set_data(a42);
29   tex42 = dpct::create_image_wrapper(res42, texDesc42);
30   /*
31   DPCT1050:0: The template argument of the image_accessor_ext could not be
32   deduced. You need to update this code.
33   */
34   q_ct1.submit([&](sycl::handler &cgh) {
35     auto tex42_acc = static_cast<dpct::image_wrapper<
36         dpct_placeholder /*Fix the type manually*/, 1> *>(tex42)
37                          ->get_access(cgh);
38
39     auto tex42_smpl = tex42->get_sampler();
40
41     cgh.parallel_for(
42         sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
43         [=](sycl::nd_item<3> item_ct1) {
44           kernel(dpct::image_accessor_ext<
45                  dpct_placeholder /*Fix the type manually*/, 1>(tex42_smpl,
46                                                                 tex42_acc));
47         });
48   });
49 }

which is rewritten to:

 1 void kernel(const dpct::image_accessor_ext<sycl::float4, 2> texObj) {}
 2
 3 void foo() {
 4   dpct::device_ext &dev_ct1 = dpct::get_current_device();
 5   sycl::queue &q_ct1 = dev_ct1.default_queue();
 6   sycl::float4 *d_data42;
 7   dpct::image_matrix_p a42;
 8   d_data42 = (sycl::float4 *)sycl::malloc_device(sizeof(sycl::float4) * 32 * 32,
 9                                                  q_ct1);
10   dpct::image_channel desc42 =
11       dpct::image_channel(32, 32, 32, 32, dpct::image_channel_data_type::fp);
12   a42 = new dpct::image_matrix(desc42, sycl::range<2>(32, 32));
13   dpct::dpct_memcpy(a42->to_pitched_data(), sycl::id<3>(0, 0, 0),
14                     dpct::pitched_data(d_data42, 32 * 32 * sizeof(sycl::float4),
15                                        32 * 32 * sizeof(sycl::float4), 1),
16                     sycl::id<3>(0, 0, 0),
17                     sycl::range<3>(32 * 32 * sizeof(sycl::float4), 1, 1));
18   dpct::image_wrapper_base_p tex42;
19   dpct::image_data res42;
20   dpct::sampling_info texDesc42;
21
22   res42.set_data(a42);
23   tex42 = dpct::create_image_wrapper(res42, texDesc42);
24
25   q_ct1.submit([&](sycl::handler &cgh) {
26     auto tex42_acc =
27         static_cast<dpct::image_wrapper<sycl::float4, 2> *>(tex42)->get_access(
28             cgh);
29
30     auto tex42_smpl = tex42->get_sampler();
31
32     cgh.parallel_for(
33         sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
34         [=](sycl::nd_item<3> item_ct1) {
35           kernel(
36               dpct::image_accessor_ext<sycl::float4, 2>(tex42_smpl, tex42_acc));
37         });
38   });
39 }