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

Change-Id: I222e7e56773ba82867bfa72ce47f2515d2cb5ed0

bin/hipcc.pl

@@ -152,7 +152,10 @@ if ($HIP_PLATFORM eq "amd") {
         $HIPCC="$HIP_CLANG_PATH/clang" . $execExtension;
         $HIPLDFLAGS = "--driver-mode=g++";
     }
-
+    # to avoid using dk linker or MSVC linker
+    if($isWindows) {
+        $HIPLDFLAGS .= " -fuse-ld=lld";
+    }
     $HIP_CLANG_VERSION = `$HIPCC --version`;
     $HIP_CLANG_VERSION=~/.*clang version (\S+).*/;
     $HIP_CLANG_VERSION=$1;

docs/markdown/hip_debugging.md

@@ -261,7 +261,7 @@ The following is the summary of the most useful environment variables in HIP.
 | AMD_SERIALIZE_KERNEL <br><sub> Serialize kernel enqueue. </sub> |  0  | 1: Wait for completion before enqueue. <br> 2: Wait for completion after enqueue. <br> 3: Both. |
 | AMD_SERIALIZE_COPY <br><sub> Serialize copies. </sub> |  0  | 1: Wait for completion before enqueue. <br> 2: Wait for completion after enqueue. <br> 3: Both. |
 | HIP_HOST_COHERENT <br><sub> Coherent memory in hipHostMalloc. </sub> |  0  |  0: memory is not coherent between host and GPU. <br> 1: memory is coherent with host. |
-| AMD_DIRECT_DISPATCH <br><sub> Enable direct kernel dispatch. </sub> | 0  | 0: Disable. <br> 1: Enable. |
+| AMD_DIRECT_DISPATCH <br><sub> Enable direct kernel dispatch. </sub> | 1  | 0: Disable. <br> 1: Enable. |
 
 
 ## General Debugging Tips

docs/markdown/hip_faq.md

@@ -57,7 +57,6 @@ The HIP API documentation describes each API and its limitations, if any, compar
 At a high-level, the following features are not supported:
 - Textures (partial support available)
 - Dynamic parallelism (CUDA 5.0)
-- Managed memory (CUDA 6.5)
 - Graphics interoperability with OpenGL or Direct3D
 - CUDA IPC Functions (Under Development)
 - CUDA array, mipmappedArray and pitched memory
@@ -70,8 +69,7 @@ See the [API Support Table](CUDA_Runtime_API_functions_supported_by_HIP.md) for
 - Virtual functions, indirect functions and try/catch (CUDA 4.0)
 - `__prof_trigger`
 - PTX assembly (CUDA 4.0).  HIP-Clang supports inline GCN assembly.
-- Several kernel features are under development.  See the [HIP Kernel Language](hip_kernel_language.md) for more information.  These include:
-  - printf
+- Several kernel features are under development.  See the [HIP Kernel Language](hip_kernel_language.md) for more information.
 
 
 ### Is HIP a drop-in replacement for CUDA?
@@ -233,6 +231,7 @@ See the [HIP Logging](hip_logging.md) for more information.
 
 ### What is maximum limit of kernel launching parameter?
 Product of block.x, block.y, and block.z should be less than 1024.
+Please note, HIP does not support kernel launch with total work items defined in dimension with size gridDim x blockDim >= 2^32, so gridDim.x * blockDim.x, gridDim.y * blockDim.y and gridDim.z * blockDim.z are always less than 2^32.
 
 ### Are __shfl_*_sync functions supported on HIP platform?
 __shfl_*_sync is not supported on HIP but for nvcc path CUDA 9.0 and above all shuffle calls get redirected to it's sync version.

docs/markdown/hip_kernel_language.md

@@ -184,7 +184,7 @@ The `__restrict__` keyword tells the compiler that the associated memory pointer
 ### Coordinate Built-Ins
 Built-ins determine the coordinate of the active work item in the execution grid. They are defined in amd_hip_runtime.h (rather than being implicitly defined by the compiler).
 In HIP, built-ins coordinate variable definitions are the same as in Cuda, for instance:
-threadIdx.x, blockIdx.y, gridDim.y, etc. 
+threadIdx.x, blockIdx.y, gridDim.y, etc.
 The products gridDim.x * blockDim.x, gridDim.y * blockDim.y and gridDim.z * blockDim.z are always less than 2^32.
 
 ### warpSize
@@ -696,7 +696,7 @@ void assert(int input)
 ```
 
 There are two kinds of implementations for assert functions depending on the use sceneries,
-- One is for the host version of assert, which is defined in assert.h, 
+- One is for the host version of assert, which is defined in assert.h,
 - Another is the device version of assert, which is implemented in hip/hip_runtime.h.
 Users need to include assert.h to use assert. For assert to work in both device and host functions, users need to include "hip/hip_runtime.h".
 

docs/markdown/hip_programming_guide.md

@@ -60,9 +60,10 @@ HIP supports Stream Memory Operations to enable direct synchronization between N
   hipStreamWriteValue64
 
 Note, CPU access to the semaphore's memory requires volatile keyword to disable CPU compiler's optimizations on memory access.
-
 For more details, please check the documentation HIP-API.pdf.
 
+Please note, HIP stream does not gurantee concurrency on AMD hardware for the case of multiple (at least 6) long running streams executing concurrently, using hipStreamSynchronize(nullptr) for synchronization.
+
 ### Coherency Controls
 ROCm defines two coherency options for host memory:
 - Coherent memory : Supports fine-grain synchronization while the kernel is running.  For example, a kernel can perform atomic operations that are visible to the host CPU or to other (peer) GPUs.  Synchronization instructions include threadfence_system and C++11-style atomic operations.
@@ -130,7 +131,10 @@ The link here(https://github.com/ROCm-Developer-Tools/HIP/blob/main/tests/src/hi
 
 ## Device-Side Malloc
 
-HIP-Clang currently doesn't supports device-side malloc and free.
+HIP-Clang now supports device-side malloc and free.
+This implementation does not require the use of `hipDeviceSetLimit(hipLimitMallocHeapSize,value)` nor respects any setting. The heap is fully dynamic and can grow until the available free memory on the device is consumed.
+
+The test codes in the link (https://github.com/ROCm-Developer-Tools/HIP/blob/develop/tests/src/deviceLib/hipDeviceMalloc.cpp) show how to implement application using malloc and free functions in device kernels.
 
 ## Use of Long Double Type
 

include/hip/hip_runtime_api.h

@@ -6177,12 +6177,15 @@ hipError_t hipGraphExecEventWaitNodeSetEvent(hipGraphExec_t hGraphExec, hipGraph
  * Memory allocation properties
  */
 typedef struct hipMemAllocationProp {
-    unsigned char compressionType;                  ///< Compression type
-    hipMemLocation location;                        ///< Memory location
-    hipMemAllocationHandleType requestedHandleType; ///< Requested handle type
     hipMemAllocationType type;                      ///< Memory allocation type
-    unsigned short usage;                           ///< Usage
+    hipMemAllocationHandleType requestedHandleType; ///< Requested handle type
+    hipMemLocation location;                        ///< Memory location
     void* win32HandleMetaData;                      ///< Metadata for Win32 handles
+    struct {
+        unsigned char compressionType;              ///< Compression type
+        unsigned char gpuDirectRDMACapable;         ///< RDMA capable
+        unsigned short usage;                       ///< Usage
+    } allocFlags;
 } hipMemAllocationProp;
 
 /**

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/hipTestMain/config/config_amd_windows.json

@@ -65,7 +65,8 @@
 	   "Unit_hipHostMalloc_NonCoherent",
 	   "Unit_hipHostMalloc_Coherent",
 	   "Unit_hipHostMalloc_Default",
-	   "Unit_hipStreamCreate_MultistreamBasicFunctionalities"
+	   "Unit_hipStreamCreate_MultistreamBasicFunctionalities",
+	   "Unit_hipEventIpc"
 	]
 
 }

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/deviceLib/CMakeLists.txt

@@ -10,17 +10,11 @@ set(TEST_SRC
     popc.cc
     ldg.cc
     threadfence_system.cc
+    syncthreadsand.cc
+    syncthreadscount.cc
+    syncthreadsor.cc
 )
 
-# skipped for windows compiler issue - Illegal instruction detected
-if(UNIX)
-    set(TEST_SRC ${TEST_SRC}
-                 syncthreadsand.cc
-                 syncthreadscount.cc
-                 syncthreadsor.cc)
-endif()
-
-
 # AMD only tests
 set(AMD_TEST_SRC
     unsafeAtomicAdd.cc

tests/catch/unit/event/CMakeLists.txt

@@ -3,15 +3,10 @@ set(TEST_SRC
     Unit_hipEvent_Negative.cc
     Unit_hipEvent.cc
     Unit_hipEventElapsedTime.cc
+    Unit_hipEventRecord.cc
+    Unit_hipEventIpc.cc
 )
 
-# skipped for windows due to duplicate symbols
-if(UNIX)
-    set(TEST_SRC ${TEST_SRC}
-                 Unit_hipEventRecord.cc
-                 Unit_hipEventIpc.cc)
-endif()
-
 hip_add_exe_to_target(NAME EventTest
                       TEST_SRC ${TEST_SRC}
                       TEST_TARGET_NAME build_tests)

tests/catch/unit/stream/CMakeLists.txt

@@ -10,15 +10,9 @@ set(TEST_SRC
     hipStreamGetCUMask.cc
     hipAPIStreamDisable.cc
     streamCommon.cc
-)
-
-#skipped in windows - duplicate HipTest::vector_square sym (compiler issue)
-if(UNIX)
-    set(TEST_SRC ${TEST_SRC}
-                 hipStreamWithCUMask.cc
-                 hipStreamACb_MultiThread.cc)
-endif()
-
+    hipStreamWithCUMask.cc
+    hipStreamACb_MultiThread.cc
+    )
 else()
 set(TEST_SRC
     hipStreamCreate.cc
@@ -31,7 +25,7 @@ set(TEST_SRC
     hipStreamCreateWithPriority.cc
     hipAPIStreamDisable.cc
     streamCommon.cc
-)
+    )
 endif()
 
 hip_add_exe_to_target(NAME StreamTest

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++)
     {