From 37745e6f1b83e595a1fd6f01e91a5c3a07bb4a41 Mon Sep 17 00:00:00 2001
From: agodavar <anusha.godavarthysurya@amd.com>
Date: Wed, 19 Feb 2020 17:31:12 +0530
Subject: [PATCH] HIP-HCC + Clang fix failing texture tests(SWDEV-223622)

Change-Id: I1c0a7059c6bba1ea2e39dc6634b32d006c07d1f5


[ROCm/clr commit: 63c1d28a4b6292fe649f923edac577acc6d5a8b8]
---
 .../hipamd/tests/src/texture/hipBindTex2DPitch.cpp    |  4 ++++
 .../tests/src/texture/hipNormalizedFloatValueTex.cpp  |  3 +++
 .../tests/src/texture/simpleTexture2DLayered.cpp      |  8 +++++++-
 .../clr/hipamd/tests/src/texture/simpleTexture3D.cpp  | 11 +++++++++++
 4 files changed, 25 insertions(+), 1 deletion(-)

diff --git a/projects/clr/hipamd/tests/src/texture/hipBindTex2DPitch.cpp b/projects/clr/hipamd/tests/src/texture/hipBindTex2DPitch.cpp
index 772d6b6050..8c57520c00 100644
--- a/projects/clr/hipamd/tests/src/texture/hipBindTex2DPitch.cpp
+++ b/projects/clr/hipamd/tests/src/texture/hipBindTex2DPitch.cpp
@@ -27,6 +27,10 @@ THE SOFTWARE.
 #define SIZE_H 8
 #define SIZE_W 12
 #define TYPE_t float
+
+#if __HIP__
+__hip_pinned_shadow__
+#endif
 texture<TYPE_t, 2, hipReadModeElementType> tex;
 
 // texture object is a kernel argument
diff --git a/projects/clr/hipamd/tests/src/texture/hipNormalizedFloatValueTex.cpp b/projects/clr/hipamd/tests/src/texture/hipNormalizedFloatValueTex.cpp
index 03af122696..f8526eb145 100644
--- a/projects/clr/hipamd/tests/src/texture/hipNormalizedFloatValueTex.cpp
+++ b/projects/clr/hipamd/tests/src/texture/hipNormalizedFloatValueTex.cpp
@@ -45,6 +45,9 @@ static float getNormalizedValue(const float value,
     }
 }
 
+#if __HIP__
+__hip_pinned_shadow__
+#endif
 texture<float, hipTextureType1D, hipReadModeElementType> textureNormalizedVal_1D;
 
 __global__ void normalizedValTextureTest(unsigned int numElements, float* pDst)
diff --git a/projects/clr/hipamd/tests/src/texture/simpleTexture2DLayered.cpp b/projects/clr/hipamd/tests/src/texture/simpleTexture2DLayered.cpp
index 186c9995d9..f4d3aac1e5 100644
--- a/projects/clr/hipamd/tests/src/texture/simpleTexture2DLayered.cpp
+++ b/projects/clr/hipamd/tests/src/texture/simpleTexture2DLayered.cpp
@@ -26,9 +26,15 @@ THE SOFTWARE.
  * HIT_END
  */
 #include "test_common.h"
+
 typedef float T;
+
 // Texture reference for 2D Layered texture
+#if __HIP__
+__hip_pinned_shadow__
+#endif
 texture<float, hipTextureType2DLayered> tex2DL;
+
 __global__ void simpleKernelLayeredArray(T* outputData,int width,int height,int layer)
 {
     unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
@@ -59,7 +65,7 @@ void runTest(int width,int height,int num_layers,texture<T, hipTextureType2DLaye
     myparms.dstPos = make_hipPos(0,0,0);
     myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height);
     myparms.dstArray = arr;
-    myparms.extent = make_hipExtent(width, height, num_layers);
+    myparms.extent = make_hipExtent(width , height, num_layers);
     //myparms.kind = hipMemcpyHostToDevice;
     HIPCHECK(hipMemcpy3D(&myparms));
 
diff --git a/projects/clr/hipamd/tests/src/texture/simpleTexture3D.cpp b/projects/clr/hipamd/tests/src/texture/simpleTexture3D.cpp
index 1b3cc49782..673e373dcf 100644
--- a/projects/clr/hipamd/tests/src/texture/simpleTexture3D.cpp
+++ b/projects/clr/hipamd/tests/src/texture/simpleTexture3D.cpp
@@ -31,8 +31,19 @@ THE SOFTWARE.
 const char *sampleName = "simpleTexture3D";
 
 // Texture reference for 3D texture
+#if __HIP__
+__hip_pinned_shadow__
+#endif
 texture<float, hipTextureType3D, hipReadModeElementType> texf;
+
+#if __HIP__
+__hip_pinned_shadow__
+#endif
 texture<int, hipTextureType3D, hipReadModeElementType> texi;
+
+#if __HIP__
+__hip_pinned_shadow__
+#endif
 texture<char, hipTextureType3D, hipReadModeElementType> texc;
 
 template <typename T>