diff --git a/docs/markdown/hip_kernel_language.md b/docs/markdown/hip_kernel_language.md
index 03bcd9810e..70a02acb29 100644
--- a/docs/markdown/hip_kernel_language.md
+++ b/docs/markdown/hip_kernel_language.md
@@ -494,7 +494,11 @@ Following is the list of supported floating-point intrinsics. Note that intrinsi
 | double __dsqrt_rn ( double  x ) <br><sub>Compute `√x` in round-to-nearest-even mode.</sub> |
 
 ## Texture Functions
-The supported Texture functions are listed in header files "texture_functions.h"(https://github.com/ROCm-Developer-Tools/HIP/blob/main/include/hip/hcc_detail/texture_functions.h) and"texture_indirect_functions.h" (https://github.com/ROCm-Developer-Tools/HIP/blob/main/include/hip/hcc_detail/texture_indirect_functions.h).
+The supported Texture functions are listed in header files "texture_fetch_functions.h"(https://github.com/ROCm-Developer-Tools/HIP/blob/main/include/hip/hcc_detail/texture_fetch_functions.h) and"texture_indirect_functions.h" (https://github.com/ROCm-Developer-Tools/HIP/blob/main/include/hip/hcc_detail/texture_indirect_functions.h).
+
+Texture functions are not supported on some devices.
+Macro __HIP_NO_IMAGE_SUPPORT == 1 can be used to check whether texture functions are not supported in device code.
+Attribute hipDeviceAttributeImageSupport can be queried to check whether texture functions are supported in host runtime code.
 
 ## Surface Functions
 Surface functions are not supported.
diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h
index f5311e767d..235a6d810a 100644
--- a/include/hip/hip_runtime_api.h
+++ b/include/hip/hip_runtime_api.h
@@ -442,7 +442,8 @@ typedef enum hipDeviceAttribute_t {
     hipDeviceAttributeIsLargeBar,                               ///< Whether it is LargeBar
     hipDeviceAttributeAsicRevision,                             ///< Revision of the GPU in this device
     hipDeviceAttributeCanUseStreamWaitValue,                    ///< '1' if Device supports hipStreamWaitValue32() and
-                                                                ///< hipStreamWaitValue64() , '0' otherwise.
+                                                                ///< hipStreamWaitValue64(), '0' otherwise.
+    hipDeviceAttributeImageSupport,                             ///< '1' if Device supports image, '0' otherwise.
 
     hipDeviceAttributeAmdSpecificEnd = 19999,
     hipDeviceAttributeVendorSpecificBegin = 20000,
diff --git a/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp b/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp
index 3bd8ed3f99..b4abf307dd 100644
--- a/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp
+++ b/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp
@@ -592,6 +592,13 @@ 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();
+  }
   bool TestPassed = true;
   if (p_tests == 0x01) {
     TestPassed = testTexType<float>(HIP_AD_FORMAT_FLOAT,
diff --git a/tests/src/runtimeApi/module/tex2d_kernel.cpp b/tests/src/runtimeApi/module/tex2d_kernel.cpp
index 5b558ffbf2..5ba0ebf89f 100644
--- a/tests/src/runtimeApi/module/tex2d_kernel.cpp
+++ b/tests/src/runtimeApi/module/tex2d_kernel.cpp
@@ -30,36 +30,44 @@ __device__ float deviceGlobalFloat;
 
 extern "C" __global__ void tex2dKernelFloat(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;
   if ((x < width) && (y < width)) {
     outputData[y * width + x] = tex2D(ftex, x, y);
   }
+#endif
 }
 
 extern "C" __global__ void tex2dKernelInt(int* 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;
   if ((x < width) && (y < width)) {
     outputData[y * width + x] = tex2D(itex, x, y);
   }
+#endif
 }
 
 extern "C" __global__ void tex2dKernelInt16(short* 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;
   if ((x < width) && (y < width)) {
     outputData[y * width + x] = tex2D(stex, x, y);
   }
+#endif
 }
 
 extern "C" __global__ void tex2dKernelInt8(char* 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;
   if ((x < width) && (y < width)) {
     outputData[y * width + x] = tex2D(ctex, x, y);
   }
+#endif
 }
diff --git a/tests/src/texture/hipGetChanDesc.cpp b/tests/src/texture/hipGetChanDesc.cpp
index b2ea1098cf..df3121cc4a 100644
--- a/tests/src/texture/hipGetChanDesc.cpp
+++ b/tests/src/texture/hipGetChanDesc.cpp
@@ -36,6 +36,13 @@ 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();
+    }
     bool testResult=runTest();
 
     if (testResult) {
diff --git a/tests/src/texture/hipNormalizedFloatValueTex.cpp b/tests/src/texture/hipNormalizedFloatValueTex.cpp
index 3ac14b314d..827ac69902 100644
--- a/tests/src/texture/hipNormalizedFloatValueTex.cpp
+++ b/tests/src/texture/hipNormalizedFloatValueTex.cpp
@@ -60,6 +60,7 @@ texture<unsigned short, hipTextureType1D, hipReadModeNormalizedFloat> texus;
 template<typename T>
 __global__ void normalizedValTextureTest(unsigned int numElements, float* pDst)
 {
+#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
     unsigned int elementID = hipThreadIdx_x;
     if(elementID >= numElements)
         return;
@@ -72,6 +73,7 @@ __global__ void normalizedValTextureTest(unsigned int numElements, float* pDst)
         pDst[elementID] = tex1D(texs, coord);
     else if(std::is_same<T, unsigned short>::value)
         pDst[elementID] = tex1D(texus, coord);
+#endif
 }
 
 bool textureVerifyFilterModePoint(float *hOutputData, float *expected, size_t size) {
@@ -171,6 +173,13 @@ 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;
     bool status = false;
diff --git a/tests/src/texture/hipTexObjPitch.cpp b/tests/src/texture/hipTexObjPitch.cpp
index 6869081a20..2e59d91d7c 100644
--- a/tests/src/texture/hipTexObjPitch.cpp
+++ b/tests/src/texture/hipTexObjPitch.cpp
@@ -29,11 +29,12 @@ THE SOFTWARE.
 // texture object is a kernel argument
 template <typename TYPE_t>
 __global__ void texture2dCopyKernel( hipTextureObject_t texObj, TYPE_t* dst,TYPE_t* A) {
-
+#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
     for(int i =0;i<SIZE_H;i++)
         for(int j = 0;j<SIZE_W;j++)
             dst[SIZE_W*i+j] = tex2D<TYPE_t>(texObj, j, i);
     __syncthreads();
+#endif
 }
 
 template <typename TYPE_t>
@@ -76,7 +77,7 @@ void texture2Dtest()
     texDescr.readMode = hipReadModeElementType;
 
     hipTextureObject_t texObj;
-    HIPCHECK( hipCreateTextureObject(&texObj, &texRes, &texDescr, NULL));
+    HIPCHECK(hipCreateTextureObject(&texObj, &texRes, &texDescr, NULL));
 
     HIPCHECK(hipMalloc((void**)&devPtrB, SIZE_W*sizeof(TYPE_t)*SIZE_H)) ;
 
@@ -95,6 +96,13 @@ void texture2Dtest()
 
 int main()
 {
+    int imageSupport = 0;
+    hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport,
+                              p_gpuDevice);
+    if (!imageSupport) {
+      printf("Texture is not support on the device\n");
+      passed();
+    }
     texture2Dtest<float>();
     texture2Dtest<int>();
     texture2Dtest<unsigned char>();
diff --git a/tests/src/texture/hipTextureMipmapObj2D.cpp b/tests/src/texture/hipTextureMipmapObj2D.cpp
index 719a427f0b..3a59a9640d 100644
--- a/tests/src/texture/hipTextureMipmapObj2D.cpp
+++ b/tests/src/texture/hipTextureMipmapObj2D.cpp
@@ -40,10 +40,12 @@ std::vector<unsigned int> mip_vector = {8, 4, 2, 1};
 __global__ void tex2DKernel(float* outputData, hipTextureObject_t textureObject, int width,
                             int height, float level) {
 #ifndef __gfx90a__
+#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] = tex2DLod<float>(textureObject, x, y, level);
 #endif
+#endif
 }
 
 bool runMipMapTest(unsigned int width, unsigned int height, unsigned int mipmap_level) {
@@ -148,6 +150,7 @@ bool runMipMapTest(unsigned int width, unsigned int height, unsigned int mipmap_
   hipDestroyTextureObject(textureObject);
   hipFree(dData);
   hipFreeArray(hipArray);
+  free(hData);
   return testResult;
 }
 
@@ -169,7 +172,13 @@ bool runTest(int argc, char** argv) {
 
 int main(int argc, char** argv) {
   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
diff --git a/tests/src/texture/hipTextureObj2D.cpp b/tests/src/texture/hipTextureObj2D.cpp
index 649c748fd2..18f40eab11 100644
--- a/tests/src/texture/hipTextureObj2D.cpp
+++ b/tests/src/texture/hipTextureObj2D.cpp
@@ -12,14 +12,23 @@
 
 __global__ void tex2DKernel(float* outputData, hipTextureObject_t textureObject, 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<float>(textureObject, x, y);
+#endif
 }
 
 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();
+    }
     int testResult = runTest(argc, argv);
 
     if (testResult) {
@@ -70,7 +79,6 @@ int runTest(int argc, char** argv) {
     // Create texture object
     hipTextureObject_t textureObject = 0;
     hipCreateTextureObject(&textureObject, &resDesc, &texDesc, NULL);
-
     float* dData = NULL;
     hipMalloc((void**)&dData, size);
 
@@ -104,5 +112,6 @@ int runTest(int argc, char** argv) {
     hipDestroyTextureObject(textureObject);
     hipFree(dData);
     hipFreeArray(hipArray);
+    free(hData);
     return testResult;
 }
diff --git a/tests/src/texture/simpleTexture3D.cpp b/tests/src/texture/simpleTexture3D.cpp
index a156ff46bf..06f2a31c97 100644
--- a/tests/src/texture/simpleTexture3D.cpp
+++ b/tests/src/texture/simpleTexture3D.cpp
@@ -42,6 +42,7 @@ __global__ void simpleKernel3DArray(T* outputData,
                                     int width,
                                     int height,int depth)
 {
+#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
     for (int i = 0; i < depth; i++) {
         for (int j = 0; j < height; j++) {
             for (int k = 0; k < width; k++) {
@@ -54,6 +55,7 @@ __global__ void simpleKernel3DArray(T* outputData,
             }
         }
     }
+#endif
 }
 
 ////////////////////////////////////////////////////////////////////////////////
@@ -127,6 +129,13 @@ 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();
+    }
     printf("%s starting...\n", sampleName);
     for(int i=1;i<25;i++)
     {