Missing tex2Dgather

[pvelesko]

From HIP unit tets HIP hipArrayCommon.hh

// 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
}

[pvelesko]

readFromTexture

[pvelesko]

tex1D

[pvelesko]

tex3D

[pvelesko]

tex2DLayered

[pvelesko]

tex3D

[pvelesko]

tex1Dfetch

[pjaaskel]

These might be a lot of work. Are these needed by apps of interest and is someone committed to implement by 1.0?

[linehill]

Dropped from Milestone 1.0: so far no demand for fixing gaps in the texture support.