SWDEV-306947 - Add mode combination tests (#2447)

Address modes are valid for array rather than linear
memory. Comment out codes of address test for linear
memory.
Add tests of mode combination(address, linear and
normorlized coordinates) in 1D, 2D and 3D.

Change-Id: I9e7e97d9bf619743c311f68f51d5bf115fdd534d

[ROCm/hip commit: e05f98967e]

projects/hip/tests/src/texture/hipBindTexRef1DFetch.cpp

@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 
 /* HIT_START
- * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_RUNTIME rocclr
+ * BUILD: %t %s ../test_common.cpp
  * TEST: %t
  * HIT_END
  */

projects/hip/tests/src/texture/hipNormalizedFloatValueTex.cpp

@@ -32,8 +32,7 @@ THE SOFTWARE.
 #include "test_common.h"
 #include <math.h>
 #define SIZE 10
-#define EPSILON     0.00001
-#define THRESH_HOLD 0.01  // For filter mode
+#include "hipTextureHelper.hpp"
 
 static float getNormalizedValue(const float value,
                                 const hipChannelFormatDesc& desc) {
@@ -104,8 +103,8 @@ bool textureVerifyFilterModeLinear(float *hOutputData, float *expected,  size_t
     bool testResult = true;
     for (int i = 0; i < size; i++) {
       float mean = (fabs(expected[i]) + fabs(hOutputData[i])) / 2;
-      float ratio = fabs(expected[i] - hOutputData[i]) / (mean + EPSILON);
-      if (ratio > THRESH_HOLD) {
+      float ratio = fabs(expected[i] - hOutputData[i]) / (mean + HIP_SAMPLING_VERIFY_EPSILON);
+      if (ratio > HIP_SAMPLING_VERIFY_RELATIVE_THRESHOLD) {
         printf("mismatch at output[%d]:%f expected[%d]:%f, ratio:%f\n", i,
                hOutputData[i], i, expected[i], ratio);
         testResult = false;
@@ -131,7 +130,7 @@ bool textureTest(texture<T, hipTextureType1D, hipReadModeNormalizedFloat> *tex)
 {
     hipChannelFormatDesc desc = hipCreateChannelDesc<T>();
     hipArray_t dData;
-    HIPCHECK(hipMallocArray(&dData, &desc, SIZE, 1, hipArrayDefault));
+    HIPCHECK(hipMallocArray(&dData, &desc, SIZE));
 
     T hData[] = {65, 66, 67, 68, 69, 70, 71, 72, 73, 74};
     HIPCHECK(hipMemcpy2DToArray(dData, 0, 0, hData, sizeof(T)*SIZE, sizeof(T)*SIZE, 1, hipMemcpyHostToDevice));
@@ -196,7 +195,8 @@ int main(int argc, char** argv)
       status = runTest<hipFilterModePoint>();
     } else if(textureFilterMode == 1) {
       printf("Test hipFilterModeLinear\n");
-      printf("THRESH_HOLD:%f, EPSILON:%f\n", THRESH_HOLD, EPSILON);
+      printf("THRESH_HOLD:%f, EPSILON:%f\n", HIP_SAMPLING_VERIFY_RELATIVE_THRESHOLD,
+             HIP_SAMPLING_VERIFY_EPSILON);
       status = runTest<hipFilterModeLinear>();
     } else {
       printf("Wrong argument!\n");

projects/hip/tests/src/texture/hipTex1DFetchCheckModes.cpp

@@ -18,7 +18,7 @@ THE SOFTWARE.
 */
 
 /* HIT_START
- * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_RUNTIME rocclr
+ * BUILD: %t %s ../test_common.cpp
  * TEST: %t
  * HIT_END
  */
@@ -38,9 +38,9 @@ int runTest(hipTextureAddressMode, hipTextureFilterMode);
 
 int main(int argc, char **argv) {
     int testResult = runTest(hipAddressModeClamp,hipFilterModePoint);
-    testResult = runTest(hipAddressModeClamp,hipFilterModeLinear);
-    testResult = runTest(hipAddressModeWrap,hipFilterModePoint);
-    testResult = runTest(hipAddressModeWrap,hipFilterModeLinear);
+    testResult = testResult & runTest(hipAddressModeClamp,hipFilterModeLinear);
+    testResult = testResult & runTest(hipAddressModeBorder,hipFilterModePoint);
+    testResult = testResult & runTest(hipAddressModeBorder,hipFilterModeLinear);
     if(testResult) {
         passed();
     } else {
@@ -84,7 +84,6 @@ int runTest(hipTextureAddressMode addressMode, hipTextureFilterMode filterMode)
     texDesc.readMode = hipReadModeElementType;
 
     texDesc.addressMode[0] = addressMode;
-    texDesc.addressMode[1] = addressMode;
     texDesc.filterMode = filterMode;
     texDesc.normalizedCoords = false;
 
@@ -107,16 +106,12 @@ int runTest(hipTextureAddressMode addressMode, hipTextureFilterMode filterMode)
             break;
         }
     }
-    if(testResult){
-        for(int i = N-offset; i < N; i++){
-           if (output[i] != 0){
-               testResult = 0;
-               break;
-           }
-        }
-    }
+    // For hipResourceTypeLinear, reading of out-of-boundary address is undefined!
+    // So we won't verify those data
+
     HIPCHECK(hipDestroyTextureObject(texObj));
     HIPCHECK(hipFree(texBuf));
     HIPCHECK(hipFree(texBufOut));
+    printf("%s(addressMode %d, filterMode %d) %s\n", __FUNCTION__, addressMode, filterMode, testResult ? "succeed" : "failed");
     return testResult;
 }

projects/hip/tests/src/texture/hipTextureHelper.hpp

@@ -0,0 +1,227 @@
+#pragma once
+
+#define HIP_SAMPLING_VERIFY_EPSILON     0.00001
+// The internal precision varies by the GPU family and sometimes within the family.
+// Thus the following threshold is subject to change.
+#define HIP_SAMPLING_VERIFY_RELATIVE_THRESHOLD  0.05  // 5% for filter mode
+#define HIP_SAMPLING_VERIFY_ABSOLUTE_THRESHOLD  0.1
+
+template<typename type, hipTextureFilterMode fMode = hipFilterModePoint>
+bool hipTextureSamplingVerify(const type outputData, const type expected) {
+  bool testResult = false;
+  if (fMode == hipFilterModePoint) {
+    testResult = outputData == expected;
+  } else if (fMode == hipFilterModeLinear) {
+    const type mean = (fabs(outputData) + fabs(expected)) / 2;
+    const type diff = fabs(outputData - expected);
+    const type ratio = diff / (mean + HIP_SAMPLING_VERIFY_EPSILON);
+    if (ratio <= HIP_SAMPLING_VERIFY_RELATIVE_THRESHOLD) {
+      testResult = true;
+    } else if (diff <= HIP_SAMPLING_VERIFY_ABSOLUTE_THRESHOLD) {
+      // Some small outputs have big ratio due to float operation difference of ALU and GPU
+      testResult = true;
+    }
+  }
+  return testResult;
+}
+
+// Simulate CTS static AddressingTable sAddressingTable
+template<hipTextureAddressMode addressMode>
+void hipTextureGetAddress(int &value, const int maxValue)
+{
+  switch(addressMode)
+  {
+    case hipAddressModeClamp:
+      value = value < 0 ? 0
+                  : (value > maxValue - 1 ? maxValue - 1 : value);
+      break;
+    case hipAddressModeBorder:
+      value = value < -1 ? -1
+                  : (value > maxValue ? maxValue : value);
+      break;
+    default:
+      break;
+  }
+}
+
+// Simlate logics in CTS read_image_pixel_float().
+// x, y and z must be returned by hipTextureGetAddress()
+template<hipTextureAddressMode addressMode>
+float hipTextureGetValue(const float *data, const int x, const int width,
+         const int y = 0, const int height = 0,const int z = 0, const int depth = 0) {
+  float result = std::numeric_limits<float>::lowest();
+  switch (addressMode) {
+    case hipAddressModeClamp:
+      if (width > 0) {
+        if (height == 0 && depth == 0) {
+          result = data[x];  // 1D
+        } else if (depth == 0) {
+          result = data[y * width + x];  // 2D
+        } else {
+          result = data[z * width * height + y * width + x];  // 3D
+        }
+      }
+      break;
+    case hipAddressModeBorder:
+      if (width > 0) {
+        if (height == 0 && depth == 0) {
+          result = (x >= 0 && x < width) ? data[x] : 0;  // 1D
+        } else if (depth == 0) {
+          result = (x >= 0 && x < width && y >= 0 && y < height) ?
+              data[y * width + x] : 0;  // 2D
+        } else {
+          result = (x >= 0 && x < width && y >= 0 && y < height && z >= 0 && z < depth) ?
+              data[z * width * height + y * width + x] : 0;  // 3D
+        }
+      }
+      break;
+    default:
+      break;
+  }
+  return result;
+}
+
+template<hipTextureAddressMode addressMode, hipTextureFilterMode filterMode>
+float getExpectedValue(const int width, float x, const float *data) {
+  float result = std::numeric_limits<float>::lowest();
+  switch (filterMode) {
+    case hipFilterModePoint: {
+      int i1 = static_cast<int>(floor(x));
+      hipTextureGetAddress < addressMode > (i1, width);
+      result = hipTextureGetValue < addressMode > (data, i1, width);
+    }
+      break;
+    case hipFilterModeLinear: {
+      x -= 0.5;
+      int i1 = static_cast<int>(floor(x));
+      int i2 = i1 + 1;
+      float a = x - i1;
+      hipTextureGetAddress < addressMode > (i1, width);
+      hipTextureGetAddress < addressMode > (i2, width);
+
+      float t1 = hipTextureGetValue < addressMode > (data, i1, width);
+      float t2 = hipTextureGetValue < addressMode > (data, i2, width);
+
+      return (1 - a) * t1 + a * t2;
+    }
+      break;
+  }
+  return result;
+}
+
+template<hipTextureAddressMode addressMode, hipTextureFilterMode filterMode>
+float getExpectedValue(const int width, const int height, float x, float y, const float *data) {
+  float result = std::numeric_limits<float>::lowest();
+  switch (filterMode) {
+    case hipFilterModePoint: {
+      int i1 = static_cast<int>(floor(x));
+      int j1 = static_cast<int>(floor(y));
+      hipTextureGetAddress < addressMode > (i1, width);
+      hipTextureGetAddress < addressMode > (j1, height);
+      result = hipTextureGetValue < addressMode > (data, i1, width, j1, height);
+    }
+      break;
+    case hipFilterModeLinear: {
+      x -= 0.5;
+      y -= 0.5;
+
+      int i1 = static_cast<int>(floor(x));
+      int j1 = static_cast<int>(floor(y));
+
+      int i2 = i1 + 1;
+      int j2 = j1 + 1;
+
+      float a = x - i1;
+      float b = y - j1;
+
+      hipTextureGetAddress < addressMode > (i1, width);
+      hipTextureGetAddress < addressMode > (i2, width);
+      hipTextureGetAddress < addressMode > (j1, height);
+      hipTextureGetAddress < addressMode > (j2, height);
+
+      float t11 = hipTextureGetValue < addressMode
+          > (data, i1, width, j1, height);
+      float t21 = hipTextureGetValue < addressMode
+          > (data, i2, width, j1, height);
+      float t12 = hipTextureGetValue < addressMode
+          > (data, i1, width, j2, height);
+      float t22 = hipTextureGetValue < addressMode
+          > (data, i2, width, j2, height);
+
+      result = (1 - a) * (1 - b) * t11 + a * (1 - b) * t21 + (1 - a) * b * t12
+          + a * b * t22;
+    }
+      break;
+  }
+  return result;
+}
+
+template<hipTextureAddressMode addressMode, hipTextureFilterMode filterMode>
+float getExpectedValue(const int width, const int height, const int depth,
+                      float x, float y, float z, const float *data) {
+  float result = std::numeric_limits<float>::lowest();
+  switch (filterMode) {
+    case hipFilterModePoint: {
+      int i1 = static_cast<int>(floor(x));
+      int j1 = static_cast<int>(floor(y));
+      int k1 = static_cast<int>(floor(z));
+
+      hipTextureGetAddress < addressMode > (i1, width);
+      hipTextureGetAddress < addressMode > (j1, height);
+      hipTextureGetAddress < addressMode > (k1, depth);
+
+      result = hipTextureGetValue < addressMode > (data, i1, width, j1, height, k1, depth);
+    }
+      break;
+    case hipFilterModeLinear: {
+      x -= 0.5;
+      y -= 0.5;
+      z -= 0.5;
+
+      int i1 = static_cast<int>(floor(x));
+      int j1 = static_cast<int>(floor(y));
+      int k1 = static_cast<int>(floor(z));
+
+      int i2 = i1 + 1;
+      int j2 = j1 + 1;
+      int k2 = k1 + 1;
+
+      float a = x - i1;
+      float b = y - j1;
+      float c = z - k1;
+
+      hipTextureGetAddress < addressMode > (i1, width);
+      hipTextureGetAddress < addressMode > (i2, width);
+      hipTextureGetAddress < addressMode > (j1, height);
+      hipTextureGetAddress < addressMode > (j2, height);
+      hipTextureGetAddress < addressMode > (k1, depth);
+      hipTextureGetAddress < addressMode > (k2, depth);
+
+      float t111 = hipTextureGetValue < addressMode
+          > (data, i1, width, j1, height, k1, depth);
+      float t211 = hipTextureGetValue < addressMode
+          > (data, i2, width, j1, height, k1, depth);
+      float t121 = hipTextureGetValue < addressMode
+          > (data, i1, width, j2, height, k1, depth);
+      float t112 = hipTextureGetValue < addressMode
+          > (data, i1, width, j1, height, k2, depth);
+      float t122 = hipTextureGetValue < addressMode
+          > (data, i1, width, j2, height, k2, depth);
+      float t212 = hipTextureGetValue < addressMode
+          > (data, i2, width, j1, height, k2, depth);
+      float t221 = hipTextureGetValue < addressMode
+          > (data, i2, width, j2, height, k1, depth);
+      float t222 = hipTextureGetValue < addressMode
+          > (data, i2, width, j2, height, k2, depth);
+
+      result =
+                (1 - a) * (1 - b) * (1 - c) * t111 + a * (1 - b) * (1 - c) * t211 +
+                (1 - a) * b * (1 - c) * t121 + a * b * (1 - c) * t221 +
+                (1 - a) * (1 - b) * c * t112 + a * (1 - b) * c * t212 +
+                (1 - a) * b * c * t122 + a * b * c * t222;
+
+    }
+      break;
+  }
+  return result;
+}

projects/hip/tests/src/texture/hipTextureObj1DCheckModes.cpp

@@ -0,0 +1,121 @@
+/* HIT_START
+ * BUILD: %t %s ../test_common.cpp
+ * TEST: %t
+ * HIT_END
+ */
+#include <stdlib.h>
+#include <stdio.h>
+#include <string.h>
+
+#include <hip/hip_runtime.h>
+#include "test_common.h"
+#include "hipTextureHelper.hpp"
+
+template<bool normalizedCoords>
+__global__ void tex1DKernel(float *outputData, hipTextureObject_t textureObject,
+                            int width, float offsetX) {
+  int x = blockIdx.x * blockDim.x + threadIdx.x;
+  outputData[x] = tex1D<float>(textureObject, normalizedCoords ? (x + offsetX) / width : x + offsetX);
+}
+
+template<hipTextureAddressMode addressMode, hipTextureFilterMode filterMode, bool normalizedCoords>
+bool runTest(const int width, const float offsetX) {
+  printf("%s(addressMode=%d, filterMode=%d, normalizedCoords=%d, width=%d, offsetX=%f)\n", __FUNCTION__,
+         addressMode, filterMode, normalizedCoords, width, offsetX);
+  bool testResult = true;
+  unsigned int size = width * sizeof(float);
+  float *hData = (float*) malloc(size);
+  memset(hData, 0, size);
+  for (int j = 0; j < width; j++) {
+    hData[j] = j;
+  }
+
+  hipChannelFormatDesc channelDesc = hipCreateChannelDesc(
+      32, 0, 0, 0, hipChannelFormatKindFloat);
+  hipArray *hipArray;
+  hipMallocArray(&hipArray, &channelDesc, width);
+
+  HIPCHECK(hipMemcpy2DToArray(hipArray, 0, 0, hData, width * sizeof(float), width * sizeof(float), 1, hipMemcpyHostToDevice));
+
+  hipResourceDesc resDesc;
+  memset(&resDesc, 0, sizeof(resDesc));
+  resDesc.resType = hipResourceTypeArray;
+  resDesc.res.array.array = hipArray;
+
+  // Specify texture object parameters
+  hipTextureDesc texDesc;
+  memset(&texDesc, 0, sizeof(texDesc));
+  texDesc.addressMode[0] = addressMode;
+  texDesc.filterMode = filterMode;
+  texDesc.readMode = hipReadModeElementType;
+  texDesc.normalizedCoords = normalizedCoords;
+
+  // Create texture object
+  hipTextureObject_t textureObject = 0;
+  hipCreateTextureObject(&textureObject, &resDesc, &texDesc, NULL);
+
+  float *dData = NULL;
+  hipMalloc((void**) &dData, size);
+
+  dim3 dimBlock(16, 1, 1);
+  dim3 dimGrid((width + dimBlock.x - 1)/ dimBlock.x, 1, 1);
+
+  hipLaunchKernelGGL(tex1DKernel<normalizedCoords>, dimGrid, dimBlock, 0, 0, dData,
+                     textureObject, width, offsetX);
+
+  hipDeviceSynchronize();
+
+  float *hOutputData = (float*) malloc(size);
+  memset(hOutputData, 0, size);
+  hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost);
+
+  for (int j = 0; j < width; j++) {
+    float expectedValue = getExpectedValue<addressMode, filterMode>(width, offsetX + j, hData);
+    if (!hipTextureSamplingVerify<float, filterMode>(hOutputData[j], expectedValue)) {
+      printf("mismatched [ %d ]:%f ----%f\n", j, hOutputData[j], expectedValue);
+      testResult = false;
+      break;
+    }
+  }
+
+  hipDestroyTextureObject(textureObject);
+  hipFree(dData);
+  hipFreeArray(hipArray);
+  free(hData);
+  free(hOutputData);
+  printf("%s %s\n", __FUNCTION__, testResult ? "succeeded":"failed");
+  return testResult;
+}
+
+int main(int argc, char **argv) {
+  bool testResult = true;
+  testResult = testResult && runTest<hipAddressModeClamp, hipFilterModePoint, false>(256, -3);
+  testResult = testResult && runTest<hipAddressModeClamp, hipFilterModePoint, false>(256, 4);
+
+  testResult = testResult && runTest<hipAddressModeBorder, hipFilterModePoint, false>(256, -8.5);
+  testResult = testResult && runTest<hipAddressModeBorder, hipFilterModePoint, false>(256, 12.5);
+
+  testResult = testResult && runTest<hipAddressModeClamp, hipFilterModeLinear, false>(256, -3);
+  testResult = testResult && runTest<hipAddressModeClamp, hipFilterModeLinear, false>(256, 4);
+
+  testResult = testResult && runTest<hipAddressModeBorder, hipFilterModeLinear, false>(256, -8.5);
+  testResult = testResult && runTest<hipAddressModeBorder, hipFilterModeLinear, false>(256, 12.5);
+
+  testResult = testResult && runTest<hipAddressModeClamp, hipFilterModePoint, true>(256, -3);
+  testResult = testResult && runTest<hipAddressModeClamp, hipFilterModePoint, true>(256, 4);
+
+  testResult = testResult && runTest<hipAddressModeBorder, hipFilterModePoint, true>(256, -8.5);
+  testResult = testResult && runTest<hipAddressModeBorder, hipFilterModePoint, true>(256, 12.5);
+
+  testResult = testResult && runTest<hipAddressModeClamp, hipFilterModeLinear, true>(256, -3);
+  testResult = testResult && runTest<hipAddressModeClamp, hipFilterModeLinear, true>(256, 4);
+
+  testResult = testResult && runTest<hipAddressModeBorder, hipFilterModeLinear, true>(256, -8.5);
+  testResult = testResult && runTest<hipAddressModeBorder, hipFilterModeLinear, true>(256, 12.5);
+
+  if (testResult) {
+    passed();
+  } else {
+    exit (EXIT_FAILURE);
+  }
+}

projects/hip/tests/src/texture/hipTextureObj2DCheckModes.cpp

@@ -0,0 +1,136 @@
+/* HIT_START
+ * BUILD: %t %s ../test_common.cpp
+ * TEST: %t
+ * HIT_END
+ */
+#include <stdlib.h>
+#include <stdio.h>
+#include <string.h>
+
+#include <hip/hip_runtime.h>
+#include "test_common.h"
+#include "hipTextureHelper.hpp"
+
+template<bool normalizedCoords>
+__global__ void tex2DKernel(float *outputData, hipTextureObject_t textureObject,
+                            int width, int height, float offsetX,
+                            float offsetY) {
+  int x = blockIdx.x * blockDim.x + threadIdx.x;
+  int y = blockIdx.y * blockDim.y + threadIdx.y;
+  outputData[y * width + x] = tex2D<float>(textureObject,
+                                           normalizedCoords ? (x + offsetX) / width : x + offsetX,
+                                           normalizedCoords ? (y + offsetY) / height : y + offsetY);
+}
+
+template<hipTextureAddressMode addressMode, hipTextureFilterMode filterMode, bool normalizedCoords>
+bool runTest(const int width, const int height, const float offsetX, const float offsetY) {
+  printf("%s(addressMode=%d, filterMode=%d, normalizedCoords=%d, width=%d, height=%d, offsetX=%f, offsetY=%f)\n",
+         __FUNCTION__, addressMode, filterMode, normalizedCoords, width, height, offsetX, offsetY);
+  bool testResult = true;
+  unsigned int size = width * height * sizeof(float);
+  float *hData = (float*) malloc(size);
+  memset(hData, 0, size);
+  for (int i = 0; i < height; i++) {
+    for (int j = 0; j < width; j++) {
+      int index = i * width + j;
+      hData[index] = index;
+    }
+  }
+
+  hipChannelFormatDesc channelDesc = hipCreateChannelDesc(
+      32, 0, 0, 0, hipChannelFormatKindFloat);
+  hipArray *hipArray;
+  hipMallocArray(&hipArray, &channelDesc, width, height);
+
+  HIPCHECK(hipMemcpy2DToArray(hipArray, 0, 0, hData, width * sizeof(float), width * sizeof(float), height, hipMemcpyHostToDevice));
+
+  hipResourceDesc resDesc;
+  memset(&resDesc, 0, sizeof(resDesc));
+  resDesc.resType = hipResourceTypeArray;
+  resDesc.res.array.array = hipArray;
+
+  // Specify texture object parameters
+  hipTextureDesc texDesc;
+  memset(&texDesc, 0, sizeof(texDesc));
+  texDesc.addressMode[0] = addressMode;
+  texDesc.addressMode[1] = addressMode;
+  texDesc.filterMode = filterMode;
+  texDesc.readMode = hipReadModeElementType;
+  texDesc.normalizedCoords = normalizedCoords;
+
+  // Create texture object
+  hipTextureObject_t textureObject = 0;
+  hipCreateTextureObject(&textureObject, &resDesc, &texDesc, NULL);
+
+  float *dData = NULL;
+  hipMalloc((void**) &dData, size);
+
+  dim3 dimBlock(16, 16, 1);
+  dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x, (height + dimBlock.y -1)/ dimBlock.y, 1);
+
+  hipLaunchKernelGGL(tex2DKernel<normalizedCoords>, dimGrid, dimBlock, 0, 0, dData,
+                     textureObject, width, height, offsetX, offsetY);
+
+  hipDeviceSynchronize();
+
+  float *hOutputData = (float*) malloc(size);
+  memset(hOutputData, 0, size);
+  hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost);
+
+  for (int i = 0; i < height; i++) {
+    for (int j = 0; j < width; j++) {
+      int index = i * width + j;
+      float expectedValue = getExpectedValue<addressMode, filterMode>(width, height,
+                                                    offsetX + j, offsetY + i, hData);
+      if (!hipTextureSamplingVerify<float, filterMode>(hOutputData[index], expectedValue)) {
+        printf("mismatched [ %d %d ]:%f ----%f\n", j, i, hOutputData[index], expectedValue);
+        testResult = false;
+        goto line1;
+      }
+    }
+  }
+line1:
+  hipDestroyTextureObject(textureObject);
+  hipFree(dData);
+  hipFreeArray(hipArray);
+  free(hData);
+  free(hOutputData);
+  printf("%s %s\n", __FUNCTION__, testResult ? "succeeded":"failed");
+  return testResult;
+}
+
+int main(int argc, char **argv) {
+  bool testResult = true;
+
+  testResult = testResult && runTest<hipAddressModeClamp, hipFilterModePoint, false>(256, 256, -3.9, 6.1);
+  testResult = testResult && runTest<hipAddressModeClamp, hipFilterModePoint, false>(256, 256, 4.4, -7.0);
+
+  testResult = testResult && runTest<hipAddressModeBorder, hipFilterModePoint, false>(256, 256, -8.5, 2.9);
+  testResult = testResult && runTest<hipAddressModeBorder, hipFilterModePoint, false>(256, 256, 12.5, 6.7);
+
+  testResult = testResult && runTest<hipAddressModeClamp, hipFilterModeLinear, false>(256, 256, -0.4, -0.4);
+  testResult = testResult && runTest<hipAddressModeClamp, hipFilterModeLinear, false>(256, 256, 4, 14.6);
+
+  // The following two cases have quite big deviation on Cpu and Gpu in 2D, so comment them out temporarily.
+  testResult = testResult && runTest<hipAddressModeBorder, hipFilterModeLinear, false>(256, 256, -0.4, 0.4);
+  testResult = testResult && runTest<hipAddressModeBorder, hipFilterModeLinear, false>(256, 256, 12.5, 23.7);
+
+  testResult = testResult && runTest<hipAddressModeClamp, hipFilterModePoint, true>(256, 256, -3, 8.9);
+  testResult = testResult && runTest<hipAddressModeClamp, hipFilterModePoint, true>(256, 256, 4, -0.1);
+
+  testResult = testResult && runTest<hipAddressModeBorder, hipFilterModePoint, true>(256, 256, -8.5, 15.9);
+  testResult = testResult && runTest<hipAddressModeBorder, hipFilterModePoint, true>(256, 256, 12.5, -17.9);
+
+  testResult = testResult && runTest<hipAddressModeClamp, hipFilterModeLinear, true>(256, 256, -3, 5.8);
+  testResult = testResult && runTest<hipAddressModeClamp, hipFilterModeLinear, true>(256, 256, 4, 9.1);
+
+  // The following two cases have quite big deviation on Cpu and Gpu in 2D, so comment them out temporarily.
+  testResult = testResult && runTest<hipAddressModeBorder, hipFilterModeLinear, true>(256, 256, -8.5, 6.6);
+  testResult = testResult && runTest<hipAddressModeBorder, hipFilterModeLinear, true>(256, 256, 12.5, 0.01);
+
+  if (testResult) {
+    passed();
+  } else {
+    exit (EXIT_FAILURE);
+  }
+}

projects/hip/tests/src/texture/hipTextureObj3DCheckModes.cpp

@@ -0,0 +1,153 @@
+/* HIT_START
+ * BUILD: %t %s ../test_common.cpp
+ * TEST: %t
+ * HIT_END
+ */
+#include <stdlib.h>
+#include <stdio.h>
+#include <string.h>
+
+#include <hip/hip_runtime.h>
+#include "test_common.h"
+#include "hipTextureHelper.hpp"
+
+
+template<bool normalizedCoords>
+__global__ void tex3DKernel(float *outputData, hipTextureObject_t textureObject,
+                            int width, int height, int depth, float offsetX,
+                            float offsetY, float offsetZ) {
+  int x = blockIdx.x * blockDim.x + threadIdx.x;
+  int y = blockIdx.y * blockDim.y + threadIdx.y;
+  int z = blockIdx.z * blockDim.z + threadIdx.z;
+  outputData[z * width * depth + y * width + x] = tex3D<float>(textureObject,
+                        normalizedCoords ? (x + offsetX) / width : x + offsetX,
+                        normalizedCoords ? (y + offsetY) / height : y + offsetY,
+                        normalizedCoords ? (z + offsetZ) / depth : z + offsetZ);
+}
+
+template<hipTextureAddressMode addressMode, hipTextureFilterMode filterMode, bool normalizedCoords>
+bool runTest(const int width, const int height, const int depth, const float offsetX, const float offsetY, const float offsetZ) {
+  printf("%s(addressMode=%d, filterMode=%d, normalizedCoords=%d, width=%d, height=%d, depth=%d, offsetX=%f, offsetY=%f, offsetZ=%f)\n",
+      __FUNCTION__, addressMode, filterMode, normalizedCoords, width, height,
+      depth, offsetX, offsetY, offsetZ);
+  bool testResult = true;
+  unsigned int size = width * height * depth * sizeof(float);
+  float *hData = (float*) malloc(size);
+  memset(hData, 0, size);
+
+  for (int i = 0; i < depth; i++) {
+    for (int j = 0; j < height; j++) {
+      for (int k = 0; k < width; k++) {
+        int index = i * width * depth + j * width + k;
+        hData[index] = index;
+      }
+    }
+  }
+
+  // Allocate array and copy image data
+  hipChannelFormatDesc channelDesc = hipCreateChannelDesc<float>();
+  hipArray *arr;
+
+  HIPCHECK(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, height, depth), hipArrayDefault));
+  hipMemcpy3DParms myparms = {0};
+  myparms.srcPos = make_hipPos(0,0,0);
+  myparms.dstPos = make_hipPos(0,0,0);
+  myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(float), width, height);
+  myparms.dstArray = arr;
+  myparms.extent = make_hipExtent(width, height, depth);
+  myparms.kind = hipMemcpyHostToDevice;
+
+  HIPCHECK(hipMemcpy3D(&myparms));
+
+  hipResourceDesc resDesc;
+  memset(&resDesc, 0, sizeof(resDesc));
+  resDesc.resType = hipResourceTypeArray;
+  resDesc.res.array.array = arr;
+
+  // Specify texture object parameters
+  hipTextureDesc texDesc;
+  memset(&texDesc, 0, sizeof(texDesc));
+  texDesc.addressMode[0] = addressMode;
+  texDesc.addressMode[1] = addressMode;
+  texDesc.addressMode[2] = addressMode;
+  texDesc.filterMode = filterMode;
+  texDesc.readMode = hipReadModeElementType;
+  texDesc.normalizedCoords = normalizedCoords;
+
+  // Create texture object
+  hipTextureObject_t textureObject = 0;
+  hipCreateTextureObject(&textureObject, &resDesc, &texDesc, NULL);
+
+  float *dData = NULL;
+  hipMalloc((void**) &dData, size);
+  hipMemset(dData, 0, size);
+  dim3 dimBlock(8, 8, 8); // 512 threads
+  dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x, (height + dimBlock.y -1)/ dimBlock.y,
+               (depth + dimBlock.z - 1) / dimBlock.z);
+
+  hipLaunchKernelGGL(tex3DKernel<normalizedCoords>, dimGrid, dimBlock, 0, 0, dData,
+                     textureObject, width, height, depth, offsetX, offsetY, offsetZ);
+
+  hipDeviceSynchronize();
+
+  float *hOutputData = (float*) malloc(size);
+  memset(hOutputData, 0, size);
+  hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost);
+  for (int i = 0; i < depth; i++) {
+    for (int j = 0; j < height; j++) {
+      for (int k = 0; k < width; k++) {
+        int index = i * width * depth + j * width + k;
+        float expectedValue = getExpectedValue<addressMode, filterMode>(
+            width, height, depth, offsetX + k, offsetY + j, offsetZ + i, hData);
+
+        if (!hipTextureSamplingVerify<float, filterMode>(hOutputData[index], expectedValue)) {
+          printf("mismatched [ %d %d %d]:%f ----%f\n", k, j, i, hOutputData[index], expectedValue);
+          testResult = false;
+          goto line1;
+        }
+      }
+    }
+  }
+line1:
+  hipDestroyTextureObject(textureObject);
+  hipFree(dData);
+  hipFreeArray(arr);
+  free(hData);
+  free(hOutputData);
+  printf("%s %s\n", __FUNCTION__, testResult ? "succeeded":"failed");
+  return testResult;
+}
+
+int main(int argc, char **argv) {
+  bool testResult = true;
+
+  testResult = testResult && runTest<hipAddressModeClamp, hipFilterModePoint, false>(256, 256, 256, -3.9, 6.1, 9.5);
+  testResult = testResult && runTest<hipAddressModeClamp, hipFilterModePoint, false>(256, 256, 256, 4.4, -7.0, 5.3);
+
+  testResult = testResult && runTest<hipAddressModeBorder, hipFilterModePoint, false>(256, 256, 256, -8.5, 2.9, 5.8);
+  testResult = testResult && runTest<hipAddressModeBorder, hipFilterModePoint, false>(256, 256, 256, 12.5, 6.7, 11.4);
+
+  testResult = testResult && runTest<hipAddressModeClamp, hipFilterModeLinear, false>(256, 256, 256, -0.4, -0.4, -0.4);
+  testResult = testResult && runTest<hipAddressModeClamp, hipFilterModeLinear, false>(256, 256, 256, 4, 14.6, -0.3);
+
+  testResult = testResult && runTest<hipAddressModeBorder, hipFilterModeLinear, false>(256, 256, 256, 6.9, 7.4, 0.4);
+  testResult = testResult && runTest<hipAddressModeBorder, hipFilterModeLinear, false>(256, 256, 256, 12.5, 23.7, 0.34);
+
+  testResult = testResult && runTest<hipAddressModeClamp, hipFilterModePoint, true>(256, 256, 256, -3, 8.9, -4);
+  testResult = testResult && runTest<hipAddressModeClamp, hipFilterModePoint, true>(256, 256, 256, 4, -0.1, 8.2);
+
+  testResult = testResult && runTest<hipAddressModeBorder, hipFilterModePoint, true>(256, 256, 256, -8.5, 15.9, 0.1);
+  testResult = testResult && runTest<hipAddressModeBorder, hipFilterModePoint, true>(256, 256, 256, 12.5, -17.9, -0.35);
+
+  testResult = testResult && runTest<hipAddressModeClamp, hipFilterModeLinear, true>(256, 256, 256, -3, 5.8, 0.89);
+  testResult = testResult && runTest<hipAddressModeClamp, hipFilterModeLinear, true>(256, 256, 256, 4, 9.1, 2.08);
+
+  testResult = testResult && runTest<hipAddressModeBorder, hipFilterModeLinear, true>(256, 256, 256, -8.5, 6.6, 3.67);
+  testResult = testResult && runTest<hipAddressModeBorder, hipFilterModeLinear, true>(256, 256, 256, 12.5, 0.01, -9.9);
+
+  if (testResult) {
+    passed();
+  } else {
+    exit (EXIT_FAILURE);
+  }
+}

projects/hip/tests/src/texture/simpleTexture3D.cpp

@@ -87,11 +87,8 @@ void runTest(int width,int height,int depth,texture<T, hipTextureType3D, hipRead
     myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height);
     myparms.dstArray = arr;
     myparms.extent = make_hipExtent(width, height, depth);
-#ifdef __HIP_PLATFORM_NVIDIA__
-    myparms.kind = cudaMemcpyHostToDevice;
-#else
     myparms.kind = hipMemcpyHostToDevice;
-#endif
+
     HIPCHECK(hipMemcpy3D(&myparms));
 
     // set texture parameters