[Level Zero][Textures] no matching function for call to 'tex1Dfetch'

[pvelesko]

From LAMMPS:

/gpfs/jlse-fs0/users/pvelesko/lammps/build/lib/gpu/zbl.cu.cpp:115:17: error: no matching function for call to 'tex1Dfetch'
    numtyp4 ix; fetch4(ix,i,pos_tex); //x_[i];
                ^~~~~~~~~~~~~~~~~~~~
/gpfs/jlse-fs0/users/pvelesko/lammps/lib/gpu/lal_pre_cuda_hip.h:138:37: note: expanded from macro 'fetch4'
  #define fetch4(ans,i,pos_tex) ans=tex1Dfetch(pos_tex, i);
                                    ^~~~~~~~~~
/gpfs/jlse-fs0/users/pvelesko/install/HIP/clang13/chip-spv-1.2/include/hip/spirv_texture_functions.h:241:30: note: candidate template ignored: couldn't infer template argument 'T'
__TEXTURE_FUNCTIONS_DECL__ T tex1Dfetch(hipTextureObject_t TexObj, int X) {
                             ^
/gpfs/jlse-fs0/users/pvelesko/install/HIP/clang13/chip-spv-1.2/include/hip/spirv_texture_functions.h:168:16: note: candidate function not viable: requires 3 arguments, but 2 were provided
DEF_TEX1D_VEC4(tex1Dfetch, float4, int, _chip_tex1dfetchf);

[pvelesko]

I changed the #define in LAMMPS and now run into this:

/gpfs/jlse-fs0/users/pvelesko/install/HIP/clang13/chip-spv-1.2/include/hip/spirv_texture_functions.h:168:16: note: candidate function not viable: no known conversion from 'texture<float4>' (aka 'texture<HIP_vector_type<float, 4>>') to 'hipTextureObject_t' (aka '__hip_texture *') for 2nd argument
DEF_TEX1D_VEC4(tex1Dfetch, float4, int, _chip_tex1dfetchf);
               ^
/gpfs/jlse-fs0/users/pvelesko/install/HIP/clang13/chip-spv-1.2/include/hip/spirv_texture_functions.h:86:35: note: expanded from macro 'DEF_TEX1D_VEC4'
  __TEXTURE_FUNCTIONS_DECL__ void _NAME(                                       \
                                  ^
/gpfs/jlse-fs0/users/pvelesko/install/HIP/clang13/chip-spv-1.2/include/hip/spirv_texture_functions.h:138:15: note: candidate function not viable: no known conversion from 'float4 *' (aka 'HIP_vector_type<float, 4> *') to 'char *' for 1st argument
DEF_TEX1D_SCL(tex1Dfetch, char, int, _chip_tex1dfetchi);

[franz]

it's also breaking some CUDA samples, i'll take a look

[franz]

I've added some prototypes into the spirv_texture_functions.h but unfortunately, the issue goes much deeper, and AFAICT requires changes in LLVM passes (possibly a new pass) as well as the runtime library.

[linehill]

A deeper look into the issue seems to be more of a C++ problem.

note: candidate template ignored: couldn't infer template argument 'T'

The compiler tries to tell here it does not know what template instance it should generate. The function being called is declared as:

template <class T>
__TEXTURE_FUNCTIONS_DECL__ T tex1Dfetch(hipTextureObject_t TexObj, int X);

To instantiate function the template argument needs to be specified. For example:

texel = tex1Dfetch<float4>(tex, pos);

Or you can write something like:

texel = tex1Dfetch<std::remove_pointer<decltype(texel)>::type>(tex, pos);

[Kerilk]

Shouldn't the template specialization be deduced from the type of texel?

[franz]

@Kerilk apparently, return types are not considered by the compiler in type deduction (or overload resolution).

[Kerilk]

In regular C++ that's clearly the case (unless something changed recently, I have issue keeping up to date). In CUDA it seems the compiler is deducing from the return type. In HIP all examples I could find use explicit specialization, but maybe implicit is supposed to work as well?

[linehill]

If you mean nvcc that does not seem to be the case. See: https://godbolt.org/z/r9rM3bnzz. The foo call gives the following error:

error: no instance of function template "foo" matches the argument list

But if you change the call to foo<int>(...) it compiles.

[Kerilk]

No I was talking specifically of the tex1Dfetch intrinsic, not in general, sorry my post was a unclear.

[Kerilk]

But maybe it was when using texRef, where the type of the texRef could be used to specialize, whereas since switching to textureObject, you most probably need to be explicit about the specialization of tex1Dfetch

[pvelesko]

Additional missing functionality from HIP Unit tests:

// read from a texture using normalized coordinates
constexpr size_t ChannelToRead = 1;
template <typename T>
__global__ void readFromTexture(T* output, hipTextureObject_t texObj, size_t width, size_t height,
                                bool textureGather) {
  #if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
  // Calculate normalized texture coordinates
  const unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
  const unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
  const float u = x / (float)width;

  // Read from texture and write to global memory
  if (height == 0) {
    output[x] = tex1D<T>(texObj, u);
  } else {
    const float v = y / (float)height;
    if (textureGather) {
      // tex2Dgather not supported on __gfx90a__
      #if !defined(__gfx90a__) && !(defined(__HIP_PLATFORM_SPIRV__))
      output[y * width + x] = tex2Dgather<T>(texObj, u, v, ChannelToRead);
      #else
      #warning("tex2Dgather not supported on gfx90a");
      #endif
    } else {
      output[y * width + x] = tex2D<T>(texObj, u, v);
    }
  }
  #endif
}