From 219343027fdf45bcb86a0070665cb58f54502c7f Mon Sep 17 00:00:00 2001
From: Aditya Atluri <Aditya.Atluri@amd.com>
Date: Fri, 17 Mar 2017 13:11:34 -0500
Subject: [PATCH] Added default module launch api functionality

    1. As in hipModuleLaunchKernel(..., kernelParams, nullptr); works with this commit
    2. Added headers AMDGPUPTNote.h, AMDGPURuntimeMetadata.h to do code object meta data parsing
    3. Changed CMake to look at llvm link libraries
    4. HIP developer should set env variable LLVM_HOME to remove link errors
    5. HIP depends on installed LLVM (not source, not build)
    6. Added sample to test out the feature
    7. Right now HCC does not support embedding metadata in code object. Use clang opencl
    8. Changed HIPCC to read LLVM_HOME env var
    9. New argument to CMake should be given -DLLVM_HOME=<where llvm 5.0 is installed>

Change-Id: Iba38194aa872d97cc2c90a8e5ff746c48055c868
---
 samples/0_Intro/module_api/Makefile          |   8 +-
 samples/0_Intro/module_api/defaultDriver.cpp |  89 +++++++++++++++++++
 samples/0_Intro/module_api/runKernel.cpp     |   4 +-
 samples/0_Intro/module_api/test.cl           |  12 +++
 samples/0_Intro/module_api/test.co           | Bin 0 -> 9824 bytes
 5 files changed, 108 insertions(+), 5 deletions(-)
 create mode 100644 samples/0_Intro/module_api/defaultDriver.cpp
 create mode 100644 samples/0_Intro/module_api/test.cl
 create mode 100755 samples/0_Intro/module_api/test.co

diff --git a/samples/0_Intro/module_api/Makefile b/samples/0_Intro/module_api/Makefile
index f2c0ce555a..632a8d3e70 100644
--- a/samples/0_Intro/module_api/Makefile
+++ b/samples/0_Intro/module_api/Makefile
@@ -5,14 +5,16 @@ endif
 HIPCC=$(HIP_PATH)/bin/hipcc
 HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --compiler)
 
-all: vcpy_kernel.code runKernel.hip.out
+all: vcpy_kernel.code runKernel.hip.out defaultDriver.hip.out
 
 runKernel.hip.out: runKernel.cpp
 	$(HIPCC) $(HIPCC_FLAGS) $< -o $@
 
+defaultDriver.hip.out: defaultDriver.cpp
+	$(HIPCC)  $(HIPCC_FLAGS) $< -o $@
+
 vcpy_kernel.code: vcpy_kernel.cpp
-	$(HIPCC) --genco $(GENCO_FLAGS) $^ -o $@
+	$(HIPCC)  --genco  $(GENCO_FLAGS) $^ -o $@
 
 clean:
 	rm -f *.code *.out
-
diff --git a/samples/0_Intro/module_api/defaultDriver.cpp b/samples/0_Intro/module_api/defaultDriver.cpp
new file mode 100644
index 0000000000..a29271a0dc
--- /dev/null
+++ b/samples/0_Intro/module_api/defaultDriver.cpp
@@ -0,0 +1,89 @@
+/*
+Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include "hip/hip_runtime.h"
+#include "hip/hip_runtime_api.h"
+#include<iostream>
+#include<fstream>
+#include<vector>
+
+#define LEN 64
+#define SIZE LEN<<2
+
+#define fileName "test.co"
+#define kernel_name "vadd"
+
+int main(){
+    float *A, *B, *C;
+    hipDeviceptr_t Ad, Bd, Cd;
+    A = new float[LEN];
+    B = new float[LEN];
+    C = new float[LEN];
+
+    for(uint32_t i=0;i<LEN;i++){
+        A[i] = i*1.0f;
+        B[i] = 1.0f;
+        C[i] = 0.0f;
+    }
+
+    hipInit(0);
+    hipDevice_t device;
+    hipCtx_t context;
+    hipDeviceGet(&device, 0);
+    hipCtxCreate(&context, 0, device);
+
+    hipMalloc((void**)&Ad, SIZE);
+    hipMalloc((void**)&Bd, SIZE);
+    hipMalloc((void**)&Cd, SIZE);
+
+    hipMemcpyHtoD(Ad, A, SIZE);
+    hipMemcpyHtoD(Bd, B, SIZE);
+    hipMemcpyHtoD(Cd, C, SIZE);
+
+    hipModule_t Module;
+    hipFunction_t Function;
+    hipModuleLoad(&Module, fileName);
+    hipModuleGetFunction(&Function, Module, kernel_name);
+
+    int n = LEN;
+    void * args[4] = {&Ad, &Bd, &Cd, &n};
+
+    hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, args, nullptr);
+
+    hipMemcpyDtoH(C, Cd, SIZE);
+    int mismatchCount = 0;
+    for(uint32_t i=0;i<LEN;i++){
+        if (A[i] + B[i] != C[i]) {
+            mismatchCount++;
+            std::cout<<"error: mismatch " << A[i]<<" + "<<B[i]<<" != "<<C[i]<<std::endl;
+        }
+    }
+
+    if (mismatchCount == 0) {
+        std::cout << "PASSED!\n";
+    } else {
+        std::cout << "FAILED!\n";
+    };
+
+    hipCtxDestroy(context);
+    return 0;
+}
diff --git a/samples/0_Intro/module_api/runKernel.cpp b/samples/0_Intro/module_api/runKernel.cpp
index b91507aaa4..201892a4a1 100644
--- a/samples/0_Intro/module_api/runKernel.cpp
+++ b/samples/0_Intro/module_api/runKernel.cpp
@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
 
 Permission is hereby granted, free of charge, to any person obtaining a copy
 of this software and associated documentation files (the "Software"), to deal
@@ -98,7 +98,7 @@ int main(){
       HIP_LAUNCH_PARAM_END
     };
 
-    hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config); 
+    hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config);
 
     hipMemcpyDtoH(B, Bd, SIZE);
     int mismatchCount = 0;
diff --git a/samples/0_Intro/module_api/test.cl b/samples/0_Intro/module_api/test.cl
new file mode 100644
index 0000000000..81b20cab0e
--- /dev/null
+++ b/samples/0_Intro/module_api/test.cl
@@ -0,0 +1,12 @@
+__kernel void memset(char in, __global int* out) {
+int tx = get_global_id(0);
+out[tx] = in;
+}
+
+
+__kernel void vadd(__global float *Ad, __global float *Bd, __global float *Cd, int N){
+int tx = get_global_id(0);
+if(tx < N){
+Cd[tx] = Ad[tx] + Bd[tx];
+}
+}
diff --git a/samples/0_Intro/module_api/test.co b/samples/0_Intro/module_api/test.co
new file mode 100755
index 0000000000000000000000000000000000000000..a3e6b991e39e55b9967ae08c86be94a583ce0dbe
GIT binary patch
literal 9824
zcmeGi+in|0bi8q#khCh|B5|uUq7Q9R;_R+>H@ikcag#Q!>Xh`B3Zc@Bch`yO`hvZV
z8c^LpLG*zK-~;#qzJcbkyp=xZp)W{%;2VH*c4oX@lU6MRTGTy~_spEjoHOT~nZ1l}
z950_xGMPCdV1(Q!8En7E33D8g)kC5IXpY!q9IpGw1j#`^8__s_jw9f5*p3DxMB;Hl
zM>6U?3@%S_j>OU*)Eh;X@NXF38R2>oOFrtMuY}?SfJWCoUKL%5rJe%xC()DhJNIko
zMf0O#;32#!$kGn_hj143B7401*jKa_*+YjBhVX+iPL^249x?ii_Df3(ks7v<{wDk|
zoJh}-u}BWOeHzdgmoIYJ?_qo#2W6+Hrzbw6tnTHP7FPUrr`Bu~>4#LGq8hzC6){eA
z*r`UfS)^P}bT6?DzuIm#!t=FH{UWuDgmAg(vS7)tH`|+nZ<+-wA#(g;XW4JhUo6t3
z>Jqz#V<Yo8Ff}!y1>6$>3okX>R~OsOZi|nutA?ErJHNPmt`r1Kcd^YH-2fcalGZOY
zLl&Ir)={uiq_&zgQ}VpF-|4VQ;KxFm3Sq%ei4d&(#Yl1{(gw<b4c$z-khBpC(d*pG
zX}{g@15lixbXo+T)6TFu%&6|yJATMz%dD~1WotlvzvVaP%XB_TA4sh>MRl3lQhSY$
zPXVM2Ao8e!B+eH4rJ%M3brBOz*BV}tikN}G3Jbdah0PWQy03+O?y9@a+I@1#b<cKL
z0P6HQOl?=wBAc=4n7~cPRt8Lq_<~k!`f#lgzH>M>4vFQ~;}Q16q1nEP0f(G_FvEDp
zdo-Z)5Dz0CW%36LY=e0o*ZEh<E1sYzubA+N9LGWfIj#mx_Mc=qE}DU)JZ~<ET`pdn
zG2s#W!fuNb)@HvbpZiq^>oA|si)SgkcMk*J>)*g9D)BtEB_2fZULMh3;u?1Z9iIpA
zhh6!I!6n3%-$4kj^8`KZq!CCXkVYVlz%vyA+{gJ3cYN5{#XkQqCrji<au=AB@a{hd
z@^u_(ll+aj&(3DExjz8<glDLc{gP~rZe_k8iqgv{*&ZP$cOV3roRaH}DdW&z*ny~I
zZxUr2=Qu)g+Zd@W;NKwM+(w;G4ge-!-5XaF(t`FV`T3V@rlf$du8?t^knisOkc;k?
zdS+Cy32E){j~Jide%hhftNgz;Fp)|4Rxk+v$NXcSDNF+M{u<yG=33+!^9oNtx}w~=
z{xRM9sYkYM9VEAI-vk{HB_kPGBcyi|+I($54Zsr<*35{quQ#ep^l~Jtk3k(}#&UKR
z>LfEcVMCSZnf+KJ%R4w0R2tOC@(z5D$wWS6@Ug44gVlun9{K%iLY{sdMXsioGy-V^
z(g>sxNF$I&AdNs8fiwbX1fJ;#$g}1LFK{;}g}3<5gAkpy16~%d`|_=Sf5a7M&d^pR
zL(u!^<orB!12{^g8~iAknt3&^QhKDqI=)An4T>Tf5VcuNGp1clS53#5wk_3hJ#SjG
zOx18bdX(%o%#rmlY;}q=Gi$YQy<5q<&HBta=K0M=bZFVEx4NM}6KUo_glcZZa~$2Q
zT2|Gr*sf{m%=dJM`F_PTD_*6l8LDBaN9p5@90VKnXkbG(w1VbY1&>t<o?0***LImv
zowa5Sk6G@lqj|3B8GAGhpIMqyFb&r>wb`of+XcJoR8-Y44cjz*%hoJI_x&fHhV8hj
zrMX^3ovr!>rn_pjU>H`V;5eq~!FhY7QgwYtJxa>8##O&vq~-DgPia5pyEX7T16EV{
z<s_GXDHtW&L-f~COo8g8NqMZh-Oejsk@TVS3X*R$L!acm%|>UlPV(!lv(9^Ao8&|P
zS{U_Mz2=fU|3gFG?RYF?1Z%18H$VV1LskL2-U-{h#W6W{OuM2*SDg%=d|X-h;7n=h
z)O^zUztO|EVT_`$N2o6w0?Yq%APWKiw=oE)i)&h{$#pLANg<4SXbZn>A;|TA0NS_*
zNPW4^C8oebkbG%h;#Yx(=S7k)_YsNZK7wN)@H*jbydD%ZxlhS+DvA65^N}h;0~`W6
z>Y@)3&?H_JUFol!zx<v#1^l5QBaKxm6o=5T?7P56nmQyW^?w^u80wM*(37zn9Af!x
x-uqS%F~>Lu>A%ch8EE~v3t2?>h5knaG;9vrUjf=~`tM<20mLx81J_~wKLJkil%fCt

literal 0
HcmV?d00001