From 09f8c228cfa32ea0ac6eda55498a5000e5a990d0 Mon Sep 17 00:00:00 2001
From: Maneesh Gupta <maneesh.gupta@amd.com>
Date: Thu, 9 May 2019 08:42:00 +0530
Subject: [PATCH] [dtests] Add hipModule test to unit tests

Change-Id: I1dac38f8580265e2e9c82d88e4f070a2ff87f60b


[ROCm/hip commit: 79843f3b1287fb31dabb9c0e7c0c422c98303410]
---
 .../tests/src/runtimeApi/module/hipModule.cpp |  61 ++++++++++--------
 .../src/runtimeApi/module/vcpy_kernel.code    | Bin 18811 -> 0 bytes
 .../src/runtimeApi/module/vcpy_kernel.cpp     |   6 +-
 3 files changed, 38 insertions(+), 29 deletions(-)
 delete mode 100755 projects/hip/tests/src/runtimeApi/module/vcpy_kernel.code

diff --git a/projects/hip/tests/src/runtimeApi/module/hipModule.cpp b/projects/hip/tests/src/runtimeApi/module/hipModule.cpp
index 22b4e7ff49..019645ea32 100644
--- a/projects/hip/tests/src/runtimeApi/module/hipModule.cpp
+++ b/projects/hip/tests/src/runtimeApi/module/hipModule.cpp
@@ -17,6 +17,13 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
 
+/* HIT_START
+ * BUILD_CMD: vcpy_kernel.code %hc --genco %S/vcpy_kernel.cpp -o vcpy_kernel.code
+ * BUILD: %t %s ../../test_common.cpp
+ * TEST: %t
+ * HIT_END
+ */
+
 #include "hip/hip_runtime.h"
 #include "hip/hip_runtime_api.h"
 #include <iostream>
@@ -33,13 +40,15 @@ THE SOFTWARE.
 #define fileName "vcpy_kernel.code"
 #define kernel_name "hello_world"
 
-__global__ void Cpy(hipLaunchParm lp, float* Ad, float* Bd) {
-    int tx = threadIdx.x;
-    Bd[tx] = Ad[tx];
-}
+#define HIP_CHECK(status)                                                                          \
+    if (status != hipSuccess) {                                                                    \
+        std::cout << "Got Status: " << status << " at Line: " << __LINE__ << std::endl;            \
+        exit(0);                                                                                   \
+    }
 
 int main() {
-    float *A, *B, *Ad, *Bd;
+    float *A, *B;
+    hipDeviceptr_t Ad, Bd;
     A = new float[LEN];
     B = new float[LEN];
 
@@ -48,32 +57,38 @@ int main() {
         B[i] = 0.0f;
     }
 
+    HIPCHECK(hipInit(0));
+
+    hipDevice_t device;
+    hipCtx_t context;
+    HIPCHECK(hipDeviceGet(&device, 0));
+    HIPCHECK(hipCtxCreate(&context, 0, device));
+
     HIPCHECK(hipMalloc((void**)&Ad, SIZE));
     HIPCHECK(hipMalloc((void**)&Bd, SIZE));
-
     HIPCHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
     HIPCHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice));
+
     hipModule_t Module;
     hipFunction_t Function;
+    hipFunction_t f;
     HIPCHECK(hipModuleLoad(&Module, fileName));
     HIPCHECK(hipModuleGetFunction(&Function, Module, kernel_name));
-    hipFunction_t f;
-    HIPCHECK(hipModuleGetFunction(&f, Module, kernel_name));
-    assert(f == Function);
+
     hipStream_t stream;
     HIPCHECK(hipStreamCreate(&stream));
-    void* args[2] = {&Ad, &Bd};
 
-    std::vector<void*> argBuffer(5);
-    memcpy(&argBuffer[3], &Ad, sizeof(void*));
-    memcpy(&argBuffer[4], &Bd, sizeof(void*));
+    struct {
+        void* _Ad;
+        void* _Bd;
+    } args;
+    args._Ad = (void*) Ad;
+    args._Bd = (void*) Bd;
+    size_t size = sizeof(args);
 
-    size_t size = argBuffer.size() * sizeof(void*);
-
-    void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &argBuffer[0], HIP_LAUNCH_PARAM_BUFFER_SIZE,
-                      &size, HIP_LAUNCH_PARAM_END};
-
-    hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, stream, NULL, (void**)&config);
+    void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
+                      HIP_LAUNCH_PARAM_END};
+    HIP_CHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, stream, NULL, (void**)&config));
 
     HIPCHECK(hipStreamDestroy(stream));
 
@@ -83,12 +98,6 @@ int main() {
         assert(A[i] == B[i]);
     }
 
-    std::vector<hipFunction_t> vec(1024 * 1024 * 64);
-    for (unsigned i = 0; i < 1024 * 1024 * 64; i++) {
-        hipFunction_t func;
-        hipModuleGetFunction(&func, Module, kernel_name);
-        vec[i] = func;
-    }
+	HIPCHECK(hipCtxDestroy(context));
     passed();
-    return 0;
 }
diff --git a/projects/hip/tests/src/runtimeApi/module/vcpy_kernel.code b/projects/hip/tests/src/runtimeApi/module/vcpy_kernel.code
deleted file mode 100755
index 4fd6472496c32a392baa767681f333d69a37ce29..0000000000000000000000000000000000000000
GIT binary patch
literal 0
HcmV?d00001

literal 18811
zcmeHPOKclO7@n+Mw++u0icpKtLQ9*X#92RLJCUfICZ#PTX$S?W<*`|N*RjDb)|*l&
zO-m8)BE(B@Kmu{%fDku&BNAs2aN*J;NEH%yD)qqs&(81cY&KO<f=ZU1DErU<&2#_n
zpV{%+GiMG=JgD^aFbXk<aP3wXDT;(kqap&L$O>!|TyJEXSpxJp$K3RyD<IVki=t|h
zn*@$9mRl(t+w7*0WTHP<ZUZI>-7nQWZaI>aj^(hg-HsJefl%aEHtUM6q;d-2|D+t<
z-^}lh$?vZyD)z})0h1JFVvj&M-kwOmR)KBt_OQbkeDYJFlO)@t#6DyDqm$#j7{(Fb
z0srU^h{xGRo)6S62Kd40$K1S!p#F^_@4&!7?<v)2n90fU$1S^4Yc)sIC)9Mms;N))
z^Mq-;)^sX|o0S%(>`9GeMzZ{G=gyt#p|LS_=CP@xhbIrH2WDoD%siP-tM_$O!%-`>
zx&<02^;(n1%qN?B`qe$@Ro9SNbq(27*D$o|8gi?yA;0Px3ahI@eM&vun@lEq(brZA
z=%d?lCJ^($v=~4p;9C*j<(5<hev0tjGX5UneKP(f;rGk<BH{fq{wv{08UK@TO~yAs
zzj!~>GQNZGtc<4#&&l{hgx}@j{h(u*CH!tz9~bxu!c`Zi{O1Y3SH`~}e5;G^67^gl
z{0<lI6ZnsW-!A9>o$x#5{8tFy?&4Wd&o)r;9|O)SxA7QbHwcN{!FfwG!_gbzDjpkL
zzX`5ec#JW0rd!~jrbPiuLNv|iH%thchxxMy($MX2j=UNEdC-qfrPBJri6i?*Cv<&E
zPi1rZ@l>j6)$1+&`IcQbQ@J_2X6kigzFDs7ZNqL%S02@Ly#l(^vQxvlUOsu!XtZ_5
zhHDFFxC1pFrXQ|MSD3C>Yi+$+)@!z2O_!}&U58oJ54So2sCBZP4uC85R?7~+>#ez<
z>~^^wklF4$XFCCSr`8NYgEKoeSYdf~!!!drQfU|gXk)$}fOh6f%TdglUDg_))!H(2
zJ~%VNLAg7I$*Um0h-74e$O4fCA`3(o_zzlu);v9YH3gsc2w0)7J?txX9?~}R71dHE
zv3%yhCifG1H!&WM2Q8(#>DM!&u*4!RNk3x>tbE`7<s7UK&jF?ac7~n*4pw>T9wo8R
zvoTR%Z2BUsAL9yOF@-Jk^dt&xo~M0L!NnkZ8S<~MKy+U#|7->RaV<_E;ag?@xA|8V
zvETcdvWOIN&XHS1S%41$@R4Huaf<hIts&z7s{F&pyoi6<mxAa2P>z50`Y#M7`Y`)I
zH2$wL{@H60|MK7mEqMInyEh7n?i(~0Sx?V@SpO~JT;^Q|C`x>Rv5P;mST8IQ!u34(
zcg8<{YqS*B6aQY`SuDQqvJ&6NwrzmdzTx7X*jX$7<?)Xz>7}rq_+M%K$Jify{9jfQ
z<LD2U<B5H-5ctJf@h`9cLNoryVf=Hxqw!DUIr#Y>QvQp;e=(Z>iF=g6{Cj8o_4xjW
zp8M==W#RH$@a|{Z26$%~&S$=t&9$2UL&N_^!2gF4|6U)F;~^jauK@q&BK~F0y;{im
ze-rqBBjVr7xEv2T|1SgoXCwY)&b?a5`9B5xzZmiFWn7Mjoc}Yx|3bvS%(+(!IsYF3
z|L;fqdl{GGA?N=z@c&Z8zs$K;3pxMq0RQLzKm2={jpC67A`3(oh%69UAhN*!!vgff
z{LjS~IJCb~Y}4%ilK&8H!h}GZa<_<$l=QvNCsdJ}ZRv`exyyB<Ij5d*w}PwLq?XiF
zwXbA!EK_YY`F`&dfT_He%M4`l`J85H1BG-YtqoPUUQQ3GdzoC#5U-}M>NxGr$l%~y
z&8g0plI2!oaK<pLR`W>PYL2xU?Rm!<<Rz1kMYW7ldbm<bo0*|}rJ$wDh0JiOY+5-j
zRT(xdOEYuXTu#&Ws@GO?y?&y>D;`Q&W+`P^rcuhK$^~;clh0<d1;Z-mhlVpHt&}NM
zhP6U!FPo?}pSA1}bz)-NwFP6z&I=94D1mfrS5_(3Y&jN7nlCgV3rkjwPSw>Mn<X9V
zq{B6%Q7f~gyB{A+(5%$v7@DGC!T!)S33wI&x`3d4DIx$mSAZHLzaae{h@fv#dGa@s
ze~B+5u^dL}_XOmJdjZ3}w^W||jpR3whJn(leUgKahQCUobn<7C<j+_K2Krvq)fmeQ
ztn22u0T_dj-SZ|~bXf%t0FGt5zl<r0j6%)|mHJEfPt80C>Ha=pk=Y_%mj{1dIHrM@
z*L+#2{C&Oxq>{va)ZM;&>Uee3Ri5@>7^QLqqfWLb4CSS9SNlmMqL1JlL-$Yo<G${o
z_AeZC%YW}{(5JeY0<Jut0?ZVbC&utD0y-T7ko?wLDh53yf0Xe!BR`XI@Mt2~lMq8s
zDS7}Ajqwo}gp;+F($I!M=Of^D{h;#^1_<Bn)}aamjuGA`<HrfVU&d|1`(^x1!hPo>
zd`7r%OltRA!hPo>{6M(xe1zW!_nnV$g>c{b2wUO)xE*}wBj7m(99N+r1UesKAK~}P
z_))@r=Oa8rxbJ*~vxNK3NBEd<;qY{x!q<fR&PQPS__3+c$-`sJbHafCD1hjM0NEnA
OlYiO%{_6EQA>ePX$&UR1

diff --git a/projects/hip/tests/src/runtimeApi/module/vcpy_kernel.cpp b/projects/hip/tests/src/runtimeApi/module/vcpy_kernel.cpp
index c99e3df985..d7c23a8810 100644
--- a/projects/hip/tests/src/runtimeApi/module/vcpy_kernel.cpp
+++ b/projects/hip/tests/src/runtimeApi/module/vcpy_kernel.cpp
@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2015 - present 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
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 #include "hip/hip_runtime.h"
 
-extern "C" __global__ void hello_world(hipLaunchParm lp, float* a, float* b) {
-    int tx = threadIdx.x;
+extern "C" __global__ void hello_world(float* a, float* b) {
+    int tx = hipThreadIdx_x;
     b[tx] = a[tx];
 }