From 1dcf618d2051673326b527eaf3ac7702f9a3f78c Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Wed, 17 Jul 2019 20:18:04 -0700 Subject: [PATCH 1/5] Fix HIP_VISIBLE_DEVICES order (#1184) * Fix HIP_VISIBLE_DEVICES order * Fix device IDs mismatch * Fix review comments- loop order and device range check * Handle incomplete VISIBLE device env variable * Revert "Handle incomplete VISIBLE device env variable" --- src/hip_hcc.cpp | 26 +++++++++++++++----------- 1 file changed, 15 insertions(+), 11 deletions(-) diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp index c0a3827376..b6c7e3d87a 100644 --- a/src/hip_hcc.cpp +++ b/src/hip_hcc.cpp @@ -1399,7 +1399,7 @@ void ihipInit() { // Make sure the hip visible devices are within the deviceCnt range for (int i = 0; i < g_hip_visible_devices.size(); i++) { - if (g_hip_visible_devices[i] >= deviceCnt) { + if ((g_hip_visible_devices[i] >= deviceCnt) ||(g_hip_visible_devices[i] < 0)){ // Make sure any DeviceID after invalid DeviceID will be erased. g_hip_visible_devices.resize(i); break; @@ -1415,17 +1415,21 @@ void ihipInit() { g_deviceArray = new ihipDevice_t*[deviceCnt]; g_deviceCnt = 0; - for (int i = 0; i < accs.size(); i++) { - // check if the device id is included in the HIP_VISIBLE_DEVICES env variable - if (!accs[i].get_is_emulated()) { - if (std::find(g_hip_visible_devices.begin(), g_hip_visible_devices.end(), (i - 1)) == - g_hip_visible_devices.end() && - g_visible_device) { - // If device is not in visible devices list, ignore - continue; + + if(g_visible_device) { + for (int i = 0; i < g_hip_visible_devices.size(); i++) { + int devIndex = g_hip_visible_devices[i]; + if (!accs[devIndex+1].get_is_emulated()) { + g_deviceArray[g_deviceCnt] = new ihipDevice_t(g_deviceCnt, deviceCnt, accs[devIndex+1]); + g_deviceCnt++; + } + } + }else { + for (int i = 0; i < accs.size(); i++) { + if (!accs[i].get_is_emulated()) { + g_deviceArray[g_deviceCnt] = new ihipDevice_t(g_deviceCnt, deviceCnt, accs[i]); + g_deviceCnt++; } - g_deviceArray[g_deviceCnt] = new ihipDevice_t(g_deviceCnt, deviceCnt, accs[i]); - g_deviceCnt++; } } From 84f84411461dfc89986e63c2f65ff16a467ed0a3 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 2/5] 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 Date: Thu, 18 Jul 2019 08:48:28 +0530 Subject: [PATCH 3/5] UChar and UShort textures as Normalized Float (#1213) * UChar and UShort textures as Normalized Float * UChar and UShort textures as Normalized Float for all float variants * Handled uninitilaized texture format value --- include/hip/hcc_detail/texture_functions.h | 30 ++++++++++++++++------ 1 file changed, 22 insertions(+), 8 deletions(-) diff --git a/include/hip/hcc_detail/texture_functions.h b/include/hip/hcc_detail/texture_functions.h index 1f22e01177..bb0a9e7223 100644 --- a/include/hip/hcc_detail/texture_functions.h +++ b/include/hip/hcc_detail/texture_functions.h @@ -22,7 +22,6 @@ THE SOFTWARE. #ifndef HIP_INCLUDE_HIP_HCC_DETAIL_TEXTURE_FUNCTIONS_H #define HIP_INCLUDE_HIP_HCC_DETAIL_TEXTURE_FUNCTIONS_H - #include #include @@ -118,8 +117,6 @@ union TData { #define TEXTURE_RETURN_UINT return texel.u.x; -#define TEXTURE_RETURN_FLOAT return texel.f.x; - #define TEXTURE_RETURN_SIGNED return texel.i.x; #define TEXTURE_RETURN_UNSIGNED return texel.u.x; @@ -136,8 +133,6 @@ union TData { #define TEXTURE_RETURN_UINT_X return make_uint1(texel.u.x); -#define TEXTURE_RETURN_FLOAT_X return make_float1(texel.f.x); - #define TEXTURE_RETURN_CHAR_XY return make_char2(texel.i.x, texel.i.y); #define TEXTURE_RETURN_UCHAR_XY return make_uchar2(texel.u.x, texel.u.y); @@ -150,8 +145,6 @@ union TData { #define TEXTURE_RETURN_UINT_XY return make_uint2(texel.u.x, texel.u.y); -#define TEXTURE_RETURN_FLOAT_XY return make_float2(texel.f.x, texel.f.y); - #define TEXTURE_RETURN_CHAR_XYZW return make_char4(texel.i.x, texel.i.y, texel.i.z, texel.i.w); #define TEXTURE_RETURN_UCHAR_XYZW return make_uchar4(texel.u.x, texel.u.y, texel.u.z, texel.u.w); @@ -164,10 +157,30 @@ union TData { #define TEXTURE_RETURN_UINT_XYZW return make_uint4(texel.u.x, texel.u.y, texel.u.z, texel.u.w); -#define TEXTURE_RETURN_FLOAT_XYZW return make_float4(texel.f.x, texel.f.y, texel.f.z, texel.f.w); +#define HIP_AD_FORMAT_NOT_INITIALIZED 0 + +#define TEXTURE_RETURN_FLOAT return (texFormatToSize[texRef.format] == 1)? texel.f.x : (float)texel.u.x/texFormatToSize[texRef.format]; + +#define TEXTURE_RETURN_FLOAT_X return (texFormatToSize[texRef.format] == 1)? make_float1(texel.f.x) : make_float1((float)texel.u.x/texFormatToSize[texRef.format]); + +#define TEXTURE_RETURN_FLOAT_XY return (texFormatToSize[texRef.format] == 1)? make_float2(texel.f.x, texel.f.y) : make_float2((float)texel.u.x/texFormatToSize[texRef.format], (float)texel.u.y/texFormatToSize[texRef.format]); + +#define TEXTURE_RETURN_FLOAT_XYZW return (texFormatToSize[texRef.format] == 1)? make_float4(texel.f.x, texel.f.y, texel.f.z, texel.f.w) : make_float4((float)texel.u.x/texFormatToSize[texRef.format], (float)texel.u.y/texFormatToSize[texRef.format], (float)texel.u.z/texFormatToSize[texRef.format], (float)texel.u.w/texFormatToSize[texRef.format]) ; extern "C" { + __device__ __constant__ static int texFormatToSize[] = { + [HIP_AD_FORMAT_NOT_INITIALIZED] = 1 , + [HIP_AD_FORMAT_UNSIGNED_INT8] = UCHAR_MAX , + [HIP_AD_FORMAT_UNSIGNED_INT16]= USHRT_MAX, + [HIP_AD_FORMAT_UNSIGNED_INT32]= 1 , + [HIP_AD_FORMAT_SIGNED_INT8] = SCHAR_MAX, + [HIP_AD_FORMAT_SIGNED_INT16] = SHRT_MAX, + [HIP_AD_FORMAT_SIGNED_INT32] = 1 , + [HIP_AD_FORMAT_HALF] = 1 , + [HIP_AD_FORMAT_FLOAT] = 1 +}; + __device__ __hip_float4_vector_value_type __ockl_image_sample_1D( unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s, @@ -3863,6 +3876,7 @@ __TEXTURE_FUNCTIONS_DECL__ float tex1D(texture texRef, TEXTURE_RETURN_FLOAT; } ////// + template __TEXTURE_FUNCTIONS_DECL__ float tex1D(texture texRef, float x) { TEXTURE_REF_PARAMETERS_INIT; From 8e496c09d9d4be2e28901f775c89e20d7f7d4af6 Mon Sep 17 00:00:00 2001 From: ansurya <50609411+ansurya@users.noreply.github.com> Date: Thu, 18 Jul 2019 08:48:50 +0530 Subject: [PATCH 4/5] Add Max Texture 1D,2D,3D device properties (#1226) * Add Max Texture 1D,2D,3D device properties * Corrected testcase to use enums defined in hipDeviceAttribute_t * Added texture 1D,2D and 3D support for NVIDIA path --- bin/hipify-perl | 6 +++++ ..._Runtime_API_functions_supported_by_HIP.md | 12 ++++----- .../src/CUDA2HIP_Driver_API_types.cpp | 12 ++++----- .../src/CUDA2HIP_Runtime_API_types.cpp | 12 ++++----- include/hip/hip_runtime_api.h | 13 +++++++++- include/hip/nvcc_detail/hip_runtime_api.h | 25 +++++++++++++++++++ src/hip_device.cpp | 18 +++++++++++++ src/hip_hcc.cpp | 14 ++++++++++- .../device/hipGetDeviceAttribute.cpp | 7 ++++++ 9 files changed, 99 insertions(+), 20 deletions(-) diff --git a/bin/hipify-perl b/bin/hipify-perl index 3e39a5a11d..e375d22928 100755 --- a/bin/hipify-perl +++ b/bin/hipify-perl @@ -321,6 +321,12 @@ while (@ARGV) { $ft{'err'} += s/\bcudaDevAttrMaxSharedMemoryPerMultiprocessor\b/hipDeviceAttributeMaxSharedMemoryPerMultiprocessor/g; $ft{'err'} += s/\bcudaDevAttrMemoryClockRate\b/hipDeviceAttributeMemoryClockRate/g; $ft{'err'} += s/\bcudaDevAttrGlobalMemoryBusWidth\b/hipDeviceAttributeMemoryBusWidth/g; + $ft{'err'} += s/\bcudaDevAttrMaxTexture1DWidth\b/hipDeviceAttributeMaxTexture1DWidth/g; + $ft{'err'} += s/\bcudaDevAttrMaxTexture2DWidth\b/hipDeviceAttributeMaxTexture2DWidth/g; + $ft{'err'} += s/\bcudaDevAttrMaxTexture2DHeight\b/hipDeviceAttributeMaxTexture2DHeight/g; + $ft{'err'} += s/\bcudaDevAttrMaxTexture3DWidth\b/hipDeviceAttributeMaxTexture3DWidth/g; + $ft{'err'} += s/\bcudaDevAttrMaxTexture3DHeight\b/hipDeviceAttributeMaxTexture3DHeight/g; + $ft{'err'} += s/\bcudaDevAttrMaxTexture3DDepth\b/hipDeviceAttributeMaxTexture3DDepth/g; #-------- $ft{'dev'} += s/\bcudaDeviceAttr\b/hipDeviceAttribute_t/g; $ft{'dev'} += s/\bcudaDeviceGetAttribute\b/hipDeviceGetAttribute/g; diff --git a/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md b/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md index 663e6ad082..afd05b7d87 100644 --- a/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md +++ b/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md @@ -519,12 +519,12 @@ | 18 |*`cudaDevAttrIntegrated`* |*`hipDeviceAttributeIntegrated`* | | 19 |*`cudaDevAttrCanMapHostMemory`* | | | 20 |*`cudaDevAttrComputeMode`* |*`hipDeviceAttributeComputeMode`* | -| 21 |*`cudaDevAttrMaxTexture1DWidth`* | | -| 22 |*`cudaDevAttrMaxTexture2DWidth`* | | -| 23 |*`cudaDevAttrMaxTexture2DHeight`* | | -| 24 |*`cudaDevAttrMaxTexture3DWidth`* | | -| 25 |*`cudaDevAttrMaxTexture3DHeight`* | | -| 26 |*`cudaDevAttrMaxTexture3DDepth`* | | +| 21 |*`cudaDevAttrMaxTexture1DWidth`* |*`hipDeviceAttributeMaxTexture1DWidth`* | +| 22 |*`cudaDevAttrMaxTexture2DWidth`* |*`hipDeviceAttributeMaxTexture2DWidth`* | +| 23 |*`cudaDevAttrMaxTexture2DHeight`* |*`hipDeviceAttributeMaxTexture2DHeight`* | +| 24 |*`cudaDevAttrMaxTexture3DWidth`* |*`hipDeviceAttributeMaxTexture3DWidth`* | +| 25 |*`cudaDevAttrMaxTexture3DHeight`* |*`hipDeviceAttributeMaxTexture3DHeight`* | +| 26 |*`cudaDevAttrMaxTexture3DDepth`* |*`hipDeviceAttributeMaxTexture3DDepth`* | | 27 |*`cudaDevAttrMaxTexture2DLayeredWidth`* | | | 28 |*`cudaDevAttrMaxTexture2DLayeredHeight`* | | | 29 |*`cudaDevAttrMaxTexture2DLayeredLayers`* | | diff --git a/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp b/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp index 51dc9665c3..cc4502a866 100644 --- a/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp +++ b/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp @@ -328,17 +328,17 @@ const std::map CUDA_DRIVER_TYPE_NAME_MAP{ // cudaDevAttrComputeMode {"CU_DEVICE_ATTRIBUTE_COMPUTE_MODE", {"hipDeviceAttributeComputeMode", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 20 // cudaDevAttrMaxTexture1DWidth - {"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_WIDTH", {"hipDeviceAttributeMaxTexture1DWidth", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 21 + {"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_WIDTH", {"hipDeviceAttributeMaxTexture1DWidth", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 21 // cudaDevAttrMaxTexture2DWidth - {"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_WIDTH", {"hipDeviceAttributeMaxTexture2DWidth", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 22 + {"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_WIDTH", {"hipDeviceAttributeMaxTexture2DWidth", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 22 // cudaDevAttrMaxTexture2DHeight - {"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_HEIGHT", {"hipDeviceAttributeMaxTexture2DHeight", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 23 + {"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_HEIGHT", {"hipDeviceAttributeMaxTexture2DHeight", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 23 // cudaDevAttrMaxTexture3DWidth - {"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_WIDTH", {"hipDeviceAttributeMaxTexture3DWidth", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 24 + {"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_WIDTH", {"hipDeviceAttributeMaxTexture3DWidth", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 24 // cudaDevAttrMaxTexture3DHeight - {"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_HEIGHT", {"hipDeviceAttributeMaxTexture3DHeight", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 25 + {"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_HEIGHT", {"hipDeviceAttributeMaxTexture3DHeight", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 25 // cudaDevAttrMaxTexture3DDepth - {"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_DEPTH", {"hipDeviceAttributeMaxTexture3DDepth", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 26 + {"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_DEPTH", {"hipDeviceAttributeMaxTexture3DDepth", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 26 // cudaDevAttrMaxTexture2DLayeredWidth {"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LAYERED_WIDTH", {"hipDeviceAttributeMaxTexture2DLayeredWidth", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 27 // cudaDevAttrMaxTexture2DLayeredHeight diff --git a/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp b/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp index 6e86cf8fda..662655edb9 100644 --- a/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp +++ b/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp @@ -255,17 +255,17 @@ const std::map CUDA_RUNTIME_TYPE_NAME_MAP { // CU_DEVICE_ATTRIBUTE_COMPUTE_MODE {"cudaDevAttrComputeMode", {"hipDeviceAttributeComputeMode", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 20 // CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_WIDTH - {"cudaDevAttrMaxTexture1DWidth", {"hipDeviceAttributeMaxTexture1DWidth", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 21 + {"cudaDevAttrMaxTexture1DWidth", {"hipDeviceAttributeMaxTexture1DWidth", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 21 // CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_WIDTH - {"cudaDevAttrMaxTexture2DWidth", {"hipDeviceAttributeMaxTexture2DWidth", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 22 + {"cudaDevAttrMaxTexture2DWidth", {"hipDeviceAttributeMaxTexture2DWidth", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 22 // CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_HEIGHT - {"cudaDevAttrMaxTexture2DHeight", {"hipDeviceAttributeMaxTexture2DHeight", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 23 + {"cudaDevAttrMaxTexture2DHeight", {"hipDeviceAttributeMaxTexture2DHeight", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 23 // CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_WIDTH - {"cudaDevAttrMaxTexture3DWidth", {"hipDeviceAttributeMaxTexture3DWidth", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 24 + {"cudaDevAttrMaxTexture3DWidth", {"hipDeviceAttributeMaxTexture3DWidth", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 24 // CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_HEIGHT - {"cudaDevAttrMaxTexture3DHeight", {"hipDeviceAttributeMaxTexture3DHeight", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 25 + {"cudaDevAttrMaxTexture3DHeight", {"hipDeviceAttributeMaxTexture3DHeight", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 25 // CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_DEPTH - {"cudaDevAttrMaxTexture3DDepth", {"hipDeviceAttributeMaxTexture3DDepth", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 26 + {"cudaDevAttrMaxTexture3DDepth", {"hipDeviceAttributeMaxTexture3DDepth", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 26 // CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LAYERED_WIDTH {"cudaDevAttrMaxTexture2DLayeredWidth", {"hipDeviceAttributeMaxTexture2DLayeredWidth", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 27 // CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LAYERED_HEIGHT diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index e3c10766e9..e6cdfdd304 100644 --- a/include/hip/hip_runtime_api.h +++ b/include/hip/hip_runtime_api.h @@ -116,7 +116,10 @@ typedef struct hipDeviceProp_t { int gcnArch; ///< AMD GCN Arch Value. Eg: 803, 701 int integrated; ///< APU vs dGPU int cooperativeLaunch; ///< HIP device supports cooperative launch - int cooperativeMultiDeviceLaunch; ///< HIP device supports cooperative launch on multiple devices + int cooperativeMultiDeviceLaunch; ///< HIP device supports cooperative launch on multiple devices + int maxTexture1D; ///< Maximum number of elements in 1D images + int maxTexture2D[2]; ///< Maximum dimensions (width, height) of 2D images, in image elements + int maxTexture3D[3]; ///< Maximum dimensions (width, height, depth) of 3D images, in image elements } hipDeviceProp_t; @@ -295,6 +298,14 @@ typedef enum hipDeviceAttribute_t { hipDeviceAttributeIntegrated, ///< iGPU hipDeviceAttributeCooperativeLaunch, ///< Support cooperative launch hipDeviceAttributeCooperativeMultiDeviceLaunch, ///< Support cooperative launch on multiple devices + + hipDeviceAttributeMaxTexture1DWidth, ///< Maximum number of elements in 1D images + hipDeviceAttributeMaxTexture2DWidth, ///< Maximum dimension width of 2D images in image elements + hipDeviceAttributeMaxTexture2DHeight, ///< Maximum dimension height of 2D images in image elements + hipDeviceAttributeMaxTexture3DWidth, ///< Maximum dimension width of 3D images in image elements + hipDeviceAttributeMaxTexture3DHeight, ///< Maximum dimensions height of 3D images in image elements + hipDeviceAttributeMaxTexture3DDepth ///< Maximum dimensions depth of 3D images in image elements + } hipDeviceAttribute_t; enum hipComputeMode { diff --git a/include/hip/nvcc_detail/hip_runtime_api.h b/include/hip/nvcc_detail/hip_runtime_api.h index aca04067e3..2c294c91a3 100644 --- a/include/hip/nvcc_detail/hip_runtime_api.h +++ b/include/hip/nvcc_detail/hip_runtime_api.h @@ -772,6 +772,13 @@ inline static hipError_t hipGetDeviceProperties(hipDeviceProp_t* p_prop, int dev p_prop->maxSharedMemoryPerMultiProcessor = cdprop.sharedMemPerMultiprocessor; p_prop->isMultiGpuBoard = cdprop.isMultiGpuBoard; + p_prop->maxTexture1D = cdprop.maxTexture1D; + p_prop->maxTexture2D[0] = cdprop.maxTexture2D[0]; + p_prop->maxTexture2D[1] = cdprop.maxTexture2D[1]; + p_prop->maxTexture3D[0] = cdprop.maxTexture3D[0]; + p_prop->maxTexture3D[1] = cdprop.maxTexture3D[1]; + p_prop->maxTexture3D[2] = cdprop.maxTexture3D[2]; + return hipCUDAErrorTohipError(cerror); } @@ -858,6 +865,24 @@ inline static hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t att case hipDeviceAttributeIntegrated: cdattr = cudaDevAttrIntegrated; break; + case hipDeviceAttributeMaxTexture1DWidth: + cdattr = cudaDevAttrMaxTexture1DWidth; + break; + case hipDeviceAttributeMaxTexture2DWidth: + cdattr = cudaDevAttrMaxTexture2DWidth; + break; + case hipDeviceAttributeMaxTexture2DHeight: + cdattr = cudaDevAttrMaxTexture2DHeight; + break; + case hipDeviceAttributeMaxTexture3DWidth: + cdattr = cudaDevAttrMaxTexture3DWidth; + break; + case hipDeviceAttributeMaxTexture3DHeight: + cdattr = cudaDevAttrMaxTexture3DHeight; + break; + case hipDeviceAttributeMaxTexture3DDepth: + cdattr = cudaDevAttrMaxTexture3DDepth; + break; default: cerror = cudaErrorInvalidValue; break; diff --git a/src/hip_device.cpp b/src/hip_device.cpp index 746eaa2ddf..5c19ad42c6 100644 --- a/src/hip_device.cpp +++ b/src/hip_device.cpp @@ -276,6 +276,24 @@ hipError_t ihipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device case hipDeviceAttributeIntegrated: *pi = prop->integrated; break; + case hipDeviceAttributeMaxTexture1DWidth: + *pi = prop->maxTexture1D; + break; + case hipDeviceAttributeMaxTexture2DWidth: + *pi = prop->maxTexture2D[0]; + break; + case hipDeviceAttributeMaxTexture2DHeight: + *pi = prop->maxTexture2D[1]; + break; + case hipDeviceAttributeMaxTexture3DWidth: + *pi = prop->maxTexture3D[0]; + break; + case hipDeviceAttributeMaxTexture3DHeight: + *pi = prop->maxTexture3D[1]; + break; + case hipDeviceAttributeMaxTexture3DDepth: + *pi = prop->maxTexture3D[2]; + break; default: e = hipErrorInvalidValue; break; diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp index b6c7e3d87a..1b11b3c67e 100644 --- a/src/hip_hcc.cpp +++ b/src/hip_hcc.cpp @@ -43,7 +43,7 @@ THE SOFTWARE. #include #include #include "hsa/hsa_ext_amd.h" - +#include "hsa/hsa_ext_image.h" #include "hip/hip_runtime.h" #include "hip_hcc_internal.h" #include "trace_helper.h" @@ -913,6 +913,18 @@ hipError_t ihipDevice_t::initProperties(hipDeviceProp_t* prop) { if(agent_profile == HSA_PROFILE_FULL) { prop->integrated = 1; } + + err = hsa_agent_get_info(_hsaAgent, (hsa_agent_info_t)HSA_EXT_AGENT_INFO_IMAGE_1D_MAX_ELEMENTS, + &prop->maxTexture1D); + DeviceErrorCheck(err); + + err = hsa_agent_get_info(_hsaAgent, (hsa_agent_info_t)HSA_EXT_AGENT_INFO_IMAGE_2D_MAX_ELEMENTS, + prop->maxTexture2D); + DeviceErrorCheck(err); + + err = hsa_agent_get_info(_hsaAgent, (hsa_agent_info_t)HSA_EXT_AGENT_INFO_IMAGE_3D_MAX_ELEMENTS, + prop->maxTexture3D); + DeviceErrorCheck(err); return e; } diff --git a/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp b/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp index ec82861013..1e1076b806 100644 --- a/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp +++ b/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp @@ -109,5 +109,12 @@ int main(int argc, char* argv[]) { CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeMaxSharedMemoryPerMultiprocessor, props.maxSharedMemoryPerMultiProcessor)); CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeIntegrated, props.integrated)); + + CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeMaxTexture1DWidth, props.maxTexture1D)); + CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeMaxTexture2DWidth, props.maxTexture2D[0])); + CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeMaxTexture2DHeight, props.maxTexture2D[1])); + CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeMaxTexture3DWidth, props.maxTexture3D[0])); + CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeMaxTexture3DHeight, props.maxTexture3D[1])); + CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeMaxTexture3DDepth, props.maxTexture3D[2])); passed(); }; From 38254caf7ac8dacc17dfa699716416b48a843394 Mon Sep 17 00:00:00 2001 From: wkwchau Date: Wed, 17 Jul 2019 23:19:29 -0400 Subject: [PATCH 5/5] Fixed bug of determine max block size in hipOccupancyMaxPotentialBlockSize (#1235) --- src/hip_module.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/hip_module.cpp b/src/hip_module.cpp index 9680299540..594555e2b8 100644 --- a/src/hip_module.cpp +++ b/src/hip_module.cpp @@ -993,7 +993,7 @@ hipError_t ihipOccupancyMaxPotentialBlockSize(uint32_t* gridSize, uint32_t* bloc tmp_min = min(maxWavesSGPRSLimited, tmp_min); activeWavefronts = min(maxWavesVGPRSLimited, tmp_min); - if (maxActivWaves < activeWavefronts) { + if (maxActivWaves <= activeWavefronts) { maxActivWaves = activeWavefronts; maxWavefronts = wavefrontsPerWG; }