submission 188799
netvope · python · License unknown
Use it
Vendorable · source mirrored · license unknownView source →
No package. Vendor the mirrored source: 122 lines, June 9 Researcher Reciprocity License v1.0.
submission.py
curl "https://kernelindex.com/api/v1/implementations/kernelbot-nvfp4-gemm-188799?include=source"interfacepython
Compatibility
measured onNVIDIA B200
declared hardwareNVIDIA B200
architecturessm_100
dtypesfp8_e4m3, nvfp4
Benchmark evidence
1 measurement across 1 GPU, fastest first.
Operation / workload
Hardware
Latency
Rank
Observed
Reported · How evidence levels are derived →
Source and license
sourceavailable
revision digestsha256:9ad8aedeb9a306ed669a4cfc7ac2f6553f39c41ff9b81c421d4d39b90e52b458
license declaredunknown
license concludedunknown
authorsnetvope
imported2026-08-26
Kernel source
submission.py122 lines
#!/usr/bin/env python3
from torch.utils import cpp_extension
import os
sources = [
"module.cpp",
"test_kernel.cu",
"minifloats.cu",
"torch.cpp",
# "nvfp4_gemv/dispatch.cu",
"nvfp4_gemm/dispatch.cu",
"nvfp4_gemm/generic.cu",
"nvfp4_gemm/task1.cu",
"nvfp4_gemm/task2.cu",
"nvfp4_gemm/task3.cu",
]
def compile_module(sources=None, cuda_sources=None, verbose=False):
if verbose:
print("compile_module()")
current_dir = os.path.dirname(os.path.abspath(__file__))
include_dirs = [
f"{current_dir}/../cutlass/include",
# f"{current_dir}/../cutlass/tools/util/include",
]
extra_cflags = [
"-fdiagnostics-color=always",
"-O2",
"-march=native",
"-Wall",
]
extra_cuda_cflags = [
"-arch=sm_100a",
"--generate-line-info",
"--resource-usage",
"--restrict",
]
extra_ldflags = [
"-lcublas",
"-lcublasLt",
]
for flag in extra_cflags:
extra_cuda_cflags.append('-Xcompiler')
extra_cuda_cflags.append(flag)
kwargs = dict(
name="nvfp4_module",
extra_include_paths=include_dirs,
extra_cflags=extra_cflags,
extra_cuda_cflags=extra_cuda_cflags,
extra_ldflags=extra_ldflags,
verbose=verbose,
)
if sources is not None:
return cpp_extension.load(**kwargs,
sources=[f"{current_dir}/{x}" for x in sources],
build_directory=os.environ.get("BUILD_DIR", None),
)
elif cuda_sources is not None:
return cpp_extension.load_inline(**kwargs,
cpp_sources='',
cuda_sources=cuda_sources,
)
else:
raise ValueError("Either sources or cuda_sources must be provided.")
import base64, gzip
cuda_sources = """ABzY8w3<h00{`uOX?NQ=vgmjJ3bf}WL-{hYC42nxY{&0tbD|S(Bgviao6g~(NXp_h7d1(FNv`K#y+3e&*{#Atf)puHww&(H$TJ;DSZV<Zg({!`Y;13nH_!ar02qevy1rK(Z>$2|#@}quaHgh3j*O|PYr)iuY}cxekL8>&GHw4D3Y*r*_AK(@qCe2uou9j{4$-u+A4Zx6Urg)PHZ1ye<Ohajdk&P=w76berLZJm9eGfiw659>t@WYP`WV+vk`(KH5?Enq`(B0Q{{6Or|Jox(pGH0jBR!gi<Q<VK9*X3j-Vx-zKQ#=<s_?qs$P@lpRs0)G0^5s53N)P+1nM$X#os1UMWD7H9FV`i3jeOWGAl%n$gA++mHC^usTSA8s_(6+gO&RckaDdsZY%0SysY4pNY=EI?rBHU)TdNE?Vdwb=$Eopt$RCCI=^%W+DW5(de!SJs>`ZAZv4&iOnbDk5n1lU(Ie|9dYD+A?pkDUydev>;d^0(GD8L=@TZ=sO()8L+^g{0ojQ>{ajXjY24Dj#ng$+GlzZ|YaTa+)Y8A3A72S%zA8tG@Nq}kJJ`+8~32FoWA5JX89zBrI@$alKBF0$vu0c!yi)(o}K*-1r!U&7A(TO>`faul=LQBIR3P?D)2Lc3ons^Vr@VK~|O7P<$Tl+yJ%K=~?U$bz&sw4)5u;P{;`o2S;t%r_f9%D`z1yds;q2W&~Q$y_+Xc06jWDzp9C!Jf%iy#GtMG&d*;ZF`(CbU_=dT5Ja=ka8uC(yl{!wvR34Oz*8gu#+Z@+Xjz1+A$uSgmrv;1X6}VNfcJG=&0yA6QGFRpAHLQ&?qI_nQD=k0cpWc)L`Rd_q2SFS*tgWcDn>4@`wO$WEnlC`v!FFOQt_e`Q#&tO&0^Nu$jx+<)+lao~IQSBut26(2c1=uy{p96JP{o{5vg&RlpMN1ue2IRI9a*K$o&mW&h@J_T&PG7ms!lgRgP$gBNoZAAX(|A$n6W0$b5Qd87y)vVo7PP4+t<6&B!WzA6}&Gk~^Ja~=20wKmw7G@4hcWGh#k}uHpauO||t2uzN@ykgy1)<P?=`XCUu#c={FOpkcOSE!<%+_V0T_BPfh)YW40-0QaMABxf>ZYlMI!KSD$d!_INQ7(^=Bjr)n(S%UmV3Kx+TlcxjPa^$*A0|=0BKkf7U-=FB8GwOO{Njm=E(o8Z<|7?>W4E)!`Y;8q~|3{CAWydexQ1GldoR#^(OgxieIsDIYiMSiXfDtc+KdB9@%)rxjK){)%TCP`u<5Q_rYnS-{<nn4N<a!KQ!pRkr(}AHls0(96byV4*JYzpjJ7IPsYpPDp0_Ts4%e9AxYsoIJk-3lV>?5;Eo&|7=VTx9g)rVek3pnKVuU!cp5km3_bTCXpsu732q24E9~|hlOA#@KCrf(!`1c06Ab@lQ9Qo1L_KcgxwTIk8LJxuLex#2`i-EDo+{?K5-qvd-&P>FL*1l6IdQS*&N#kYY;)xk1(~3ik(K6ovb#DlT%ELWp~{&m+`rO{xJ9n5NE=Q^BP&osu#(V5gpbPMv!_h=C8kV3aRbI8G_Gx0tC}~rw_g5lgKUwOKY0l3>v2RBqeALC^<DA?KmI|^f9|%s4br;kU0(DWgYLySc}*H8C*9L-W6<eWNyBjntr~*Y9a_PyWmXYn|FZqdn^W7cywG~nHNi@?N45p?O0(a7v-_sy=+n@Gf>^d^L0>|{4W|(<H|U;8AQL*0g#I*uuEUg}?dib-8To+=cI}-VjS1)z{Pd?0LUDZ)P#8MGP>CK`WMT!b9YvN24w!!n{bG`648uGcTL8y#{5xFdK%>{jN+GS}TG0Wgt5(gB77~9XXfk{g3QxlbR~vMo6`|@w{}!`Eb9#ORO`t-a9U??9vLl2c<Mhl7L4eV4bo4r^oRIn~A%IY75kUy>X-<K-IkX@=BxhM#2xrmsjj0RWN!uBLL8Wcz2_G^8;({j{*t!!YO-{Rt)|40!ob6W0xn<L;ScLkoKofpXC?ETd2?abqDMWk1j%cF-0}QSogh0Rp8Cobmz=}^mYeTAqq5=fDeq<5WMj)>M(*|)RBglwbCG<zp9m)ZhEWSQw8|cKNzy&B#9+NG)ek@oSeCYN`|KeousnP2Y_<q^DfYH-wljg6G*?~cH`D?HH{=<NLxHxThdVSJ3Z$s+&px13)4K5&Iv(bl|n+Oi`8t1=~&M%j}PQOnsdZc@HdD;bdfS}hnA9Oo?m2}Trr&sOn`FoWB403)kAgA54?f{AnE>y&ez-DWbixYCz>9sz<&qfo5^WayC@}xUBM|3BEuR$&wy+OBib=v5W%PUX;{SE=<5MsOAZ=E)}XPtHxXaKyV^K<8XK>8mVr>D8>A_pmEn;jsv(LAN_C_Z4f-R*T+1LQ0D)`C_9Do<6=DxFprzjS`-085SDuPO)I@BH%$3PC2Nu5s3Q4@@cxng_b51=IBzQVtE$ziRde-NDtML*8Fpw6W2tF6;Dg4Yq$sPA~ej<*xc26_5-X6g@zJW`k__-n{B}X#;l82c2H;>JoKk1)AU!v=4>WYC!EaZEVn{$TPIoMekRHi|s-CLM5L*bRfNljZ0gtfo;);)@u!<Vh9khLVFG*X32Ty{b~1o=e*Uy+zSNssoU>Vp!d3cEYM|`KQ#dD73Cef3<#qDPtwAy(!M3#6Vhn^+(qJfc@T(xmkSSVzSalc^n6gf_}@0nkw1~&Mz&*7OX9a-U|S<{71<8xG*A;G9QiOeb4|~TANoW0d=I7zhqR^xi?E5C;368Uv|GIgWnz1d?@s_ipu1sp1MsDH$|7sq(H}rhRL7IaaW)Gz6M(RSZ9NF|2U-S*#JxMnJ~*JM=}PEsVZB@WkMA&TIX>o<m@kM~2L~6RgLF`OXwJOwpKMPGKk?gl!3%8nY1f3YMjoQohts{4?3+N_QKGtDG&tT6%M8*6$_8t32I~lmsk{^5`%&%qc*E6iEDe*ilr5($sz^eS&>!m)>j;clmH!8Mn53P3JY~Bx=%Lt#29#QKWfS}gJJg`0bsgxAc1z(Ot5lRX8A>*_kM&zLgNP_f1r;VMm;+Igamov=h?o!1Pbk#RLzVO?GA`L6&i^EwP9}Z;7EyJR8q#c2RKuwjojXiotbp--v~#SmtO|GhfO10*BSxf>6s$qPm`@6qZ@{qXsSU0rJ4$Y8l~*d&Q8^WV*p8}|?Wmm+9rgXk^veurx=*&hr$tX5*$IU85T705=(E#~gZfkW{!WHlmy6TVBcjNqk8wl#n^;kmYAr9v=mE^3^AY~D9B=`a#I_UHenDlntdP!wuvTw7Ya#t(75R1GPbUS6oE0oNQj)@1m7M(ADQmNp1h&%gBXRp@SM?I9&$Q{5>f&@y<o8kOE-ne|2k+wAa$WrAvj)-cmfYTuW?oFA$uvrwU*>Ng91JWk^n>&Ehfgi?=9os%yxA3gywddss#a6A9U2XC?gP{01i+6bfj`+cp-W71V+Ee&sBweivLB$Bx@_$bx`^Ro3*5K>dj6l6|NOYkaQcC11qht4ng$k-Z;8$|eBH4^gF*T)l!)KA@2Zr9yQPpLh8X|HL9fp-Q?!N1ZmL-ClSM3##?`XL&3LJoAh}FuGXW2)xkimj_|{7gT6ds5(GV{;prx6wEG!t$QZ2fgf@6BOu0^UKnU8CqO3`d_h9fRZ5oZhMExJ%S)%?ky)1FGx!?+{I@5t_yW+wr6tRS_MXsL|iK*h+bSLoJU%%RjypJ{<NwJ+q|JFu7%odlAnS-LUiI)`?AX#IY~%GZy-J*wpM>n|)c+ZLu)E()6hm%H+_{U3mFi)6|{BD0A&y``r_g?A|8Dd_WqgPTt{b}Qun#hcMG@|`c_AUA_E{G~|8;#h4Mk$p=yRY*iN8&Q-ym5fzls>G|PL<3`VBs3UT9~)eON|lTBq@*S5v6fUmspl2Yla=SWG^z2*M7`wV+U%e&2{CWzv(mfqF!t^YUMM|uTk!MZPdakb(hnSbu1n);zBsEH8!rY?`;skLD=4{BD5GU$>_d@|eRvr$dnsXCR^;{)ksJT@We!;Qyjw=-iuC(#fo7Kl(9RP;#9V)4e-R4&>5|x<BZ}>p6vmeo$z=ueABkX^e?7^Z6-{K~^WJl0bML9LNm1`E6i$@S6WfU}pn*^0H!4xU6+!1SP)SrNp2m5_rKa#0dZn_8#q;9rlc!qlb0C(P#BajCyx0N@&z;2)#7<@wujS6-Sh7|;i!Ttd5A$UVB>08J?4^WlvB)Kp_)YKw$Q<xcZWf;-bU{Yyij-=7p;E1>Fs~P(K&NfNl)pf1+l%FvYToZ4x-Tohi^VvZ@^7LaL5`7$a#Q{sL5|8w@@@%9-p!ZfsGKMl`1&ZOvcJ!a1bL~D*X>5c4vFqSLuc#^fKgGY(5Ed8GXvx~3kH--Tr{ZQawH}1sfLS0E`vttT+QBU(IUd~hWJ8=KjuQjQ&u8y`;>81Y7K36V9j5?SZw+v3mEZ!YFV1vGEMCCEQc6NCJs9=uBxTU-dTd|+Lw77xldnv;p6qgvLp7xq9IEy!S6F_Uux_wV-k8lg7;ivVdoiRc-bj6Junuo2Z=d#WkY+M9os?a!9AWC-OF0WWo%=LMi>+_`zu})i*-bMSuN2nqb8Uk@*OqB%j%0|EbxDyMYk8w8+%XJ8>zLrl`9T&Msq~vRfkz#ca&3{`Q8Y_Q&xG<5Q+c1{AbD=@xD%e&-X$$<g*B6y#@;IdxU{68r>ome6iL3`^{f|?EZdJe&sROHU$84Ruh*uM*k{eAenxXA3zR(hw`g|xq>jEqO`OmwEy~|AkF7re*{@TCd#j<<qE@BCQtd&AZ8Rg*Ao^3Q@9zX5*DIYVwLw}tDh6-m+XHrP%pX%wwhSB((r{$e*Wz04?|0vdU&q+%a{Csx7~8FZtJbJKV)kSq_;Oud~reL<t6z0$nh$qd9R!_uXpYKvb%OTA_fbFL^4Hux1GE1GB#E*LjF;1gse9N|I8bL^K7NeD7aa|R%-P>=8UY^P`Y{6m|ASbEOU@&K?qz~_-k8frM6P7w5_zWdRqx5hGOw$hSG~l@b{79)k*W(R(h&Dug+H5c~M&lO{KX*A~i+4<f!<bV`FtAWNj-weT=NoR@!-vt@ItvqGqm)<b!c#_eHwCId!b6F`2CT&6df-(Duw)ZCiZ7hg&vM%+MYe*&r<R63t-hX^|C18r$8W4=UKu{4CSDw!;W7KlcOZZBZ9`obsWO*Dgxk<*&`|dAnBA&Mw+lryXT*(QAFsI=>7$=l$-*xpv+->!`%7u<J*$8LFl=Qa16v$Y@g~uP1?zx4~-k`T&h4tE?owJU>RO2iH&w6yRp1IRDNM9MC)VS8KDv#)Lq}Sw)N)wpFr;zxWmxA4(bS-T2XrRY-z_ylPBWbX({)i3MVFY)`7=&AhjJH0|2)hoGG_c8}_8_tRku(8RL^a*0=Fg+5qdv*ws1L=*~rvf~=YIPDRIDIuCTbbRARPL&`W9UaAWRH@|g@v&ml!a!Z>9{e_aM@YU@zaq-vMZh`Pl>BU4m{Zq%5NuF^m3f@WU5Y=s_B;kxDrc_k+4MbuaAjhac_?Z$*~>f<#ZqhgR0U;Q7>eK5>EUA0kQ?~qk9VuF9DJHUzqZ76zFROus<gydcSiLsbdgnlv5W1|HmhuW6A9)$s5si1UGu(rPqx^Pw(VB&Ih<%5Sh`8GM9czKIf@}Ad;L`O>20HL$<U72b)(*ADE>thNvhX&D)Cr>La)p(_+Z2RtDX8i8!P6$O11(EC;MOm^BOx`O&KIpFSM^c%cL*8Y#l4FIU26KV%7=!_Z@32<2K?A7Wp(k(C_LBuL?$>cK);0sh=^)BiL#*2dYyrEEqt@Hth|($rxZ^$0^B`&ScBP9F=t*{ZZRzPtbT%7py!x2&IuA3w`}<G8mE(qQuppLz3dqKm3g%IgJ9Ck(hadhjLrc(s@z_4Ne)XsQER8+_7S{eO@l%9VwJ9Zs)g-pF&-F^Eg#QE1WiPD*A5rgFK-mCBn#=$)t2`TlW!rF)FPPSmxv*8t3QGJ2%rAw$GkW0R=+K2semOTa&GiDO2dHNrbo4n;L$q8^yPxGYOyu`mJRE%p-Aa)@$Y%sANuUPD39IaXlEIrEXGt{hi3UIYAw}zc-Hr95ej6I~Yp%&<!0o?6PxHk91@L(r09lmOueSCaJ)&iLC`K1U?>LE%Ybs4rdJ5fvIu3chPb=K6cSZe33Cv?6!qJAgadfu*SY%$x4MXRi!K|^k0!e0V%vvJN4@Z2Zfq2T)`{x@v0a8+!!YUM9R9)a^(5m8~@J3TUt4DlmP=ACm(roEPCt#uktp1<_V3VuwX|Nct<{ep{i^P-x=86B5hw31!Z%GS6O?1j|gC5x1+e0>jw{Hi@th8N#ZUDX9@3#)<E41713L&kR6Uv6-~(o{c5!9VG7LKLNITOzyNJ{VK;x+hl%Q2&DwK!M_zV($#mxA(63r-AhyPqal<a$$2d6OqpA?{m<zyEqExX8&M|68r2-ar*OS(|sZ$J2F9*TSr58bAl2}2nR*%Hs+taMNTN`=WeLV&)prWRDOI>Vk+tb%G!6Cv4o>Krte4^9K#O-T;8uD}@R0)2NE18&E+q#Q|{tEp4x0NabG6g7#G+tp|VDbXehltq{iNHVBLpnFpEYW4GqHdZ%S@Iqq*dlep@3FT$dg%{d{SJPUhxP=Advdu)ZgppBg|XiclVq48ze<yoN)?nMBb6ppC)04O4FQf)j~k2+&eR^oQW_-IAt|K^b84K+G>AkVPhqxvRvv5nd^IOQYq^u?fj$=z%43oCz|1(5@$dnSf++#cq{qc!RG|kTFH=ah=-3yJy7W5#yz2Hkc)N_ro*prM&_>C-C_k~RN}Nd0PWA8DWW}Fh)#%eD%4W~{eQF0(Hpob~9m~Y$O8!P69LA4`&?i91-5BzTvQxo#B(S3nxcR^u@H2G__Sm5LLLWAuviaGi%ZSbwoXdtK>(rz=;V2iA&^XhtS!5_&sIQ<pySIP%0IfQ*0?Wf00v{-o9vr|oM#{slJq12=xfQt62wS%-qFvBBpvVJsV%_Ac)GEVgLiKW>8hM~<`<WK5VeyTvjH*YEHhD%&nLHdEL{oe$l)a~dqR(8bbTu7DTT<fl&X+1RQc{JEm7D=POjoRxt_U5Hu3BHBY6k1>oT|dUmJ!E&mIkx89GGDNn17UJ#7I}%FRaM@TD=O>1x%4Yy(7EizyF))@uc}v!K8_$e$r!&n@rWhD=yYbshM1lXWiX4k@ba`+qvnt8AqhkYgaS$4$ILyd_KLyC(#R{k)fET<k2ifHYJ#rdM?R0AhJZi<w)ieJ)9wWSc>T3Du^B~O*CI+k;h@ez&>wvDXUELOT`OVqU*}#yCNCeg1!9`?BUy*GKa%r4u{KeI9xS{!)I}b>=_xcGh{|G)*$7=c)q+DGtvgI^97ABAL8q?I<<)Ty~^~3vLvyFw>Z<KGhx)-kXK+?pdXS{I87?lD{^5-WSgON=D7$=_=J+k+NtDM*^#T%QdMdNRcaNn>&FT6aBWm(PgEJkV4(t7OQfEc&8R-jX2~oKC6<N-EDh5v4JDR_1uPA7SrWOZQ3YDabyHR2%!E1cH*!nl@kcXd4rOFpM9R)^naW93tYs?J3M&?HJCwQgGL?r5xgE~qc9`XMnBsO=$n7x8?J&jdu#nsNeQOl=tRZt+%gf4R(MU6BNbDH}%o%Cc42dzLfGs1Jsl?VkI1o=3vZWKI*Qz?z9Foh)l81EM&LlIdnA#JwSwq4)VI~XhNMkk`RiDehqn_~ichr9Tz)KYs+iCnpqs&eEB}TD(O5I}PknO4C!%1`84|nZ;ti<<NJ05GtW9@i6-|_eZypET5Ii%kBp3cV0dm8_4ZiZy*Cx%Y#$8$DnKe)3|`?qp7YCnjx;T1a@wLhP;;jY%%kkoyl*D<HeGu@Bctoz~SyC3*y3im+R4(&HjCAbyxSgA&V?u8RSgjQC0mr34AVZ7gAjsqi?(DHp4+TeN6vf%{`tRlx%mWaQl;CY0r6dZ#=uF_A|hM7oN8c?!X$<3piM+JE%+ej~pe#w>=5g@e2PAVi<+KokmrJpL<#5j8^i0D3cVthwFBtUVo#27<*n}+l@7t&k8%?Jrh3QpqIiZJ0MC5}1HatPhLSz(!CvpEaSodM_Ox0ZmFYA0EA1fZ1YErhkY5VVAWZ87j(AheVo$~cyv9hOIYq=z@PIC_eXt~m6V={g!0^O^{a@Hh`rePxh#R|4tU@~pr%mVq@>P)n5Da_Cu*J_|~_wjymf9gVC&;m<(HQL<cnc~i#$JtFLIK}uRFXVG&I(-32^z-ZSZdKSc#&h6lFH5R<e^5I`7^NrT8x+()y^~9GO@#XZ9$IYTgKr6jt5WdywZ&mVtYWwhSZGZRe9-R&1JGH1xQCmON);xgrF}XTy7Utg0|J{!vXF%4U29g1*@9gD(++74x6y|~?7thLts3?M(%T{@uIV+Qh&`YQqPS;BWvGa2g1p(}fquqy&9-c<&R98<u2G1f3Q<1M|CLAYWLIy_CwJ0n+<~0qKkbYE1i9Z(i^Nl?7@3%O3)sP6mQ4e46g)C0eVOcPAu>tFX``x-~o;Lbgf6y3Q!N*mr)#>*ubBO*e6-q0@>lqjt`!o$K^2)?GxL4u9EAs$8XtX_I#0dD4?rBHU@Y`wkoPJ|$y|*I98J@@Ut|^6;IC@I_S@(8?2REM4ijU9lj1qGMJ%t!djfi}t!?o*;{P^4*lg%?da;Hw)3JvH?pPpm1eGh-5t)1xaRE@yCKcz=dp-#grU$(hK+17%xK=oxDsRiXHN!ZrhN*}Cg0G6|D^sI2|Af|CT1!od5PbAHz2q+(P&vX!>!@sg%T9?PvvX7hial|`6xS?Itm!?Oy4&s>TINcnkl7_QM*a|Kuh~#L4*{)>an9A!^f2v}fwJy#suLd1$w|=6s^9eot9v><ZEp{Pl0~=J&Be@*L#iceloL_pFDt&&q1pR$0YD6FnL+})pTqaHj+F4_8c6F+?JN=e6=rzv!4X}P0ovii!k-~mfDwUVV-W1q}Y3#$wOv4IfOycq607^=OemV0^D{i<14}A2xzi9nd0~FWKomOpMyXx$n?H-l@+<X?`muRpC$9Hzo?m)98oX1U<B)xR!HPf^-XK87sX_3aJc1S<ACeX02;Rl!>)b%10pPtkB9&LDe(LEn@dU#Yerg-wb8f_7B{(&L&EJEC+l1Wgv2lgcLgHXlO7WlyJweOqkFbBr2n^pw#c*q{34e@XW@htX;EgpW?y$6_vVSvFL&j}9`WRX5}EE?x3&cL%X^oH+(7p${*bbge6<WD^luk58*ePbHO`yM0!$5Av154N`<KdgGUwrT5C0Je>G-!^N{ZF*d#I*wfDZ!E=$T-w5AIwGF=gzBc#YOJ33PUwj}8c7W^4a^)RJ9q|evk<O_aEDKY+e*Q$?VnIWDNI8eHHjoh^Yxg{qlf(AVeF=<V`4b=Oc!JKgZD_gH$K}G68L}&T+@Oc$G4K5YZMd5GN9`d%y+iq&<UA|*-$r~2Pn&qSi<Z9D!j$1z>`uThJV|J9X%jX^yy^w2s4yzsx9^^vva3Egpd!A1ti6)URM*WrX2tsfC&T!00+KHj~$=sj^i7ZlA|;+bwq<F?N@%RCs0x&J^?j5+%J2b6Zrm%_UWScv47cUb+qQMgH9iYe9E+mhkx4ODV+NjnqlGRUA&Np+G2^uWeW?I$;;rXVOi!IAQWKV$B320fXk>nKK8^=Nqxv|g%S8!s&NG_O=HqM==01l5Fhp{;v(%GonH*VL%O`Y=nXpU81f_koHb=<T2<UH{>b{}Uy_L{nWsRhBw&ShH?&t2_z!YfnOWLqm0Sy7)MQ2EO3lTUnlqIeiz_u|Dy6!;lH(GIs_BBM_LDsXdfqNcc#VtN?DXGZrtE7ja(?h`Q8ng@a5wHfICw1|OpE&{<HyaoC0Q?sE!)MTGn<{~u0WaH`qk+v*&?-_`X2w_<#9pI)h9>v3th5cD|6zXW0jSediL*A3%3)F5+09_`F$EK9R!}M2#GmK6pM$o(1e2HWT0gZT{BgyBnnH-xMgugv6O(df{46=X-HthA9M{tB!Dhy(PBSa0Uo`>b6`?loSK%Gf8}#2wMNx&EIm+)m^le7D^ap#bia_|Y&$*9Han9=DP=~6<-8doE~=Q3e!6-r4V)bRIcXVshVaIE5R1TU=L+RbVFt!c+yv}qA=G+ubB@3;Qvg6}=-YJRNiJqVsfw(L`5Lwzy*UmoJVjRps+Ad)%aG#2icBDrdRgM7>4wvtl|lhoZYTtNNsoTxw=@MQ>hV%9!$Lm!xO4qmT$CO8iHHb}7|djKSEv|zg;c780bf>DXNwUdx<8bc{rT_5jk22*9CQjb00FLL0uDFwS_gNHqx))Bq6&5bggJtz&GE=Bcy_qWFw*hde@GN?0>JD<*T4X)^cu`mymt-s2Q$=`cBYWRH^)!&=NQT1C|S;kQ(_nE1qo^v5;UA4NLu8eaZX?hPsX@hR6CRshv3sC7TY4y6`|OmcjdLVS^9aw*A%JfYF(7&)*MS6id1IYiJC~6C}M~rExBSQRa9hY86&HT{)b8lm!_Q5Pi8Xn2oHI0to%0J3k{lhthj&aK~p57f7qZY8pVIAK{L!9H2R-*&=l@7#=9UdvEfL!gI{2~5$0{{-moT7dJ~b}Gq}w^vpbBDwaqUc7*8g7_4r*B;GoS^rTNyxcKqupfT1^FZNtz8^HY<S%JS=_qCB*T@7Tse5e>Ag;XAk>W8YfRS8X`89n&hLKBn$Twb>%#cIei2b~Jh!6zVR^)Yt{8bIE0~N%T8fbGKm5>yD<I`XoYxh4q;7&eWTWiPwm1@ZUlh;26IZXcIlqUE19X16^3CMCxivW^}WR4!gskL8J%G@U|;<b|E^76KM)P%%DXyz9Ln>xomwNJ=GA4zA$@@?CUAhLF5e$-l7X}bi;v8D;B@igdrAxF$3`Tj^7x{JkA(3F7Ilf-a_bHelcoHph_|nPBGAnkw`MA*pfkGpkEW2U2)-V&%Zm<|H}^wwtj(f27oO$<04|rJ=MM*u13;X+`6C8du2!iR}DFP=x;B#s|*;=*wsY3f%jyRYd|t}^uUfDwsip-m%?L0Y>$9kvTN(aT{1+OPH+jewoir+krh^W-!)g&cg@xJU4GX!i@L7mJ5$#y?7f_>YnJZ1=F4_n^C?}|+HU7|-b}mk9%{!u#M%w-8`7e~5?pWqNyD`>4M!A=r>qoVZZXVg{vIyZe&;@EF;k;J8d|H9hSsb!<Y8>D4r8_%H;M(YjVNpTOGc#3NMIs$6n`<k7NI8TO3}5r;4+VdjHaG}@~S&rerM)0!<d4_OYhyri>pvQ94<x8jNsBuTg0df*z$l8;_AB~a2A8|J@7UHV@#z=`jZ4qXZH`HjdM63F9NYW(>4G;-k&$N=pC@`)Fu5hc4r>nFb!!YlI59}Nk$g98i5s4B<?$J(HrCoD7u)?>?WFBbO+rO7k?eb6vY>(v%Bac{W>iKmsT+1-4BzFi>4-}To1r?6|as8s?J<}05(*g-fbjPg_p@=;*nPW$G)oV@2U9ZctLBOoi+XrB;iF_=(rKQwe)^}?;ffqs;uQqF$#Y!P8I)OJ1(P2862No^u`QwX11d0*P)dtHX!A`eqR%szE7{IBFu1?_qd6$o{8f5>@rj~Q*+my-`AFKb27yx>o}4}?o$B@=3!ds;NX<vYbEH8HqsM{n*PnlO*S>Z6eP;#St~aAJA$Qmboijrn@42gBT&uDzc-n?APD6HOHIqOAV?M`HPz8-^3?7;yn<AO(wIhmhEqv@sd;JHc9NzvRno+2w$-JHNGjJ}lIT8nL@42lHJa!Z;lB{(5OSOuxuU7Za>sy9$m{Gg-S!;cpX7En^MY2;(W32D$@-UCu8}F%R3%eWD%YITzsrl%EO&j-j#kT`0%YiumLA+%K*SM$W`dr)$89%71fPipn;kn;H3gR0XEz;JD#r!rnHC<n`X@b@(A!yI%KmDZ^}c(4JKCdj&K#-WN<2>{ihCc|GkL{#S$delKRCcLSHiazoy^oNZa<3u_%<X^(q*%r6_U%qA431o!eWDMM&8t!9MN<l2aVO^4AryrvNePZy$!a6aFj55EDjp=;*c1A1@8OK*;(cmLOK#+yzJeO2Ay@{9)C}fC-B8vDvd+gL*v3utD2VK(Bsh3s~ya6O}*oRji&6V@V|;GP=9-5>v1yB{z+p#J+n;?^NFaMm<QBsASZ0HVhdSe>UosFvEY|U_<9a$1_=>R0EQR7oHpW5iBvPq)-2a-&75XyKC{`H`OVgRdb2fOwAo_z8TqYpE!WCvxz;mVu9a>%me?*wPdkU6_A}{eGkP|pmbbF@Ju{~1t;6Wi8FZ;6<y5eW?+R8sIN1$6|CE9rf>j*ocohNF<U9>go?Pxp4g}PbRsolkuvGQKT&T_GL2Zim>w@g*USwl{Ab*mITlvsK0NzBzr~~;UR$&qvW3U}BEt2-d2&!SWaWeFM2l~K{_!Neoo9=`>sv$#rBe<Tj#T8wATNp$W=f_bIN@ecG9lzpNqDhWENRtXIT8DAo$PU6t3#MM`oYu?{FWisMDI~{!uxB^Y;0no2uLYdribMVFVQ%IbWa2Tk{M?Upv10e_UK|@JrZN8HQ);QY=-se*{!z)dGnIU#hv5!W0Gk3}vpm2i1$dqhuq^<#%L8l+3ePz}4w36Gw|0NubhwAyVy^DegDDS+pgft9nL^Xp?3`ES;DByIfbhC_M4G*d2E}J}AY%ucC<!k~Cb1Y%b2&uKMTj`2%yy}o*(4>&u*CpU>F8R@I!jq)7~T&C=KJ{YH9Iz!pnsH;ad|-aCNVDQjM&UO4MKe$9JKXFAFyC_e8c9LZwC)6YApR~2}sRCNO}UPSq2iDdeSg(w_KIv<>_#A8BQ}f|CVXZ*nv!X?hVHL^W|MaUo@ylRD2_R<C=E;qf8)I$uQ_1%h279sWp<8Ym4x%#~G2T#(Q*CSHIW7JNxTb+__@`lRI*#`&B5`+2YzN`F&~y4|I-Z(e24>60jTGO`6Hqe!+u`Yz#wZ#=k+AM-nf3HgHoA@}BR4M_M{diAhNX#AWIqV^nhSHAe${6T`aK>B18yJ0x;^<=8gH`0fL~MEGDui4$YRhb?&h&D80D28L?l^q({fJ0Y7CSt7cVJ>8T-z~7}{OKcP74YCD3pf))#0zCkv$03;Pbxej7wMW&^JuT)*YSR?1Zaj%+tr@y^mdHA8k|YeLPY>lFPmBXa60B5EsKl@z-~j{=SKPM=c=GOaLhMLhi^)ZT20`9Qo=o!>y{?vDIzM_<wnikIrO*;VOM5yDLGI8HhVrN_=mCxC#H8=u$U|;s0J~guSMkxDM&7eJCDirde6Y<qVCU)X1z_3Sy{sx<2(-C0&^$`TiXsKto&!==xhZ{_^d284EDC}S7xYa>pxOcIn|J^$kYOecg@dG&>?aYj%fzzsFxkW_iGyU}OX{+O$CnCGs;@(o)*(uNc!*Mc9ip@jQCf#6twWU7Axi5IrM>SPp|rP-P+CVQ{k0;L@P@bV9idcH>HpUeO6v%vKWT*0x!-p!Hxq?Y>$~g7q;+IcRwmYwN$bd@b!5^yGHD%|w2n+#M<%T!lV&56@J{aH$RrlGv<^&K2PUlplh%Pr>%gRSVA47;X&sof4oq4HCanXLRuh=CyADiR2PXaDfl0gTz@&9x(mF6{9hkHZOj<55Y2H9yU!1hQIB6Z0v<^#JU!1fKOZxG{l5*ou)-g%zn51<~(mEz-9h0<<Nm|Dwtz(kbF-hx~r03q2RNvWK$0V&|lGZUv>zJf<Owu|gX&sZaj!9a_B&}nT{vBhI{uhxhL%t#-0RR"""
cuda_sources = gzip.decompress(base64.b85decode(cuda_sources)).decode("utf-8")
module = compile_module(cuda_sources=cuda_sources)
import torch
from task import input_t, output_t
def custom_kernel(data: input_t) -> output_t:
"""
Execute the block-scaled GEMM kernel.
This is the main entry point called by the evaluation framework.
It converts PyTorch tensors to CuTe tensors, launches the kernel,
and returns the result.
Args:
data: Tuple of (a, b, sfa_ref, sfb_ref, sfa_permuted, sfb_permuted, c) PyTorch tensors
a: [m, k, l] - Input matrix in float4e2m1fn
b: [n, k, l] - Input vector in float4e2m1fn
sfa_ref: [m, k, l] - Scale factors in float8_e4m3fn, used by reference implementation
sfb_ref: [n, k, l] - Scale factors in float8_e4m3fn, used by reference implementation
sfa_permuted: [32, 4, rest_m, 4, rest_k, l] - Scale factors in float8_e4m3fn
sfb_permuted: [32, 4, rest_n, 4, rest_k, l] - Scale factors in float8_e4m3fn
c: [m, n, l] - Output vector in float16
Returns:
Output tensor c with computed results
"""
a, b, sfa, sfb, sfa_permuted, sfb_permuted, c = data
m, k_div_16, l = sfa.shape
n, k_div_16, l = sfb.shape
k = k_div_16 * 16
assert a.shape == (m, k//2, l), a.shape
assert b.shape == (n, k//2, l), b.shape
a = module.fp4_to_fp16(a.permute(2, 0, 1)).reshape(l, m, k_div_16, 16)
b = module.fp4_to_fp16(b.permute(2, 0, 1)).reshape(l, n, k_div_16, 16)
a *= sfa.to(torch.float16).permute(2, 0, 1).reshape(l, m, k_div_16, 1)
b *= sfb.to(torch.float16).permute(2, 0, 1).reshape(l, n, k_div_16, 1)
a = a.reshape(l, m, k)
b = b.reshape(l, n, k)
torch.matmul(a, b.transpose(1, 2), out=c.permute(2, 0, 1))
# torch.einsum('lmki,lnki->lmn', a, b)
return c
custom_kernel = module.nvfp4_gemm
scrolls · 122 lines total
Source code from GPU Mode and the KernelBot dataset · June 9 Researcher Reciprocity License v1.0
Best evidence level for this revision: reported
JSON