SWDEV-299127 - Merge 'develop' into 'amd-staging'

Change-Id: Ie58d3b7ff26fcdb906a49c92d01f31cf8a07b759

bin/hipvars.pm

@@ -26,7 +26,7 @@ use Cwd;
 use File::Basename;
 
 $HIP_BASE_VERSION_MAJOR = "5";
-$HIP_BASE_VERSION_MINOR = "0";
+$HIP_BASE_VERSION_MINOR = "1";
 $HIP_BASE_VERSION_PATCH = "0";
 
 #---

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

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");

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

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;
+}

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);
+  }
+}

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);
+  }
+}

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);
+  }
+}

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