From 6ca7a87e0e733ab0459399c9e979e86ceb1e2e4d Mon Sep 17 00:00:00 2001
From: Aditya Atluri <Aditya.Atluri@amd.com>
Date: Fri, 2 Sep 2016 13:17:17 -0500
Subject: [PATCH] corrected offline kernel compilation on hipcc path

1. hipgenisa.sh now adds int main(){} during kernel compilation. User does not have to put it there
2. Renamed vcpy_isa.cpp to vcpy_kernel.cpp
3. Removed vcpy_isa.cu as the kernel code should be common for both paths
4. Changed Makefile and runkernel.cpp to work with above changes

Change-Id: I9f8c84706b44bb500bc493a68e959762b55a0142
---
 samples/0_Intro/module_api/Makefile                    | 10 +++++-----
 samples/0_Intro/module_api/runKernel.cpp               |  8 ++++----
 samples/0_Intro/module_api/vcpy_isa.cu                 |  6 ------
 .../module_api/{vcpy_isa.cpp => vcpy_kernel.cpp}       |  1 -
 4 files changed, 9 insertions(+), 16 deletions(-)
 delete mode 100644 samples/0_Intro/module_api/vcpy_isa.cu
 rename samples/0_Intro/module_api/{vcpy_isa.cpp => vcpy_kernel.cpp} (91%)

diff --git a/samples/0_Intro/module_api/Makefile b/samples/0_Intro/module_api/Makefile
index 8981938e91..99cff6bc0a 100644
--- a/samples/0_Intro/module_api/Makefile
+++ b/samples/0_Intro/module_api/Makefile
@@ -11,8 +11,8 @@ ifeq (${HIP_PLATFORM}, hcc)
 
 all: runKernel.hip.out
 
-vcpy_isa.compile: vcpy_isa.cpp
-	$(HIPCC) --genco $(GENCODEOBJECT_FLAGS) vcpy_isa.cpp -o vcpy_isa.co
+vcpy_kernel.compile: vcpy_kernel.cpp
+	$(HIPCC) --genco $(GENCODEOBJECT_FLAGS) vcpy_kernel.cpp -o vcpy_kernel.co
 
 clean:
 	rm -f *.co *.out
@@ -21,15 +21,15 @@ endif
 
 ifeq (${HIP_PLATFORM}, nvcc)
 
-vcpy_isa.compile: vcpy_isa.cu
-	$(HIPCC) --genco vcpy_isa.cu -o vcpy_isa.ptx
+vcpy_kernel.compile: vcpy_kernel.cpp
+	$(HIPCC) --genco vcpy_kernel.cpp -o vcpy_kernel.ptx
 
 clean:
 	rm -f *.ptx *.out
 
 endif
 
-all: vcpy_isa.compile runKernel.hip.out
+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 dcbaa4cd35..1156b8ddbf 100644
--- a/samples/0_Intro/module_api/runKernel.cpp
+++ b/samples/0_Intro/module_api/runKernel.cpp
@@ -27,12 +27,12 @@ THE SOFTWARE.
 #define SIZE LEN<<2
 
 #ifdef __HIP_PLATFORM_HCC__ 
-#define fileName "vcpy_isa.co"
+#define fileName "vcpy_kernel.co"
 #define kernel_name "ZN12_GLOBAL__N_146_Z11hello_world16grid_launch_parmPfS0__functor19__cxxamp_trampolineEiiiiiiPKfPf"
 #endif
 
 #ifdef __HIP_PLATFORM_NVCC__
-#define fileName "vcpy_isa.ptx"
+#define fileName "vcpy_kernel.ptx"
 #define kernel_name "hello_world"
 #endif
 
@@ -67,8 +67,8 @@ int main(){
     hipModuleGetFunction(&Function, Module, kernel_name);
 
 #ifdef __HIP_PLATFORM_HCC__
-    uint32_t len = LEN;
-    uint32_t one = 1;
+		uint32_t len = LEN;
+		uint32_t one = 1;
 
     std::vector<void*>argBuffer(5);
     uint32_t *ptr32_t = (uint32_t*)&argBuffer[0];
diff --git a/samples/0_Intro/module_api/vcpy_isa.cu b/samples/0_Intro/module_api/vcpy_isa.cu
deleted file mode 100644
index d2a0838604..0000000000
--- a/samples/0_Intro/module_api/vcpy_isa.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];
-}
diff --git a/samples/0_Intro/module_api/vcpy_isa.cpp b/samples/0_Intro/module_api/vcpy_kernel.cpp
similarity index 91%
rename from samples/0_Intro/module_api/vcpy_isa.cpp
rename to samples/0_Intro/module_api/vcpy_kernel.cpp
index ead3253115..640cf5b1bb 100644
--- a/samples/0_Intro/module_api/vcpy_isa.cpp
+++ b/samples/0_Intro/module_api/vcpy_kernel.cpp
@@ -6,4 +6,3 @@ __global__ void hello_world(hipLaunchParm lp, float *a, float *b)
     b[tx] = a[tx];
 }
 
-int main(){}