From 3169f3a8811d0d03fb95297a4a2c541bca3ffc37 Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu@amd.com>
Date: Wed, 3 Jul 2019 04:51:59 -0400
Subject: [PATCH] Add __hip_pinned_shadow__ for hip-clang and fix texture
 reference tests (#1200)

---
 hipamd/include/hip/hcc_detail/host_defines.h      | 1 +
 hipamd/tests/src/texture/hipBindTexRef1DFetch.cpp | 3 +++
 hipamd/tests/src/texture/hipTextureRef2D.cpp      | 3 +++
 3 files changed, 7 insertions(+)

diff --git a/hipamd/include/hip/hcc_detail/host_defines.h b/hipamd/include/hip/hcc_detail/host_defines.h
index 689a4530da..11bd577f08 100644
--- a/hipamd/include/hip/hcc_detail/host_defines.h
+++ b/hipamd/include/hip/hcc_detail/host_defines.h
@@ -72,6 +72,7 @@ THE SOFTWARE.
 
 #define __noinline__ __attribute__((noinline))
 #define __forceinline__ inline __attribute__((always_inline))
+#define __hip_pinned_shadow__ __attribute__((hip_pinned_shadow))
 
 #else
 
diff --git a/hipamd/tests/src/texture/hipBindTexRef1DFetch.cpp b/hipamd/tests/src/texture/hipBindTexRef1DFetch.cpp
index 546e9775e8..52a0d99ac1 100644
--- a/hipamd/tests/src/texture/hipBindTexRef1DFetch.cpp
+++ b/hipamd/tests/src/texture/hipBindTexRef1DFetch.cpp
@@ -32,6 +32,9 @@ THE SOFTWARE.
 
 #define N 512
 
+#if __HIP__
+__hip_pinned_shadow__
+#endif
 texture<float, 1, hipReadModeElementType> tex;
 
 __global__ void kernel(float *out) {
diff --git a/hipamd/tests/src/texture/hipTextureRef2D.cpp b/hipamd/tests/src/texture/hipTextureRef2D.cpp
index 2b19d1b5b3..2fb4097b0a 100644
--- a/hipamd/tests/src/texture/hipTextureRef2D.cpp
+++ b/hipamd/tests/src/texture/hipTextureRef2D.cpp
@@ -9,6 +9,9 @@
 
 #include <hip/hip_runtime.h>
 #include "test_common.h"
+#if __HIP__
+__hip_pinned_shadow__
+#endif
 texture<float, 2, hipReadModeElementType> tex;
 
 __global__ void tex2DKernel(float* outputData,