From 2b6a2d84fc56971a87b09a1fc3a04f39d0cd71d6 Mon Sep 17 00:00:00 2001
From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com>
Date: Thu, 8 Sep 2022 01:40:20 +0530
Subject: [PATCH] SWDEV-346579 - MI2xx skip build/execution
 hipArrayCreate.cc/hipMallocArray.cc (#2917)

Change-Id: I3a39d541e8ac8b565cccc1b08403e0110edc471d
---
 .../config/config_amd_linux_MI2xx.json        |  8 ++
 .../config/config_amd_windows_MI2xx.json      | 76 +++++++++++++++++++
 catch/unit/memory/hipArrayCommon.hh           | 14 +++-
 3 files changed, 96 insertions(+), 2 deletions(-)
 create mode 100644 catch/hipTestMain/config/config_amd_linux_MI2xx.json
 create mode 100644 catch/hipTestMain/config/config_amd_windows_MI2xx.json

diff --git a/catch/hipTestMain/config/config_amd_linux_MI2xx.json b/catch/hipTestMain/config/config_amd_linux_MI2xx.json
new file mode 100644
index 0000000000..05276e392b
--- /dev/null
+++ b/catch/hipTestMain/config/config_amd_linux_MI2xx.json
@@ -0,0 +1,8 @@
+{
+    "DisabledTests":
+	[
+		"Unit_hipStreamPerThread_DeviceReset_1",
+		"Unit_hipMallocManaged_OverSubscription"
+	]
+
+}
diff --git a/catch/hipTestMain/config/config_amd_windows_MI2xx.json b/catch/hipTestMain/config/config_amd_windows_MI2xx.json
new file mode 100644
index 0000000000..e6b4364a87
--- /dev/null
+++ b/catch/hipTestMain/config/config_amd_windows_MI2xx.json
@@ -0,0 +1,76 @@
+{
+    "DisabledTests":
+	[
+		"Unit_hipMalloc_CoherentTst",
+		"Unit_printf_flags",
+		"Unit_printf_specifier",
+		"Unit_hipTextureMipmapObj2D_Check",
+		"Unit_hipGraphAddHostNode_ClonedGraphwithHostNode",
+		"Unit_hipEventIpc",
+		"Unit_hipMalloc3D_Negative",
+		"Unit_hipPointerGetAttribute_MappedMem",
+		"Unit_hipStreamValue_Write",
+		"Unit_hipMemPoolApi_Basic",
+		"Unit_hipMemPoolApi_BasicAlloc",
+		"Unit_hipMemPoolApi_BasicTrim",
+		"Unit_hipMemPoolApi_BasicReuse",
+		"Unit_hipMemPoolApi_Opportunistic",
+		"Unit_hipMemPoolApi_Default",
+		"Unit_hipDeviceGetUuid",
+		"Unit_hipGraphMemcpyNodeSetParams_Functional",
+		"Unit_hipMalloc3D_ValidatePitch",
+		"Unit_hipArrayCreate_happy",
+		"Unit_hipHostRegister_Negative - int",
+		"Unit_hipHostRegister_Negative - float",
+		"Unit_hipHostRegister_Negative - double",
+		"Unit_hipMemAllocPitch_ValidatePitch",
+		"Unit_hipArrayCreate_happy - int",
+		"Unit_hipArrayCreate_happy - int4",
+		"Unit_hipArrayCreate_happy - short2",
+		"Unit_hipArrayCreate_happy - char",
+		"Unit_hipArrayCreate_happy - char4",
+		"Unit_hipArrayCreate_happy - float",
+		"Unit_hipArrayCreate_happy - float2",
+		"Unit_hipArrayCreate_happy - float4",
+		"Unit_hipMemVmm_Basic",
+		"Unit_hipGraphExecMemcpyNodeSetParamsToSymbol_Functional",
+		"Unit_hipMallocManaged_MultiChunkMultiDevice",
+		"Unit_hipMallocManaged_TwoPointers - int",
+		"Unit_hipMallocManaged_TwoPointers - float",
+		"Unit_hipMallocManaged_TwoPointers - double",
+		"Unit_hipMallocManaged_DeviceContextChange - unsigned char",
+		"Unit_hipMallocManaged_DeviceContextChange - int",
+		"Unit_hipMallocManaged_DeviceContextChange - float",
+		"Unit_hipMallocManaged_DeviceContextChange - double",
+		"Unit_hipGraphNodeGetDependentNodes_Functional",
+		"Unit_hipGraphNodeGetDependentNodes_ParamValidation",
+		"Unit_hipGraphNodeGetDependencies_Functional",
+		"Unit_hipGraphNodeGetDependencies_ParamValidation",
+		"Unit_hipMemGetInfo_DifferentMallocSmall",
+		"Unit_hipMemGetInfo_MallocArray - int",
+		"Unit_hipMemGetInfo_MallocArray - int4",
+		"Unit_hipMemGetInfo_MallocArray - char",
+		"Unit_hipMemGetInfo_Malloc3D",
+		"Unit_hipMemGetInfo_Malloc3DArray - char",
+		"Unit_hipMemGetInfo_Malloc3DArray - int",
+		"Unit_hipMemGetInfo_Malloc3DArray - int4",
+		"Unit_hipMemGetInfo_ParaSmall",
+		"Unit_hipMemGetInfo_ParaMultiSmall",
+		"Unit_hipFreeMultiTDev - char",
+		"Unit_hipFreeMultiTDev - int",
+		"Unit_hipFreeMultiTDev - float2",
+		"Unit_hipFreeMultiTDev - float4",
+		"Unit_hipFreeMultiTHost - char",
+		"Unit_hipFreeMultiTHost - int",
+		"Unit_hipFreeMultiTHost - float2",
+		"Unit_hipFreeMultiTHost - float4",
+		"Unit_hipFreeMultiTArray - char",
+		"Unit_hipFreeMultiTArray - int",
+		"Unit_hipFreeMultiTArray - float2",
+		"Unit_hipFreeMultiTArray - float4",
+		"Unit_hipStreamSynchronize_FinishWork",
+		"Unit_hipStreamSynchronize_NullStreamAndStreamPerThread",
+		"Unit_hipMultiThreadDevice_NearZero",
+		"Unit_hipStreamPerThread_DeviceReset_1"
+	]
+}
diff --git a/catch/unit/memory/hipArrayCommon.hh b/catch/unit/memory/hipArrayCommon.hh
index 8502ac218d..b40014b490 100644
--- a/catch/unit/memory/hipArrayCommon.hh
+++ b/catch/unit/memory/hipArrayCommon.hh
@@ -91,6 +91,7 @@ constexpr size_t ChannelToRead = 1;
 template <typename T>
 __global__ void readFromTexture(T* output, hipTextureObject_t texObj, size_t width, size_t height,
                                 bool textureGather) {
+  #if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
   // Calculate normalized texture coordinates
   const unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
   const unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
@@ -101,9 +102,18 @@ __global__ void readFromTexture(T* output, hipTextureObject_t texObj, size_t wid
     output[x] = tex1D<T>(texObj, u);
   } else {
     const float v = y / (float)height;
-    output[y * width + x] =
-        textureGather ? tex2Dgather<T>(texObj, u, v, ChannelToRead) : tex2D<T>(texObj, u, v);
+    if (textureGather) {
+      // tex2Dgather not supported on __gfx90a__
+      #if !defined(__gfx90a__)
+      output[y * width + x] = tex2Dgather<T>(texObj, u, v, ChannelToRead);
+      #else
+      #warning("tex2Dgather not supported on gfx90a");
+      #endif
+    } else {
+      output[y * width + x] = tex2D<T>(texObj, u, v);
+    }
   }
+  #endif
 }
 
 template <typename T> void checkDataIsAscending(const std::vector<T>& hostData) {