DPCT1084#

Message#

The function call <function name> has multiple migration results in different template instantiations that could not be unified. You may need to adjust the code.

Detailed Help#

SYCLomatic was unable to migrate the code correctly. Modify the code manually.

The following example shows original code, migrated code, and the manual changes made to correct the migrated code.

For example, this original CUDA* code:

 1  __constant__ int4 example_i[32];
 2  __constant__ float4 example_f[32];
 3
 4  struct example_int {
 5    __device__ int4 foo(int idx) const { return example_i[idx]; }
 6  };
 7
 8  struct example_float {
 9    __device__ float4 foo(int idx) const { return example_f[idx]; }
10  };
11
12  template <typename T> __global__ void example_kernel() {
13    T example_v;
14    int idx = blockIdx.x * blockDim.x + threadIdx.x;
15    float j = example_v.foo(idx).x;
16  }
17
18  void foo() {
19    example_kernel<example_int><<<1, 1>>>();
20    example_kernel<example_float><<<1, 1>>>();
21  }

results in the following migrated SYCL* code:

 1  static dpct::constant_memory<sycl::int4, 1> example_i(32);
 2  static dpct::constant_memory<sycl::float4, 1> example_f(32);
 3
 4  struct example_int {
 5    sycl::int4 foo(int idx, sycl::int4 const *example_i) const {
 6      return example_i[idx];
 7    }
 8  };
 9
10  struct example_float {
11    sycl::float4 foo(int idx, sycl::float4 const *example_f) const {
12      return example_f[idx];
13    }
14  };
15
16  template <typename T> void example_kernel(const sycl::nd_item<3> &item_ct1,
17                                            sycl::float4 const *example_f) {
18    T example_v;
19    int idx = item_ct1.get_group(2) * item_ct1.get_local_range(2) +
20              item_ct1.get_local_id(2);
21    /*
22    DPCT1084:0: The function call "example_int::foo" has multiple migration
23    results in different template instantiations that could not be unified. You
24    may need to adjust the code.
25    */
26    float j = example_v.foo(idx).x();
27  }
28
29  void foo() {
30    dpct::device_ext &dev_ct1 = dpct::get_current_device();
31    sycl::queue &q_ct1 = dev_ct1.in_order_queue();
32    q_ct1.submit([&](sycl::handler &cgh) {
33      example_f.init();
34
35      auto example_f_ptr_ct1 = example_f.get_ptr();
36
37      cgh.parallel_for(
38          sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
39          [=](sycl::nd_item<3> item_ct1) {
40            example_kernel<example_int>(item_ct1, example_f_ptr_ct1);
41          });
42    });
43    q_ct1.submit([&](sycl::handler &cgh) {
44      example_f.init();
45
46      auto example_f_ptr_ct1 = example_f.get_ptr();
47
48      cgh.parallel_for(
49          sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
50          [=](sycl::nd_item<3> item_ct1) {
51            example_kernel<example_float>(item_ct1, example_f_ptr_ct1);
52          });
53    });
54  }

which is manually adjusted to:

 1  static dpct::constant_memory<sycl::int4, 1> example_i(32);
 2  static dpct::constant_memory<sycl::float4, 1> example_f(32);
 3
 4  struct example_int {
 5    typedef sycl::int4 data_type;
 6    sycl::int4 foo(int idx, sycl::int4 const *example_i) const {
 7      return example_i[idx];
 8    }
 9  };
10
11  struct example_float {
12    typedef sycl::float4 data_type;
13    sycl::float4 foo(int idx, sycl::float4 const *example_f) const {
14      return example_f[idx];
15    }
16  };
17
18  template <typename T> void example_kernel(const sycl::nd_item<3> &item_ct1,
19                                            typename T::data_type const *example) {
20    T example_v;
21    int idx = item_ct1.get_group(2) * item_ct1.get_local_range(2) +
22              item_ct1.get_local_id(2);
23    float j = example_v.foo(idx, example).x();
24  }
25
26  void foo() {
27    dpct::device_ext &dev_ct1 = dpct::get_current_device();
28    sycl::queue &q_ct1 = dev_ct1.in_order_queue();
29    q_ct1.submit([&](sycl::handler &cgh) {
30      example_i.init();
31
32      auto example_i_ptr_ct1 = example_i.get_ptr();
33
34      cgh.parallel_for(
35          sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
36          [=](sycl::nd_item<3> item_ct1) {
37            example_kernel<example_int>(item_ct1, example_i_ptr_ct1);
38          });
39    });
40    q_ct1.submit([&](sycl::handler &cgh) {
41      example_f.init();
42
43      auto example_f_ptr_ct1 = example_f.get_ptr();
44
45      cgh.parallel_for(
46          sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
47          [=](sycl::nd_item<3> item_ct1) {
48            example_kernel<example_float>(item_ct1, example_f_ptr_ct1);
49          });
50    });
51  }

Suggestions to Fix#

Code requires manual adjustment.