DPCT1059#

Message#

SYCL only supports 4-channel image format. Adjust the code.

Detailed Help#

SYCL* supports only 4-channel image format. The warning is emitted, when the tool generates code with unsupported image format, which corresponds to the original code. You can fix the resulting code by changing the image format. Note: suggested workaround may impact code performance.

For example, this original CUDA* code:

1static texture<uint2, 2> tex;
2
3__global__ void test_image() {
4  uint2 tex_data = tex2D(tex, 0, 0);
5}
6void foo(uint2 *image_data) {
7  ...
8  test_image<<<1, 1>>>();
9}

results in the following migrated SYCL code:

 1/*
 2DPCT1059:0: SYCL only supports 4-channel image format. Adjust the code.
 3*/
 4dpct::image_wrapper<sycl::uint2, 2> tex;
 5
 6void test_image(dpct::image_accessor_ext<sycl::uint2, 2> tex) {
 7  sycl::uint2 tex_data = tex.read(0, 0);
 8}
 9void foo(sycl::uint2 *image_data) {
10  ...
11  dpct::get_default_queue().submit([&](sycl::handler &cgh) {
12    auto tex_acc = tex.get_access(cgh);
13
14    auto tex_smpl = tex.get_sampler();
15
16    cgh.parallel_for(
17        sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
18        [=](sycl::nd_item<3> item_ct1) {
19          test_image(
20              dpct::image_accessor_ext<sycl::uint2, 2>(tex_smpl, tex_acc));
21        });
22  });
23}

which is rewritten to:

 1dpct::image_wrapper<sycl::uint4, 2> tex;
 2
 3void test_image(dpct::image_accessor_ext<sycl::uint4, 2> tex) {
 4  sycl::uint4 tex_data = tex.read(0, 0);
 5}
 6void foo(sycl::uint2 *image_data) {
 7  // (1) Please allocate sycl::uint4 type new_image_data and initialize the value with image_data.
 8  // (2) Bind new_image_data to the tex object instead of the image_data.
 9  // (3) Free the memory of new_image_data after usage.
10  ...
11  dpct::get_default_queue().submit([&](sycl::handler &cgh) {
12    auto tex_acc = tex.get_access(cgh);
13
14    auto tex_smpl = tex.get_sampler();
15
16    cgh.parallel_for(
17        sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
18        [=](sycl::nd_item<3> item_ct1) {
19          test_image(
20              dpct::image_accessor_ext<sycl::uint4, 2>(tex_smpl, tex_acc));
21        });
22  });
23}

Suggestions to Fix#

You may need to rewrite this code.