module_api: workaround to use vcpy_kernel.cpp for NV path

Change-Id: Ib4868bf02c64070e846c19427c39289609909466

hipamd/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
-

hipamd/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
 
 

hipamd/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];

hipamd/samples/0_Intro/module_api/vcpy_kernel.cu

@@ -1,6 +0,0 @@
-extern "C" __global__ void hello_world(float *a, float *b)
-{
-    int tx = threadIdx.x;
-    b[tx] = a[tx];
-
-}