From 56b8b0d80ea87097355f7eee04f2410b2c96e557 Mon Sep 17 00:00:00 2001
From: vsytch <vladislav.sytchenko@amd.com>
Date: Fri, 14 Feb 2020 23:22:25 -0500
Subject: [PATCH] Add missing __hip_pinned_shadow__ attributes to the texture
 global vars. (#1866)

---
 tests/src/texture/hipBindTex2DPitch.cpp          |  4 ++++
 tests/src/texture/hipNormalizedFloatValueTex.cpp |  3 +++
 tests/src/texture/simpleTexture2DLayered.cpp     |  6 ++++++
 tests/src/texture/simpleTexture3D.cpp            | 11 +++++++++++
 4 files changed, 24 insertions(+)

diff --git a/tests/src/texture/hipBindTex2DPitch.cpp b/tests/src/texture/hipBindTex2DPitch.cpp
index 905e24bce3..b01402c91d 100644
--- a/tests/src/texture/hipBindTex2DPitch.cpp
+++ b/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/tests/src/texture/hipNormalizedFloatValueTex.cpp b/tests/src/texture/hipNormalizedFloatValueTex.cpp
index 3179f7412e..609f6916f8 100644
--- a/tests/src/texture/hipNormalizedFloatValueTex.cpp
+++ b/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/tests/src/texture/simpleTexture2DLayered.cpp b/tests/src/texture/simpleTexture2DLayered.cpp
index e3ba6d9afe..fa545cb4ca 100644
--- a/tests/src/texture/simpleTexture2DLayered.cpp
+++ b/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;
diff --git a/tests/src/texture/simpleTexture3D.cpp b/tests/src/texture/simpleTexture3D.cpp
index 741b1671af..06833761c8 100644
--- a/tests/src/texture/simpleTexture3D.cpp
+++ b/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>