From 8bbd449f9fb0031ca81f65e8dd814ccbaefa4e30 Mon Sep 17 00:00:00 2001
From: Aditya Atluri <Aditya.Atluri@amd.com>
Date: Thu, 1 Sep 2016 10:39:14 -0500
Subject: [PATCH] Fixed offline kernel compilation

1. Removed vcpy_isa.ptx as it should be generated during make
2. Made argument padding specific to hcc path
3. Renamed --gencodeobject to --genco
4. Changed Makefile to work on both nvcc and hcc path

Change-Id: Ifd053d541085d9ce4fd37bc21b07674786c7163e


[ROCm/clr commit: f22fda129111152362c9dbb94da357511ae020c9]
---
 projects/clr/hipamd/bin/hipcc                 |  4 +-
 .../samples/0_Intro/module_api/Makefile       | 24 +++++++++---
 .../samples/0_Intro/module_api/runKernel.cpp  | 11 +++++-
 .../samples/0_Intro/module_api/vcpy_isa.ptx   | 38 -------------------
 4 files changed, 29 insertions(+), 48 deletions(-)
 delete mode 100644 projects/clr/hipamd/samples/0_Intro/module_api/vcpy_isa.ptx

diff --git a/projects/clr/hipamd/bin/hipcc b/projects/clr/hipamd/bin/hipcc
index 8fde7f426d..422dba4e54 100755
--- a/projects/clr/hipamd/bin/hipcc
+++ b/projects/clr/hipamd/bin/hipcc
@@ -181,7 +181,7 @@ my $ISACMD="";
 if($HIP_PLATFORM eq "hcc"){
   $ISACMD .= "$HIP_PATH/bin/hipgenisa.sh ";
   $ISACMD .= $ROCM_PATH;
-  if($ARGV[0] eq "--gencodeobject"){
+  if($ARGV[0] eq "--genco"){
     foreach $isaarg (@ARGV[1..$#ARGV]){
       $ISACMD .= " ";
       $ISACMD .= $isaarg;
@@ -196,7 +196,7 @@ if($HIP_PLATFORM eq "hcc"){
 
 if($HIP_PLATFORM eq "nvcc"){
   $ISACMD .= "$HIP_PATH/bin/hipcc -ptx ";
-  if($ARGV[0] eq "--gencodeobject"){
+  if($ARGV[0] eq "--genco"){
     foreach $isaarg (@ARGV[1..$#ARGV]){
       $ISACMD .= " ";
       $ISACMD .= $isaarg;
diff --git a/projects/clr/hipamd/samples/0_Intro/module_api/Makefile b/projects/clr/hipamd/samples/0_Intro/module_api/Makefile
index 25ae7b8411..db270beaa0 100644
--- a/projects/clr/hipamd/samples/0_Intro/module_api/Makefile
+++ b/projects/clr/hipamd/samples/0_Intro/module_api/Makefile
@@ -6,16 +6,28 @@ HIPCC=$(HIP_PATH)/bin/hipcc
 HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --compiler)
 
 ifeq (${HIP_PLATFORM}, hcc)
-	GENCODEOBJECT_FLAGS=--target-isa-fiji
+	GENCODEOBJECT_FLAGS=--target-isa=fiji
+
+vcpy_isa.compile: vcpy_isa.cpp
+	$(HIPCC) --genco $(GENCODEOBJECT_FLAGS) vcpy_isa.cpp -o vcpy_isa.co
+
+clean:
+	rm -f *.co *.out
+
+endif
+
+ifeq (${HIP_PLATFORM}, nvcc)
+
+vcpy_isa.compile: vcpy_isa.cu
+	$(HIPCC) --genco vcpy_isa.cu -o vcpy_isa.ptx
+
+clean:
+	rm -f *.ptx *.out
+
 endif
 
 all: vcpy_isa.compile runKernel.hip.out
 
-vcpy_isa.compile: vcpy_isa.cpp
-	$(HIPCC) --gencodeobject $(GENCODEOBJECT_FLAGS) vcpy_isa.cpp -o vcpy_isa.co
-
 runKernel.hip.out: runKernel.cpp
 	$(HIPCC) runKernel.cpp -o runKernel.hip.out
 
-clean:
-	rm -f *.co *.out
diff --git a/projects/clr/hipamd/samples/0_Intro/module_api/runKernel.cpp b/projects/clr/hipamd/samples/0_Intro/module_api/runKernel.cpp
index e4fa1b6d93..dcbaa4cd35 100644
--- a/projects/clr/hipamd/samples/0_Intro/module_api/runKernel.cpp
+++ b/projects/clr/hipamd/samples/0_Intro/module_api/runKernel.cpp
@@ -66,8 +66,9 @@ int main(){
     hipModuleLoad(&Module, fileName);
     hipModuleGetFunction(&Function, Module, kernel_name);
 
-  uint32_t len = LEN;
-  uint32_t one = 1;
+#ifdef __HIP_PLATFORM_HCC__
+    uint32_t len = LEN;
+    uint32_t one = 1;
 
     std::vector<void*>argBuffer(5);
     uint32_t *ptr32_t = (uint32_t*)&argBuffer[0];
@@ -79,7 +80,13 @@ int main(){
     memcpy(ptr32_t + 5, &one, sizeof(uint32_t));
     memcpy(&argBuffer[3], &Ad, sizeof(void*));
     memcpy(&argBuffer[4], &Bd, sizeof(void*));
+#endif
 
+#ifdef __HIP_PLATFORM_NVCC__
+    std::vector<void*>argBuffer(2);
+    memcpy(&argBuffer[0], &Ad, sizeof(void*));
+    memcpy(&argBuffer[1], &Bd, sizeof(void*));
+#endif
 
 
     size_t size = argBuffer.size()*sizeof(void*);
diff --git a/projects/clr/hipamd/samples/0_Intro/module_api/vcpy_isa.ptx b/projects/clr/hipamd/samples/0_Intro/module_api/vcpy_isa.ptx
deleted file mode 100644
index 62eb3f63df..0000000000
--- a/projects/clr/hipamd/samples/0_Intro/module_api/vcpy_isa.ptx
+++ /dev/null
@@ -1,38 +0,0 @@
-//
-// Generated by NVIDIA NVVM Compiler
-//
-// Compiler Build ID: CL-19856038
-// Cuda compilation tools, release 7.5, V7.5.17
-// Based on LLVM 3.4svn
-//
-
-.version 4.3
-.target sm_20
-.address_size 64
-
-	// .globl	hello_world
-
-.visible .entry hello_world(
-	.param .u64 hello_world_param_0,
-	.param .u64 hello_world_param_1
-)
-{
-	.reg .f32 	%f<2>;
-	.reg .b32 	%r<2>;
-	.reg .b64 	%rd<8>;
-
-
-	ld.param.u64 	%rd1, [hello_world_param_0];
-	ld.param.u64 	%rd2, [hello_world_param_1];
-	cvta.to.global.u64 	%rd3, %rd2;
-	cvta.to.global.u64 	%rd4, %rd1;
-	mov.u32 	%r1, %tid.x;
-	mul.wide.s32 	%rd5, %r1, 4;
-	add.s64 	%rd6, %rd4, %rd5;
-	ld.global.f32 	%f1, [%rd6];
-	add.s64 	%rd7, %rd3, %rd5;
-	st.global.f32 	[%rd7], %f1;
-	ret;
-}
-
-