SWDEV-305992 - Improve texture tests (#2646)

Enable more texture supports for hip-rocclr.
Skip texture tests on device that doesn't support images.
SWDEV-322257 - Fix issues of image tests skipped on
NVidia devices.

Change-Id: Ia99d06b1e97fc945f1b740e47710f4dcd70f38ca

samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp

@@ -25,7 +25,9 @@ THE SOFTWARE.
 texture<float, 2, hipReadModeElementType> tex;
 
 extern "C" __global__ void tex2dKernel(float* outputData, int width, int height) {
+#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
     int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
     int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
     outputData[y * width + x] = tex2D(tex, x, y);
+#endif
 }

samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp

@@ -127,7 +127,20 @@ bool runTest(int argc, char** argv) {
     return testResult;
 }
 
+inline bool isImageSupported() {
+    int imageSupport = 1;
+#ifdef __HIP_PLATFORM_AMD__
+    HIPCHECK(hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport,
+                              0));
+#endif
+  return imageSupport != 0;
+}
+
 int main(int argc, char** argv) {
+    if (!isImageSupported()) {
+      printf("Texture is not support on the device. Skipped.\n");
+      return 0;
+    }
     hipInit(0);
     testResult = runTest(argc, argv);
     printf("%s ...\n", testResult ? "PASSED" : "FAILED");

tests/catch/include/hip_test_common.hh

@@ -115,4 +115,22 @@ static inline int RAND_R(unsigned* rand_seed)
       return rand_r(rand_seed);
   #endif
 }
+
+inline bool isImageSupported() {
+    int imageSupport = 1;
+#ifdef __HIP_PLATFORM_AMD__
+    int device;
+    HIP_CHECK(hipGetDevice(&device));
+    HIPCHECK(hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport,
+                                   device));
+#endif
+  return imageSupport != 0;
 }
+
+}
+
+// This must be called in the beginning of image test app's main() to indicate whether image
+// is supported.
+#define checkImageSupport()                                                                \
+    if (!HipTest::isImageSupported())                                                      \
+        { printf("Texture is not support on the device. Skipped.\n"); return; }

tests/catch/unit/texture/hipCreateTextureObject_ArgValidation.cc

@@ -25,6 +25,8 @@ THE SOFTWARE.
  * Validate argument list of texture object api.
  */
 TEST_CASE("Unit_hipCreateTextureObject_ArgValidation") {
+  checkImageSupport();
+
   float *texBuf;
   hipError_t ret;
   constexpr int xsize = 32;

tests/catch/unit/texture/hipCreateTextureObject_Array.cc

@@ -23,6 +23,8 @@ THE SOFTWARE.
  * Validates Array Resource texture object with negative/functional tests.
  */
 TEST_CASE("Unit_hipCreateTextureObject_ArrayResource") {
+  checkImageSupport();
+
   hipError_t ret;
   hipResourceDesc resDesc;
   hipTextureDesc texDesc;
@@ -47,6 +49,8 @@ TEST_CASE("Unit_hipCreateTextureObject_ArrayResource") {
  * with negative/functional tests.
  */
 TEST_CASE("Unit_hipCreateTextureObject_MmArrayResource") {
+  checkImageSupport();
+
   hipError_t ret;
   hipResourceDesc resDesc;
   hipTextureDesc texDesc;

tests/catch/unit/texture/hipCreateTextureObject_Linear.cc

@@ -26,6 +26,8 @@ THE SOFTWARE.
  * Validates Linear Resource texture object with negative/functional tests.
  */
 TEST_CASE("Unit_hipCreateTextureObject_LinearResource") {
+  checkImageSupport();
+
   float *texBuf;
   hipError_t ret;
   constexpr int xsize = 32;

tests/catch/unit/texture/hipCreateTextureObject_Pitch2D.cc

@@ -29,6 +29,8 @@ THE SOFTWARE.
  * Validates Pitch2D Resource texture object with negative and functional tests
  */
 TEST_CASE("Unit_hipCreateTextureObject_Pitch2DResource") {
+  checkImageSupport();
+
   hipError_t ret;
   hipResourceDesc resDesc;
   hipTextureDesc texDesc;

tests/catch/unit/texture/hipTextureObjFetchVector.cc

@@ -23,10 +23,12 @@ THE SOFTWARE.
 
 template <typename T>
 __global__ void tex1dKernelFetch(T *val, hipTextureObject_t obj, int N) {
+#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
   int k = blockIdx.x * blockDim.x + threadIdx.x;
   if (k < N) {
     val[k] = tex1Dfetch<T>(obj, k);
   }
+#endif
 }
 
 template <typename T>
@@ -203,6 +205,8 @@ bool runTest(const char *description) {
 }
 
 TEST_CASE("Unit_hipTextureFetch_vector") {
+  checkImageSupport();
+
   // test for char
   runTest<char1>("char1");
   runTest<char2>("char2");

tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp

@@ -592,13 +592,8 @@ bool testTexSingleStreamMultGPU(unsigned int numOfGPUs,
 
 int main(int argc, char** argv) {
   HipTest::parseStandardArguments(argc, argv, true);
-  int imageSupport = 0;
-  hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport,
-                            p_gpuDevice);
-  if (!imageSupport) {
-    printf("Texture is not support on the device\n");
-    passed();
-  }
+  checkImageSupport();
+
   bool TestPassed = true;
   if (p_tests == 0x01) {
     TestPassed = testTexType<float>(HIP_AD_FORMAT_FLOAT,

tests/src/test_common.h

@@ -564,5 +564,20 @@ struct MemTraits<MemcpyAsync> {
     }
 };
 
+inline bool isImageSupported() {
+    int imageSupport = 1;
+#ifdef __HIP_PLATFORM_AMD__
+    HIPCHECK(hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport,
+                              p_gpuDevice));
+#endif
+  return imageSupport != 0;
+}
+
 };  // namespace HipTest
+
+// This must be called in the beginning of image test app's main() to indicate whether image
+// is supported.
+#define checkImageSupport()                                                                \
+    if (!HipTest::isImageSupported())                                                      \
+        { printf("Texture is not support on the device. Skipped.\n"); passed(); }
 #endif //__cplusplus

tests/src/texture/hipBindTex2DPitch.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
  */
@@ -32,16 +32,18 @@ texture<TYPE_t, 2, hipReadModeElementType> tex;
 
 // texture object is a kernel argument
 __global__ void texture2dCopyKernel( TYPE_t* dst) {
-
+#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
     int x = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
     int y = hipThreadIdx_y + hipBlockIdx_y * hipBlockDim_y;
     if ( (x< SIZE_W) && (y< SIZE_H) ){
         dst[SIZE_W*y+x] = tex2D(tex, x, y);
     }
+#endif
 }
 
 int main (void)
 {
+    checkImageSupport();
     TYPE_t* B;
     TYPE_t* A;
     TYPE_t* devPtrB;
@@ -49,8 +51,8 @@ int main (void)
 
     B = new TYPE_t[SIZE_H*SIZE_W];
     A = new TYPE_t[SIZE_H*SIZE_W];
-    for(size_t i=1; i <= (SIZE_H*SIZE_W); i++){
-        A[i-1] = i;
+    for (size_t i = 0; i < (SIZE_H * SIZE_W); i++) {
+      A[i] = i + 1;
     }
 
     size_t devPitchA, tex_ofs;
@@ -58,12 +60,14 @@ int main (void)
     HIPCHECK(hipMemcpy2D(devPtrA, devPitchA, A, SIZE_W*sizeof(TYPE_t),
             SIZE_W*sizeof(TYPE_t), SIZE_H, hipMemcpyHostToDevice));
 
+    tex.addressMode[0] = hipAddressModeClamp;
+    tex.addressMode[1] = hipAddressModeClamp;
     tex.normalized = false;
     HIPCHECK(hipBindTexture2D(&tex_ofs, &tex, devPtrA, &tex.channelDesc,
                                        SIZE_W, SIZE_H, devPitchA));
     HIPCHECK(hipMalloc((void**)&devPtrB, SIZE_W*sizeof(TYPE_t)*SIZE_H)) ;
 
-    hipLaunchKernelGGL(texture2dCopyKernel, dim3(4,4,1), dim3(32,32,1), 0, 0, devPtrB);
+    hipLaunchKernelGGL(texture2dCopyKernel, dim3(3, 2, 1), dim3(4, 4, 1), 0, 0, devPtrB);
     hipDeviceSynchronize();
     HIPCHECK(hipMemcpy2D(B, SIZE_W*sizeof(TYPE_t), devPtrB, SIZE_W*sizeof(TYPE_t),
             SIZE_W*sizeof(TYPE_t), SIZE_H, hipMemcpyDeviceToHost));

tests/src/texture/hipBindTexRef1DFetch.cpp

@@ -35,15 +35,19 @@ THE SOFTWARE.
 texture<float, 1, hipReadModeElementType> tex;
 
 __global__ void kernel(float *out) {
+#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
   int x = blockIdx.x * blockDim.x + threadIdx.x;
   if(x<N){
       out[x] = tex1Dfetch(tex, x);
   }
+#endif
 }
 
 int runTest(void);
 
 int main(int argc, char **argv) {
+    checkImageSupport();
+
     int testResult = runTest();
     if (testResult) {
         passed();
@@ -70,7 +74,6 @@ int runTest() {
     HIPCHECK(hipMemcpy(texBuf, val, N * sizeof(float), hipMemcpyHostToDevice));
 
     tex.addressMode[0] = hipAddressModeClamp;
-    tex.addressMode[1] = hipAddressModeClamp;
     tex.filterMode = hipFilterModePoint;
     tex.normalized = 0;
 

tests/src/texture/hipGetChanDesc.cpp

@@ -36,13 +36,8 @@ using namespace std;
 bool runTest(void);
 
 int main(int argc, char** argv) {
-    int imageSupport = 0;
-    hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport,
-                              p_gpuDevice);
-    if (!imageSupport) {
-      printf("Texture is not support on the device\n");
-      passed();
-    }
+    checkImageSupport();
+
     bool testResult=runTest();
 
     if (testResult) {
@@ -60,7 +55,7 @@ hipArray *hipArray;
 HIPCHECK(hipMallocArray(&hipArray, &chan_desc,C,R,0));
 HIPCHECK(hipGetChannelDesc(&chan_test,hipArray));
 
-if((chan_test.x == 32)&&(chan_test.y == 0)&&(chan_test.z == 0)&&(chan_test.f == 0))
+if((chan_test.x == 32)&&(chan_test.y == 0)&&(chan_test.z == 0)&&(chan_test.f == hipChannelFormatKindSigned))
 	testResult=true;
 else
 	testResult=false;

tests/src/texture/hipNormalizedFloatValueTex.cpp

@@ -135,6 +135,7 @@ bool textureTest(texture<T, hipTextureType1D, hipReadModeNormalizedFloat> *tex)
     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));
 
+    tex->addressMode[0] = hipAddressModeClamp;
     tex->normalized = true;
     tex->channelDesc = desc;
     tex->filterMode = fMode;
@@ -172,15 +173,10 @@ bool runTest() {
 
 int main(int argc, char** argv)
 {
-    int imageSupport = 0;
-    hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport,
-                              p_gpuDevice);
-    if (!imageSupport) {
-      printf("Texture is not support on the device\n");
-      passed();
-    }
     HipTest::parseStandardArguments(argc, argv, true);
-    int device = 0;
+    checkImageSupport();
+
+    int device = p_gpuDevice;
     bool status = false;
     HIPCHECK(hipSetDevice(device));
     hipDeviceProp_t props;

tests/src/texture/hipTex1DFetchCheckModes.cpp

@@ -29,14 +29,18 @@ THE SOFTWARE.
 #define N 16
 #define offset 3
 __global__ void tex1dKernel(float *val, hipTextureObject_t obj) {
+#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
     int k = blockIdx.x * blockDim.x + threadIdx.x;
     if (k < N)
         val[k] = tex1Dfetch<float>(obj, k+offset);
+#endif
 }
 
 int runTest(hipTextureAddressMode, hipTextureFilterMode);
 
 int main(int argc, char **argv) {
+    checkImageSupport();
+
     int testResult = runTest(hipAddressModeClamp,hipFilterModePoint);
     testResult = testResult & runTest(hipAddressModeClamp,hipFilterModeLinear);
     testResult = testResult & runTest(hipAddressModeBorder,hipFilterModePoint);

tests/src/texture/hipTexObjPitch.cpp

@@ -96,13 +96,8 @@ void texture2Dtest()
 
 int main()
 {
-    int imageSupport = 0;
-    hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport,
-                              p_gpuDevice);
-    if (!imageSupport) {
-      printf("Texture is not support on the device\n");
-      passed();
-    }
+    checkImageSupport();
+
     texture2Dtest<float>();
     texture2Dtest<int>();
     texture2Dtest<unsigned char>();

tests/src/texture/hipTextureMipmapObj2D.cpp

@@ -171,14 +171,9 @@ bool runTest(int argc, char** argv) {
 }
 
 int main(int argc, char** argv) {
+  checkImageSupport();
+
   bool testResult = true;
-  int imageSupport = 0;
-  hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport,
-                            p_gpuDevice);
-  if (!imageSupport) {
-    printf("Texture is not support on the device\n");
-    passed();
-  }
 #ifdef _WIN32
   testResult = runTest(argc, argv);
 #else

tests/src/texture/hipTextureObj1DCheckModes.cpp

@@ -14,8 +14,10 @@
 template<bool normalizedCoords>
 __global__ void tex1DKernel(float *outputData, hipTextureObject_t textureObject,
                             int width, float offsetX) {
+#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
   int x = blockIdx.x * blockDim.x + threadIdx.x;
   outputData[x] = tex1D<float>(textureObject, normalizedCoords ? (x + offsetX) / width : x + offsetX);
+#endif
 }
 
 template<hipTextureAddressMode addressMode, hipTextureFilterMode filterMode, bool normalizedCoords>
@@ -88,6 +90,8 @@ bool runTest(const int width, const float offsetX) {
 }
 
 int main(int argc, char **argv) {
+  checkImageSupport();
+
   bool testResult = true;
   testResult = testResult && runTest<hipAddressModeClamp, hipFilterModePoint, false>(256, -3);
   testResult = testResult && runTest<hipAddressModeClamp, hipFilterModePoint, false>(256, 4);

tests/src/texture/hipTextureObj1DFetch.cpp

@@ -21,7 +21,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
  */
@@ -32,14 +32,18 @@ THE SOFTWARE.
 #define N 512
 
 __global__ void tex1dKernel(float *val, hipTextureObject_t obj) {
+#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
     int k = blockIdx.x * blockDim.x + threadIdx.x;
     if (k < N)
         val[k] = tex1Dfetch<float>(obj, k);
+#endif
 }
 
 int runTest(void);
 
 int main(int argc, char **argv) {
+    checkImageSupport();
+
     int testResult = runTest();
     if(testResult) {
         passed();
@@ -72,6 +76,7 @@ int runTest() {
     hipTextureDesc texDesc;
     memset(&texDesc, 0, sizeof(texDesc));
     texDesc.readMode = hipReadModeElementType;
+    texDesc.addressMode[0]= hipAddressModeClamp;
 
     // Creating texture object
     hipTextureObject_t texObj = 0;

tests/src/texture/hipTextureObj2D.cpp

@@ -22,13 +22,8 @@ __global__ void tex2DKernel(float* outputData, hipTextureObject_t textureObject,
 int runTest(int argc, char** argv);
 
 int main(int argc, char** argv) {
-    int imageSupport = 0;
-    hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport,
-                              p_gpuDevice);
-    if (!imageSupport) {
-      printf("Texture is not support on the device\n");
-      passed();
-    }
+    checkImageSupport();
+
     int testResult = runTest(argc, argv);
 
     if (testResult) {
@@ -70,8 +65,8 @@ int runTest(int argc, char** argv) {
     // Specify texture object parameters
     hipTextureDesc texDesc;
     memset(&texDesc, 0, sizeof(texDesc));
-    texDesc.addressMode[0] = hipAddressModeWrap;
-    texDesc.addressMode[1] = hipAddressModeWrap;
+    texDesc.addressMode[0] = hipAddressModeClamp;
+    texDesc.addressMode[1] = hipAddressModeClamp;
     texDesc.filterMode = hipFilterModePoint;
     texDesc.readMode = hipReadModeElementType;
     texDesc.normalizedCoords = 0;

tests/src/texture/hipTextureObj2DCheckModes.cpp

@@ -15,11 +15,13 @@ template<bool normalizedCoords>
 __global__ void tex2DKernel(float *outputData, hipTextureObject_t textureObject,
                             int width, int height, float offsetX,
                             float offsetY) {
+#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
   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);
+#endif
 }
 
 template<hipTextureAddressMode addressMode, hipTextureFilterMode filterMode, bool normalizedCoords>
@@ -100,6 +102,8 @@ line1:
 }
 
 int main(int argc, char **argv) {
+  checkImageSupport();
+
   bool testResult = true;
 
   testResult = testResult && runTest<hipAddressModeClamp, hipFilterModePoint, false>(256, 256, -3.9, 6.1);

tests/src/texture/hipTextureObj3DCheckModes.cpp

@@ -16,6 +16,7 @@ template<bool normalizedCoords>
 __global__ void tex3DKernel(float *outputData, hipTextureObject_t textureObject,
                             int width, int height, int depth, float offsetX,
                             float offsetY, float offsetZ) {
+#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
   int x = blockIdx.x * blockDim.x + threadIdx.x;
   int y = blockIdx.y * blockDim.y + threadIdx.y;
   int z = blockIdx.z * blockDim.z + threadIdx.z;
@@ -23,6 +24,7 @@ __global__ void tex3DKernel(float *outputData, hipTextureObject_t textureObject,
                         normalizedCoords ? (x + offsetX) / width : x + offsetX,
                         normalizedCoords ? (y + offsetY) / height : y + offsetY,
                         normalizedCoords ? (z + offsetZ) / depth : z + offsetZ);
+#endif
 }
 
 template<hipTextureAddressMode addressMode, hipTextureFilterMode filterMode, bool normalizedCoords>
@@ -119,6 +121,8 @@ line1:
 }
 
 int main(int argc, char **argv) {
+  checkImageSupport();
+
   bool testResult = true;
 
   testResult = testResult && runTest<hipAddressModeClamp, hipFilterModePoint, false>(256, 256, 256, -3.9, 6.1, 9.5);

tests/src/texture/hipTextureRef2D.cpp

@@ -1,5 +1,5 @@
 /* HIT_START
- * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_RUNTIME rocclr
+ * BUILD: %t %s ../test_common.cpp
  * TEST: %t
  * HIT_END
  */
@@ -14,14 +14,18 @@ texture<float, 2, hipReadModeElementType> tex;
 
 __global__ void tex2DKernel(float* outputData,
                             int width, int height) {
+#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
     int x = blockIdx.x * blockDim.x + threadIdx.x;
     int y = blockIdx.y * blockDim.y + threadIdx.y;
     outputData[y * width + x] = tex2D(tex, x, y);
+#endif
 }
 
 int runTest(int argc, char** argv);
 
 int main(int argc, char** argv) {
+    checkImageSupport();
+
     int testResult = runTest(argc, argv);
     if (testResult) {
         passed();
@@ -54,8 +58,8 @@ int runTest(int argc, char** argv) {
 
     hipMemcpyToArray(hipArray, 0, 0, hData, size, hipMemcpyHostToDevice);
 
-    tex.addressMode[0] = hipAddressModeWrap;
-    tex.addressMode[1] = hipAddressModeWrap;
+    tex.addressMode[0] = hipAddressModeClamp;
+    tex.addressMode[1] = hipAddressModeClamp;
     tex.filterMode = hipFilterModePoint;
     tex.normalized = 0;
 

tests/src/texture/simpleTexture2DLayered.cpp

@@ -21,7 +21,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
  */
@@ -34,9 +34,11 @@ texture<float, hipTextureType2DLayered> tex2DL;
 
 __global__ void simpleKernelLayeredArray(T* outputData,int width,int height,int layer)
 {
+#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
     unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
     unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
     outputData[layer*width*height + y*width + x] = tex2DLayered(tex2DL, x, y, layer);
+#endif
 }
 
 ////////////////////////////////////////////////////////////////////////////////
@@ -67,8 +69,8 @@ void runTest(int width,int height,int num_layers,texture<T, hipTextureType2DLaye
     HIPCHECK(hipMemcpy3D(&myparms));
 
     // set texture parameters
-    tex->addressMode[0] = hipAddressModeWrap;
-    tex->addressMode[1] = hipAddressModeWrap;
+    tex->addressMode[0] = hipAddressModeClamp;
+    tex->addressMode[1] = hipAddressModeClamp;
     tex->filterMode = hipFilterModePoint;
     tex->normalized = false;
 
@@ -103,6 +105,8 @@ void runTest(int width,int height,int num_layers,texture<T, hipTextureType2DLaye
 ////////////////////////////////////////////////////////////////////////////////
 int main(int argc, char **argv)
 {
+    checkImageSupport();
+
     runTest(512,512,5,&tex2DL);
     passed();
 }

tests/src/texture/simpleTexture3D.cpp

@@ -92,8 +92,8 @@ void runTest(int width,int height,int depth,texture<T, hipTextureType3D, hipRead
     HIPCHECK(hipMemcpy3D(&myparms));
 
     // set texture parameters
-    tex->addressMode[0] = hipAddressModeWrap;
-    tex->addressMode[1] = hipAddressModeWrap;
+    tex->addressMode[0] = hipAddressModeClamp;
+    tex->addressMode[1] = hipAddressModeClamp;
     tex->filterMode = hipFilterModePoint;
     tex->normalized = false;
 
@@ -126,13 +126,8 @@ void runTest(int width,int height,int depth,texture<T, hipTextureType3D, hipRead
 ////////////////////////////////////////////////////////////////////////////////
 int main(int argc, char **argv)
 {
-    int imageSupport = 0;
-    hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport,
-                              p_gpuDevice);
-    if (!imageSupport) {
-      printf("Texture is not support on the device\n");
-      passed();
-    }
+    checkImageSupport();
+
     printf("%s starting...\n", sampleName);
     for(int i=1;i<25;i++)
     {