From 688ce62b493b9afa9ca84d9278e3fcf459dda5a6 Mon Sep 17 00:00:00 2001
From: Yaxun Sam Liu <yaxun.liu@amd.com>
Date: Tue, 2 Jul 2019 16:36:53 -0400
Subject: [PATCH] Fix sample module_api_global for hip-clang

module_api_global relies on a HCC only feature which allows host code
to write to device variables. This feature does not exist in CUDA
or hip-clang, which causes the sample not working in CUDA or hip-clang.

This patch fixes the sample by using standard features of CUDA and
hip-clang. The fixed sample works in HCC, CUDA and hip-clang.
---
 samples/0_Intro/module_api_global/runKernel.cpp   | 8 +++++---
 samples/0_Intro/module_api_global/vcpy_kernel.cpp | 3 +--
 2 files changed, 6 insertions(+), 5 deletions(-)

diff --git a/samples/0_Intro/module_api_global/runKernel.cpp b/samples/0_Intro/module_api_global/runKernel.cpp
index 1e10dba5c6..7031b9f898 100644
--- a/samples/0_Intro/module_api_global/runKernel.cpp
+++ b/samples/0_Intro/module_api_global/runKernel.cpp
@@ -31,7 +31,6 @@ THE SOFTWARE.
 #define SIZE LEN * sizeof(float)
 
 #define fileName "vcpy_kernel.code"
-float myDeviceGlobalArray[16];
 #define HIP_CHECK(cmd)                                                                             \
     {                                                                                              \
         hipError_t status = cmd;                                                                   \
@@ -71,14 +70,17 @@ int main() {
     float* deviceGlobal;
     size_t deviceGlobalSize;
     HIP_CHECK(hipModuleGetGlobal((void**)&deviceGlobal, &deviceGlobalSize, Module, "myDeviceGlobal"));
-    *deviceGlobal = 42.0;
+    HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(deviceGlobal), &myDeviceGlobal_h, deviceGlobalSize));
 
 #define ARRAY_SIZE 16
 
     float myDeviceGlobalArray_h[ARRAY_SIZE];
+    float *myDeviceGlobalArray;
+    size_t myDeviceGlobalArraySize;
+    HIP_CHECK(hipModuleGetGlobal((void**)&myDeviceGlobalArray, &myDeviceGlobalArraySize, Module, "myDeviceGlobalArray"));
     for (int i = 0; i < ARRAY_SIZE; i++) {
         myDeviceGlobalArray_h[i] = i * 1000.0f;
-        myDeviceGlobalArray[i] = i * 1000.0f;
+        HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(myDeviceGlobalArray), &myDeviceGlobalArray_h, myDeviceGlobalArraySize));
     }
 
     struct {
diff --git a/samples/0_Intro/module_api_global/vcpy_kernel.cpp b/samples/0_Intro/module_api_global/vcpy_kernel.cpp
index 662dcd0cc3..9e8d7c9e29 100644
--- a/samples/0_Intro/module_api_global/vcpy_kernel.cpp
+++ b/samples/0_Intro/module_api_global/vcpy_kernel.cpp
@@ -25,8 +25,7 @@ THE SOFTWARE.
 #define ARRAY_SIZE (16)
 
 __device__ float myDeviceGlobal;
-extern float myDeviceGlobalArray[16];
-;
+__device__ float myDeviceGlobalArray[16];
 
 extern "C" __global__ void hello_world(const float* a, float* b) {
     int tx = hipThreadIdx_x;