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 SYCLomatic because the variable of this type was not used directly in the code. SYCLomatic 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 }