DPCT1062#

Message#

SYCL Image doesn’t support normalized read mode.

Detailed Help#

This warning is emitted when cudaReadModeNormalizedFloat is used as the third argument of texture in the original code. Since SYCL* Image doesn’t support normalized read mode, cudaReadModeNormalizedFloat will be ignored during migration.

It may cause errors in the resulting code, for example, redefinition of overloaded functions, if the overloaded functions are differentiated based on the texture type in the original code.

For example, this original CUDA* code:

1__device__ void foo(const texture<char4, 2, cudaReadModeNormalizedFloat> tex) {
2  float4 f = tex2D(tex, 0.5f, 0.5f);
3}
4__device__ void foo (const texture<char4, 2, cudaReadModeElementType> tex) {
5  char4 c = tex2D(tex, 0.5f, 0.5f);
6}

results in the following migrated SYCL code:

1/*
2DPCT1062:0: SYCL Image doesn't support normalized read mode.
3*/
4void foo(dpct::image_accessor_ext<sycl::char4, 2> tex) {
5  sycl::float4 f = tex.read(0.5f, 0.5f);
6}
7void foo(dpct::image_accessor_ext<sycl::char4, 2> tex) {
8  sycl::char4 c = tex.read(0.5f, 0.5f);
9}

which is rewritten to:

1void foo1(dpct::image_accessor_ext<sycl::char4, 2> tex) {
2  sycl::char4 temp_c = tex.read(0.5f, 0.5f);
3  sycl::float4 f{temp_c.x(), temp_c.y(), temp_c.z(), temp_c.w()};
4  f = f / 255.f;
5}
6void foo2(dpct::image_accessor_ext<sycl::char4, 2> tex) {
7  sycl::char4 c = tex.read(0.5f, 0.5f);
8}

Suggestions to Fix#

Review the code and update as needed.