From 53aeee42f9df67267269669d3002fa7c293a2ce3 Mon Sep 17 00:00:00 2001 From: ansurya <50609411+ansurya@users.noreply.github.com> Date: Thu, 18 Jul 2019 08:48:15 +0530 Subject: [PATCH] HIPCommander code cleanup (#1207) * HIPCommander code cleanup * Removed non-used headerfiles * Removed empty ifdef --- samples/1_Utils/hipCommander/c.cmd | 4 +- samples/1_Utils/hipCommander/classic.cmd | 2 +- samples/1_Utils/hipCommander/hipCommander.cpp | 164 ++---------------- samples/1_Utils/hipCommander/l2.hcm | 4 +- samples/1_Utils/hipCommander/loop.hcm | 4 +- samples/1_Utils/hipCommander/loop2.hcm | 4 +- samples/1_Utils/hipCommander/nullkernel.hsaco | Bin 10265 -> 39333 bytes .../hipCommander/perf/latency_hostsync.hcm | 8 +- .../hipCommander/perf/latency_nosync.hcm | 8 +- .../hipCommander/perf/latency_nullstream.hcm | 8 +- .../perf/modulelaunch_latency.hcm | 2 +- samples/1_Utils/hipCommander/setstream.hcm | 4 +- 12 files changed, 34 insertions(+), 178 deletions(-) diff --git a/samples/1_Utils/hipCommander/c.cmd b/samples/1_Utils/hipCommander/c.cmd index db11071203..4cb980eccb 100644 --- a/samples/1_Utils/hipCommander/c.cmd +++ b/samples/1_Utils/hipCommander/c.cmd @@ -1,3 +1,3 @@ -loop,1000; H2D; NullKernel; D2H; endloop; +loop(1000); H2D; NullKernel; D2H; endloop; streamsync; -printTiming, 1000 +printTiming(1000) diff --git a/samples/1_Utils/hipCommander/classic.cmd b/samples/1_Utils/hipCommander/classic.cmd index c149eec5f7..7b1e4273c9 100644 --- a/samples/1_Utils/hipCommander/classic.cmd +++ b/samples/1_Utils/hipCommander/classic.cmd @@ -1 +1 @@ -H2D; NullKernel, D2H, streamsync +H2D; NullKernel; D2H; streamsync diff --git a/samples/1_Utils/hipCommander/hipCommander.cpp b/samples/1_Utils/hipCommander/hipCommander.cpp index 457cdfa7d3..21b5505623 100644 --- a/samples/1_Utils/hipCommander/hipCommander.cpp +++ b/samples/1_Utils/hipCommander/hipCommander.cpp @@ -7,11 +7,6 @@ #include #include -#ifdef __HIP_PLATFORM_HCC__ -#include -#include -#include -#endif #include @@ -283,100 +278,6 @@ class Command { #define FILENAME "nullkernel.hsaco" #define KERNEL_NAME "NullKernel" - -#ifdef __HIP_PLATFORM_HCC__ -//================================================================================================= -// Use Aql to launch the NULL kernel. -class AqlKernelCommand : public Command { - public: - AqlKernelCommand(CommandStream* cmdStream, const std::vector args) - : Command(cmdStream, args) { - hc::accelerator_view* av; - HIPCHECK(hipHccGetAcceleratorView(cmdStream->currentStream(), &av)); - - hc::accelerator acc = av->get_accelerator(); - - hsa_region_t systemRegion = *(hsa_region_t*)acc.get_hsa_am_system_region(); - - _hsaAgent = *(hsa_agent_t*)acc.get_hsa_agent(); - - std::ifstream file(FILENAME, std::ios::binary | std::ios::ate); - std::streamsize fsize = file.tellg(); - file.seekg(0, std::ios::beg); - - std::vector buffer(fsize); - if (file.read(buffer.data(), fsize)) { - uint64_t elfSize = ElfSize(&buffer[0]); - - assert(fsize == elfSize); - - // TODO - replace module load code with explicit module load and unload. - - hipModule_t module; - HIPCHECK(hipModuleLoadData(&module, &buffer[0])); - HIPCHECK(hipModuleGetFunction(&_function, module, KERNEL_NAME)); - - } else { - failed("could not open code object '%s'\n", FILENAME); - } - }; - - ~AqlKernelCommand(){}; - - void run() override { -#define LEN 64 - uint32_t len = LEN; - uint32_t one = 1; - - float* Ad = NULL; - - size_t argSize = 36; - char argBuffer[argSize]; - *(uint32_t*)(&argBuffer[0]) = len; - *(uint32_t*)(&argBuffer[4]) = one; - *(uint32_t*)(&argBuffer[8]) = one; - *(uint32_t*)(&argBuffer[12]) = len; - *(uint32_t*)(&argBuffer[16]) = one; - *(uint32_t*)(&argBuffer[20]) = one; - *(float**)(&argBuffer[24]) = Ad; // Ad pointer argument - - - void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &argBuffer[0], - HIP_LAUNCH_PARAM_BUFFER_SIZE, &argSize, HIP_LAUNCH_PARAM_END}; - - hipModuleLaunchKernel(_function, len, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config); - }; - - - public: - hsa_queue_t _hsaQueue; - hsa_agent_t _hsaAgent; - - hipFunction_t _function; - - private: - static uint64_t ElfSize(const void* emi) { - const Elf64_Ehdr* ehdr = (const Elf64_Ehdr*)emi; - const Elf64_Shdr* shdr = (const Elf64_Shdr*)((char*)emi + ehdr->e_shoff); - - uint64_t max_offset = ehdr->e_shoff; - uint64_t total_size = max_offset + ehdr->e_shentsize * ehdr->e_shnum; - - for (uint16_t i = 0; i < ehdr->e_shnum; ++i) { - uint64_t cur_offset = static_cast(shdr[i].sh_offset); - if (max_offset < cur_offset) { - max_offset = cur_offset; - total_size = max_offset; - if (SHT_NOBITS != shdr[i].sh_type) { - total_size += static_cast(shdr[i].sh_size); - } - } - } - return total_size; - } -}; -#endif - //================================================================================================= // HCC optimizes away fully NULL kernel calls, so run one that is nearly null: class ModuleKernelCommand : public Command { @@ -391,26 +292,18 @@ class ModuleKernelCommand : public Command { void run() override { #define LEN 64 - uint32_t len = LEN; - uint32_t one = 1; + float *X = NULL; + HIPCHECK(hipMalloc((void**)&X, sizeof(float))); + struct { + float *Ad; + }args; + args.Ad = X; + size_t argSize = sizeof(args); - float* Ad = NULL; - - size_t argSize = 36; - char argBuffer[argSize]; - *(uint32_t*)(&argBuffer[0]) = len; - *(uint32_t*)(&argBuffer[4]) = one; - *(uint32_t*)(&argBuffer[8]) = one; - *(uint32_t*)(&argBuffer[12]) = len; - *(uint32_t*)(&argBuffer[16]) = one; - *(uint32_t*)(&argBuffer[20]) = one; - *(float**)(&argBuffer[24]) = Ad; // Ad pointer argument - - - void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &argBuffer[0], + void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE, &argSize, HIP_LAUNCH_PARAM_END}; - hipModuleLaunchKernel(_function, len, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config); + hipModuleLaunchKernel(_function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config); }; @@ -447,33 +340,6 @@ class KernelCommand : public Command { hipStream_t _stream; }; - -#ifdef __HIP_PLATFORM_HCC__ -//================================================================================================= -class PfeCommand : public Command { - public: - PfeCommand(CommandStream* cmdStream, const std::vector args, - hipStream_t stream = 0) - : Command(cmdStream, args) { - HIPCHECK(hipHccGetAcceleratorView(stream, &_av)); - } - - ~PfeCommand() {} - - - void run() override { - static const int gridX = 64; - static const int groupX = 64; - auto cf = hc::parallel_for_each(*_av, hc::extent<1>(gridX).tile(groupX), - [=](hc::index<1>& idx) __HC__ {}); - } - - private: - hc::accelerator_view* _av; -}; -#endif - - //================================================================================================= class CopyCommand : public Command { enum MemType { PinnedHost, UnpinnedHost, Device }; @@ -708,10 +574,8 @@ CommandStream::CommandStream(std::string commandStreamString, int iterations) std::vector tokens; tokenize(commandStreamString, ';', tokens); - - std::for_each(tokens.begin(), tokens.end(), [&](const std::string s) { this->parse(s); }); - setStream(0); + std::for_each(tokens.begin(), tokens.end(), [&](const std::string s) { this->parse(s); }); } @@ -831,14 +695,6 @@ void CommandStream::parse(std::string fullCmd) { } else if (c == "vectoraddkernel") { cmd = new KernelCommand(cmdStream, args, KernelCommand::VectorAdd); -#ifdef __HIP_PLATFORM_HCC__ - } else if (c == "nullpfe") { - cmd = new PfeCommand(cmdStream, args); - - } else if (c == "aqlkernel") { - cmd = new AqlKernelCommand(cmdStream, args); -#endif - } else if (c == "devicesync") { cmd = new DeviceSyncCommand(cmdStream, args); diff --git a/samples/1_Utils/hipCommander/l2.hcm b/samples/1_Utils/hipCommander/l2.hcm index b541bd6a66..6b14f7b829 100644 --- a/samples/1_Utils/hipCommander/l2.hcm +++ b/samples/1_Utils/hipCommander/l2.hcm @@ -1,3 +1,3 @@ -setstream,1; +setstream(1); NullKernel; streamsync; -loop,10000; H2D; NullKernel; streamsync; endloop,1; +loop(10000); H2D; NullKernel; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/loop.hcm b/samples/1_Utils/hipCommander/loop.hcm index db11071203..4cb980eccb 100644 --- a/samples/1_Utils/hipCommander/loop.hcm +++ b/samples/1_Utils/hipCommander/loop.hcm @@ -1,3 +1,3 @@ -loop,1000; H2D; NullKernel; D2H; endloop; +loop(1000); H2D; NullKernel; D2H; endloop; streamsync; -printTiming, 1000 +printTiming(1000) diff --git a/samples/1_Utils/hipCommander/loop2.hcm b/samples/1_Utils/hipCommander/loop2.hcm index b8a14aa156..ae753d0722 100644 --- a/samples/1_Utils/hipCommander/loop2.hcm +++ b/samples/1_Utils/hipCommander/loop2.hcm @@ -1,2 +1,2 @@ -setstream,1; -loop,1000; NullKernel; syncstream; endloop,1, +setstream(1); +loop(1000); NullKernel; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/nullkernel.hsaco b/samples/1_Utils/hipCommander/nullkernel.hsaco index 585b55cce587ed928acc1308d8a76361cdd87e57..da6a3e6823371f3c3cee3eb9f4e014d6a509bf81 100755 GIT binary patch literal 39333 zcmeHQTW=Fb6rPxnh7_rkr(RTb1Zo46b=H?SHZRSMgbF7FQlJP`?s|9ZxLNO7Ya5aZ z0ewKI571Q8YkSEBbBRJY2YJuTIdgVqzMWZRcRprI zrJ3T?;(Y1K++6X>)NE<``r>TyVySeLDfwhTC$xHB>lMG*GS>3BQZ8$>y;aY@>ltp% zYp-Dd|5TOy2m@@}uO5-19yGKW>YrMe4Z|P}_7COnh_(;1x8y%1 zSXP#wU%IZ!!sK~dm)&tLNI#tSo@~pAE^8RZ$PE!RYrc1Zqul1H^W5ZLk6adkC){So zGs3yS7p=PJm}Jmm+jUh@>F0cTy*yDNE zFI(<(yHXKBbb_$dm+OtFg1JmOa!xr;u-ve0akXu^aGH7HiJ)eO1!jClECqfe>W~ft zJguO*EUIhmtP>&x>?*P@n207}P=9!Ew;ELdxEuGPBf2dk>h;U!3q=SGsrBn-hF1;+R}nECoX zgZC4#Oc_Fe5Fi8y0YZQfAOr{jLSR1;fWK4K z->dL+Cv#8z;TmtvPwa^_9Sfh}IzD{v1

w)Y$P24h{}&N$FS)Q0nO9<^%R*Bf)<8 z6*_-E@z>)a_IN{ra(^o^c%QK|(*ucw)Q#UmpxE~t2dMSb6heRyAOr{jLVyq;1PB2_ zfDj-A2mwOinGlH9n)P}+tl@&Zr<7W2m+2Ml$3?T_+W+@|TRtRB?T586kR_d0xz4dz z^HxQ0fftLjypS}LCO(t;QOtyyK>>K2`iUlwqI|wTSA)cx)BK1 zNGHd~#99khDc!LZ$-rpI@~-RfiHSnWDpv}*lvNSw@j}YUO_XyLONe~h%sPc^!OUe+ zj(wJ0nwjCYE7yVf9kl}e3R@;PCq^7)LN%@oX(?G!9qR$sQ7!r{Kh z0i0AalS~!Tsk9*~`K++V4c!q#pOK&KT-ohY>UGi(GvDp?pM-2ezH{Qec6bZ5-MTAF zdVWi=q_gfd*Xt}<_1ei5tGUAT5$vsiC0lY#gr%%{&1Ojn0W4p$qzr0VWvN$WF_lu9 zO;4MpH0xTPEzRCqoLcx|rnh0wnS=gApWq!^D45V1`55TN`Ga|gg*k~~(S+3V%hUxS z3w7apIGn%lNZUVKTRz6)JVni=9|ZjvudIv#uLI7Few^2+u^#ke`%!-^V}V10e!P#M z#`}mY5{eoxm}6{ICczfnr}TP!d|f_@h@HUDvCxK4c&|nt4&JV?|5!~<7Kc+J&N`Yk zxJvyNs}HlNKbOU@?-&dE(kYeuMyc=`LmRPv#C2ZUqPFX@pnlZxx4LYQ#Cp+uRVKEJ zA2_W_0LQ0W;_?0R!_LW*9b|r(sXK@o#}SyrF~&6-F^%H{HO9fU?I#Ye?_uPJVc%f9 z0hg>lewd^^$~d@h5I+pqVcjG!;)h*`$Xha3&$7aewfSTsH~Jn{FJ#1GqFe%N_=9f%($LEvgYBA*O3 z&d@ANCOT4tpTMq`kxT9#K?+5Fi8y0YZQfAOr{jLVyq;1PB2_ z;2c9oNr3U>FC}xSzO#$`2dVo6JR8!v}mVXfxmnoG<7v zq`@Blz9jO)0FO97?3A?a{(VWgt>hjn9+uXJAJI!F4xjh4ISz^!FukUZci( iFdpYW>iG91@jin3MZYf@@qDnsvG

f7pxdAND^s?cSgO literal 10265 zcmeHNO>7%Q6n-9W){e7rDA1N4A(nuuP*E$pN!(QGAwP9XMB;{~MIu7lxa&>h;Gf9e zRH;N?J4e){oH(I02gIQ#BrYIvLE_l(-t3J3NKt=6)k-_E=gs?@ zo&DZ>E6;i^Cq_@QKmcS5;-JO>3<$-{8`6Sg{Sd;w9U`EjSi$drP&d;xF`1-cVMk)} zyI<=RN=e~4&yxIFXejt*dOYOWNhdcY$op{GQ5pX!{nq9&&*`z#++QqalPqA>_pO> zv5lj4Avc$G+)N>lv$^v{CoyiAWIf>)om9?p<U<1^Z(sJ!sobsZ=PAn7uXqa53X% z>{NCrRdOtMBst;sSeDV3n|GasL?SVM)|wcya*5=*Y;vO8G7LA9b<$QQpLP~TMl%MQ zB&^++pBuOB-IqT5^e=sOwi;c3;3ME8;3ME8;3ME8;3ME8;3ME8@T?;swu88Lqn#`6 z>ZIU4H?{=tzzx*3;oiT#y;#TW=X2ErilXf0mbNzW9*h{UP1@I9Vd^GLcNwm4z)o~! zX9pl04S#{CnCrkgBk;@OtiE=u99~;H2WuNqO(`#&gwLu%#VEL9i;PS%*?aek}+q+w3*kPhaH};D@b>%AgWfdococln}#6tZWBckhX-) zV2cvRz4jf}imN4Gd;t=y`u?=V(3h!dZ&dwi5%1+rWjR3aNRZydXyV*q21;A_5EJ7i zC1T!I!Y0kx(h{M4JfAzNhU4U~gdz}CRa347<9Aq5)x${RNaGdt(?!x%*v$a04d=PK z9aV7sfy%9QCETe&#B673IS|tNSpdvHE9+NTv>bt-7Xha3w%|LOS*WbeAx34mckQp8 zYeZMe;d`q9OB)E`6KgWvA2k1$A0Yn^Kx@PLpS}Foy!_X8&wmT^U#rXiE02=@FN^#~ z`pEf@zDtYKP|40gfr}70iFFF7GwWdtMK1RD)hE z{&)UkePaG=SYs-DwvEpOtm_rdN41FA!L)LPDdkYR*1xWXJ45jK_Sd{tmTv{CwV{*& z=I&y-1Jy*REn=$sw0?j;ws>{2<<&A1b;9>wbO{vLU|F*D~{|R>qy=PVHNl$xYG9Mf} zLBF|qBl)ZNrC*99SMWpu*-?`+l%Ei$H#jE$k`eNO>-gO@T zu1)JGAd~KxfuDHKG4n`Y`dF{ilRnaKPxtkhaeZ|3gE7ykbTR)nK&*5r=cdjhb&DRI zk)!mbJSrhJlPb+Xj5xOlG5j0Y3^n9MnicmN%x^lVpDJW zx9)MddLfRbpA-`>V=rYyzkSQ&WFFG{Fv48oFY}Myi0pR#xkH{my{W2iO>DiYVeHu{ zaQ*oPv52(-XPe~h+)RP*_BYyl$6XZoq)E