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


[ROCm/clr commit: 8f536d8580]

projects/clr/hipamd/bin/hipgenisa.sh

@@ -22,10 +22,15 @@ HIP_PATH="$( command cd -P "$( dirname "$SOURCE" )/.." && pwd )"
 export KMDUMPISA=1
 export KMDUMPLLVM=1
 mkdir /tmp/hipgenisa
-$HIP_PATH/bin/hipcc $FILE_NAMES -o /tmp/hipgenisa/a.out
+cp $FILE_NAMES $FILE_NAMES.kernel.tmp.cpp
+echo "
+int main(){}
+" >> $FILE_NAMES.kernel.tmp.cpp
+$HIP_PATH/bin/hipcc $FILE_NAMES.kernel.tmp.cpp -o /tmp/hipgenisa/a.out
 mv dump.* /tmp/hipgenisa/
 $ROCM_PATH/hcc-lc/bin/llvm-mc -arch=amdgcn -mcpu=$TARGET -filetype=obj /tmp/hipgenisa/dump.isa -o /tmp/hipgenisa/dump.o
 $ROCM_PATH/llvm/bin/clang -target amdgcn--amdhsa /tmp/hipgenisa/dump.o -o $OUTPUT_FILE
+rm $FILE_NAMES.kernel.tmp.cpp
 rm -r /tmp/hipgenisa
 export KMDUMPISA=0
 export KMDUMPLLVM=0

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

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

projects/clr/hipamd/samples/0_Intro/module_api/vcpy_isa.cu

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

projects/clr/hipamd/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(){}