From 07badd6f4e9567f2503abf61431cc0771752f464 Mon Sep 17 00:00:00 2001
From: Maneesh Gupta <maneesh.gupta@amd.com>
Date: Sun, 4 Sep 2016 12:35:08 +0530
Subject: [PATCH] module_api: workaround to use vcpy_kernel.cpp for NV path

Change-Id: Ib4868bf02c64070e846c19427c39289609909466
---
 samples/0_Intro/module_api/Makefile        | 19 ++++++++-----------
 samples/0_Intro/module_api/runKernel.cpp   |  9 ++++++---
 samples/0_Intro/module_api/vcpy_kernel.cpp |  8 +++++++-
 samples/0_Intro/module_api/vcpy_kernel.cu  |  6 ------
 4 files changed, 21 insertions(+), 21 deletions(-)
 delete mode 100644 samples/0_Intro/module_api/vcpy_kernel.cu

diff --git a/samples/0_Intro/module_api/Makefile b/samples/0_Intro/module_api/Makefile
index 2570290f07..bda0a7dd93 100644
--- a/samples/0_Intro/module_api/Makefile
+++ b/samples/0_Intro/module_api/Makefile
@@ -6,13 +6,15 @@ HIPCC=$(HIP_PATH)/bin/hipcc
 OPT=
 HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --compiler)
 
-ifeq (${HIP_PLATFORM}, hcc)
-	GENCODEOBJECT_FLAGS=--target-isa=fiji
+all: vcpy_kernel.compile runKernel.hip.out
 
-all: runKernel.hip.out
+runKernel.hip.out: runKernel.cpp
+	$(HIPCC) $(OPT) runKernel.cpp -o runKernel.hip.out
+
+ifeq (${HIP_PLATFORM}, hcc)
 
 vcpy_kernel.compile: vcpy_kernel.cpp
-	$(HIPCC) --genco $(GENCODEOBJECT_FLAGS) vcpy_kernel.cpp -o vcpy_kernel.co
+	$(HIPCC) --genco --target-isa=fiji vcpy_kernel.cpp -o vcpy_kernel.co
 
 clean:
 	rm -f *.co *.out
@@ -21,16 +23,11 @@ endif
 
 ifeq (${HIP_PLATFORM}, nvcc)
 
-vcpy_kernel.compile: vcpy_kernel.cu
-	$(HIPCC) --genco vcpy_kernel.cu -o vcpy_kernel.ptx
+vcpy_kernel.compile: vcpy_kernel.cpp
+	$(HIPCC) --genco vcpy_kernel.cpp -o vcpy_kernel.ptx
 
 clean:
 	rm -f *.ptx *.out
 
 endif
 
-all: vcpy_kernel.compile runKernel.hip.out
-
-runKernel.hip.out: runKernel.cpp
-	$(HIPCC) $(OPT) runKernel.cpp -o runKernel.hip.out
-
diff --git a/samples/0_Intro/module_api/runKernel.cpp b/samples/0_Intro/module_api/runKernel.cpp
index 1156b8ddbf..ea372773c3 100644
--- a/samples/0_Intro/module_api/runKernel.cpp
+++ b/samples/0_Intro/module_api/runKernel.cpp
@@ -83,9 +83,12 @@ int main(){
 #endif
 
 #ifdef __HIP_PLATFORM_NVCC__
-    std::vector<void*>argBuffer(2);
-    memcpy(&argBuffer[0], &Ad, sizeof(void*));
-    memcpy(&argBuffer[1], &Bd, sizeof(void*));
+	uint32_t one = 1;
+    std::vector<void*>argBuffer(3);
+    uint32_t *ptr32_t = (uint32_t*)&argBuffer[0];
+    memcpy(ptr32_t + 0, &one, sizeof(uint32_t));
+    memcpy(&argBuffer[1], &Ad, sizeof(void*));
+    memcpy(&argBuffer[2], &Bd, sizeof(void*));
 #endif
 
 
diff --git a/samples/0_Intro/module_api/vcpy_kernel.cpp b/samples/0_Intro/module_api/vcpy_kernel.cpp
index 640cf5b1bb..ebb6d066a9 100644
--- a/samples/0_Intro/module_api/vcpy_kernel.cpp
+++ b/samples/0_Intro/module_api/vcpy_kernel.cpp
@@ -1,6 +1,12 @@
 #include<hip_runtime.h>
 
-__global__ void hello_world(hipLaunchParm lp, float *a, float *b)
+#ifdef __HIP_PLATFORM_NVCC__
+#define EXTERN_C extern "C"
+#else
+#define EXTERN_C
+#endif
+
+EXTERN_C __global__ void hello_world(hipLaunchParm lp, float *a, float *b)
 {
     int tx = hipThreadIdx_x;
     b[tx] = a[tx];
diff --git a/samples/0_Intro/module_api/vcpy_kernel.cu b/samples/0_Intro/module_api/vcpy_kernel.cu
deleted file mode 100644
index ead7d70311..0000000000
--- a/samples/0_Intro/module_api/vcpy_kernel.cu
+++ /dev/null
@@ -1,6 +0,0 @@
-extern "C" __global__ void hello_world(float *a, float *b)
-{
-    int tx = threadIdx.x;
-    b[tx] = a[tx];
-
-}