From 888f4e46ae2e37014b5a70cf976caf08d8aa1f79 Mon Sep 17 00:00:00 2001 From: Megvii Engine Team Date: Thu, 10 Feb 2022 20:24:43 +0800 Subject: [PATCH] feat(dnn/cuda): add implicit bmm large kernel dwconv2d dgrad kernels GitOrigin-RevId: fcb7974d62f4cc032544d507144e121bcf01c565 --- dnn/scripts/cutlass_generator/BUILD | 4 + .../cutlass_generator/conv2d_operation.py | 3 +- dnn/scripts/cutlass_generator/gen_list.py | 12 ++ dnn/scripts/cutlass_generator/generator.py | 30 ++-- dnn/scripts/cutlass_generator/list.bzl | Bin 98224 -> 118391 bytes dnn/src/CMakeLists.txt | 2 + dnn/src/cuda/conv_bias/algo.cpp | 8 +- .../conv_bias/cutlass_convolution_base.cpp | 4 +- ...mplicit_batched_gemm_float16_nchw_hmma.cpp | 1 + .../cuda/convolution/backward_data/algo.cpp | 34 ++++ dnn/src/cuda/convolution/backward_data/algo.h | 83 +++++++++- ...mplicit_batched_gemm_float16_nchw_hmma.cpp | 146 ++++++++++++++++++ ...implicit_batched_gemm_float32_nchw_fma.cpp | 141 +++++++++++++++++ .../implicit_gemm_int8_nchw4_dp4a.cpp | 2 +- dnn/src/cuda/convolution/opr_impl.h | 2 + dnn/src/cuda/cutlass/initialize_all.cu | 4 + dnn/test/common/checker.h | 1 + dnn/test/cuda/chanwise_convolution.cpp | 139 +++++++++++++++-- 18 files changed, 583 insertions(+), 33 deletions(-) create mode 100644 dnn/src/cuda/convolution/backward_data/implicit_batched_gemm_float16_nchw_hmma.cpp create mode 100644 dnn/src/cuda/convolution/backward_data/implicit_batched_gemm_float32_nchw_fma.cpp diff --git a/dnn/scripts/cutlass_generator/BUILD b/dnn/scripts/cutlass_generator/BUILD index 361bb8856..64e61884c 100644 --- a/dnn/scripts/cutlass_generator/BUILD +++ b/dnn/scripts/cutlass_generator/BUILD @@ -13,6 +13,10 @@ genrule( CUTLASS_WITH_LONG_PATH=true python3 $$GEN --operations conv2d --type simt $(@D) CUTLASS_WITH_LONG_PATH=true python3 $$GEN --operations conv2d --type tensorop8816 $(@D) CUTLASS_WITH_LONG_PATH=true python3 $$GEN --operations conv2d --type tensorop8832 $(@D) + CUTLASS_WITH_LONG_PATH=true python3 $$GEN --operations dwconv2d_fprop --type simt $(@D) + CUTLASS_WITH_LONG_PATH=true python3 $$GEN --operations dwconv2d_fprop --type tensorop884 $(@D) + CUTLASS_WITH_LONG_PATH=true python3 $$GEN --operations dwconv2d_dgrad --type simt $(@D) + CUTLASS_WITH_LONG_PATH=true python3 $$GEN --operations dwconv2d_dgrad --type tensorop884 $(@D) """, tools = ["//brain/megbrain/dnn/scripts/cutlass_generator:generator.py"], visibility = ["//visibility:public"], diff --git a/dnn/scripts/cutlass_generator/conv2d_operation.py b/dnn/scripts/cutlass_generator/conv2d_operation.py index 5bb0bd8fe..89b6d2582 100644 --- a/dnn/scripts/cutlass_generator/conv2d_operation.py +++ b/dnn/scripts/cutlass_generator/conv2d_operation.py @@ -545,8 +545,9 @@ def GenerateConv2d( epilogue: EpilogueFunctor, conv_kind: ConvKind ) -> bool: return ( - conv_kind == ConvKind.Dgrad + (conv_kind == ConvKind.Dgrad or conv_kind == ConvKind.Wgrad) and epilogue != EpilogueFunctor.BiasAddLinearCombinationClamp + and epilogue != EpilogueFunctor.BiasAddLinearCombination ) # loop over all tile descriptions diff --git a/dnn/scripts/cutlass_generator/gen_list.py b/dnn/scripts/cutlass_generator/gen_list.py index fc5f980bf..c67586950 100644 --- a/dnn/scripts/cutlass_generator/gen_list.py +++ b/dnn/scripts/cutlass_generator/gen_list.py @@ -3,6 +3,8 @@ from generator import ( GenerateGemvOperations, GenerateConv2dOperations, GenerateDeconvOperations, + GenerateDwconv2dFpropOperations, + GenerateDwconv2dDgradOperations, ) @@ -21,6 +23,12 @@ def write_op_list(f, gen_op, gen_type): operations = GenerateConv2dOperations(GenArg(gen_op, gen_type)) elif gen_op == "deconv": operations = GenerateDeconvOperations(GenArg(gen_op, gen_type)) + elif gen_op == "dwconv2d_fprop": + operations = GenerateDwconv2dFpropOperations(GenArg(gen_op, gen_type)) + elif gen_op == "dwconv2d_dgrad": + operations = GenerateDwconv2dDgradOperations(GenArg(gen_op, gen_type)) + elif gen_op == "dwconv2d_wgrad": + pass for op in operations: f.write(' "%s.cu",\n' % op.procedural_name()) if gen_op != "gemv": @@ -40,4 +48,8 @@ if __name__ == "__main__": write_op_list(f, "conv2d", "simt") write_op_list(f, "conv2d", "tensorop8816") write_op_list(f, "conv2d", "tensorop8832") + write_op_list(f, "dwconv2d_fprop", "simt") + write_op_list(f, "dwconv2d_fprop", "tensorop884") + write_op_list(f, "dwconv2d_dgrad", "simt") + write_op_list(f, "dwconv2d_dgrad", "tensorop884") f.write("]") diff --git a/dnn/scripts/cutlass_generator/generator.py b/dnn/scripts/cutlass_generator/generator.py index 6b4343161..68c2dfc5f 100644 --- a/dnn/scripts/cutlass_generator/generator.py +++ b/dnn/scripts/cutlass_generator/generator.py @@ -1056,7 +1056,8 @@ def GenerateGemm_Simt(args): return operations -def GenerateDwconv2dFprop_Simt(args): +# +def GenerateDwconv2d_Simt(args, conv_kind): ################################################################################ # warps per threadblock ################################################################################ @@ -1121,10 +1122,10 @@ def GenerateDwconv2dFprop_Simt(args): tile_descriptions = [ TileDescription([128, 128, 8], 2, [4, 2, 1], math_inst, min_cc, max_cc), TileDescription([128, 64, 8], 2, [2, 2, 1], math_inst, min_cc, max_cc), - TileDescription([64, 128, 8], 2, [1, 4, 1], math_inst, min_cc, max_cc), + TileDescription([64, 128, 8], 2, [2, 2, 1], math_inst, min_cc, max_cc), TileDescription([128, 32, 8], 2, [2, 1, 1], math_inst, min_cc, max_cc), TileDescription([32, 128, 8], 2, [1, 2, 1], math_inst, min_cc, max_cc), - TileDescription([64, 64, 8], 2, [1, 2, 1], math_inst, min_cc, max_cc), + TileDescription([64, 64, 8], 2, [2, 1, 1], math_inst, min_cc, max_cc), TileDescription([32, 64, 8], 2, [1, 1, 1], math_inst, min_cc, max_cc), TileDescription([64, 32, 8], 2, [1, 1, 1], math_inst, min_cc, max_cc), TileDescription([32, 32, 8], 2, [1, 1, 1], math_inst, min_cc, max_cc), @@ -1232,7 +1233,7 @@ def GenerateDwconv2dFprop_Simt(args): for alignment_src in alignment_constraints: operations += GenerateConv2d( ConvType.DepthwiseConvolution, - ConvKind.Fprop, + conv_kind, [tile], layout[0], layout[1], @@ -1249,7 +1250,7 @@ def GenerateDwconv2dFprop_Simt(args): # -def GenerateDwconv2dFprop_TensorOp_884(args): +def GenerateDwconv2d_TensorOp_884(args, conv_kind): layouts = [(LayoutType.TensorNCHW, LayoutType.TensorNCHW)] math_instructions = [ @@ -1296,7 +1297,7 @@ def GenerateDwconv2dFprop_TensorOp_884(args): for alignment_src in alignment_constraints: operations += GenerateConv2d( ConvType.DepthwiseConvolution, - ConvKind.Fprop, + conv_kind, tile_descriptions, layout[0], layout[1], @@ -1574,13 +1575,24 @@ def GenerateDeconvOperations(args): def GenerateDwconv2dFpropOperations(args): if args.type == "simt": - return GenerateDwconv2dFprop_Simt(args) + return GenerateDwconv2d_Simt(args, ConvKind.Fprop) else: assert args.type == "tensorop884", ( "operation dwconv2d fprop only support" "simt, tensorop884. (got:{})".format(args.type) ) - return GenerateDwconv2dFprop_TensorOp_884(args) + return GenerateDwconv2d_TensorOp_884(args, ConvKind.Fprop) + + +def GenerateDwconv2dDgradOperations(args): + if args.type == "simt": + return GenerateDwconv2d_Simt(args, ConvKind.Dgrad) + else: + assert args.type == "tensorop884", ( + "operation dwconv2d fprop only support" + "simt, tensorop884. (got:{})".format(args.type) + ) + return GenerateDwconv2d_TensorOp_884(args, ConvKind.Dgrad) def GenerateGemmOperations(args): @@ -1655,7 +1667,7 @@ if __name__ == "__main__": elif args.operations == "dwconv2d_fprop": operations = GenerateDwconv2dFpropOperations(args) elif args.operations == "dwconv2d_dgrad": - pass + operations = GenerateDwconv2dDgradOperations(args) elif args.operations == "dwconv2d_wgrad": pass diff --git a/dnn/scripts/cutlass_generator/list.bzl b/dnn/scripts/cutlass_generator/list.bzl index 596b4561b633d4834ec189c030e5aa1fadf6718e..76b877b84a2253c35be0ee8d8d37f289b5bc3fdc 100644 GIT binary patch literal 118391 zcmcg#(QYI;j(qR0X!N}fW?;A1sF!=bk2o9(-m<%GjK?-y?(yv3AC+08#FSJiCK)wo zw6(L7!5||p6`yw{m=P7 zczyqaM*HWd)UVL|Ni*){PeN^bN_bte)s(R@N5RSy7fjhBhUr6H?Abj2;Z;0 zS(*{(0%w+Hgxec$mSzM};_7Z->S%sIoL8&;>S*t&lbZn$ONG%?RbPy|J=(2r8{Kp^ZW95L6nQ z8kL@Agw!dWg4!X7rMI|th(yy_T{{GMfEx6?Y;U)b5rS%CyrShJBSfl=SJVv{#M0|u zJ4Cv|X(c;^%+n9=r|m(2yyy;IP`SvAKo^|q>Ii8=pc#7Ve_KPE5$FON$EvVq1ZoM+ zEX@eTn!B~Kb_gmAH=(UT?GRKNr=Dhn)G3{U+98Ogx43qQMAKPaI|O-v8uYYmJ8Wcx zpxPL(sQQ8(BGtw#>IRIo{5kz=he%gAtz?JD^K|3>&k9g1x*H>!5$FQ0>-LZ~1e&4j zpN%2S2y}ssV>La^2sEFaS(*{bHFslW?GRKLX+j%=+99YkPCd;CsZ%-ywL=g~Z*lDq ziKer2y3`l7QWn#JNc8o4D4`^3PZ@VA~ ziq3z^)&R;CtE4nwbSvfoWnxv;BZzrGyHa}FMN>)xC|9hix&dQcN&_eptEz6m=o0gQ zR#Ni8XJ~`n0Lm2;DK}trE9L=ZVp8=;Vjj@0l-_osW+i890Og8RRX1RaOKAXQVpY`* z7+qo>(5{r;cF~m50Lm4sq%>f3E9L=ZVpY{6hfOe(ywu`2e22ie8RdoZ#xReG^ zCRSD5fYBvp1AUb9S5U9*L-b7#sfwP7|&YI5b=PeYo(grig>`9)jT`x&s4+% z)~x1MA|9}2H9bTKP#^9Of;Fpob{boy83b!qbMV?~`g*{c(s+ne(hP!WX-4)v!_wHm zbTw_)u1I48)6&?kZ4$}S*ua|749>WUhzG1G%~m2Fu%GYF=n;j^QV9hwnLRnvIwd1;JbN*d#}O(G(V5v(cA;Ebz?c)*&{ zY$f6WYf9rGQb{ui)~x2>j2n|?5Ug3v!D}lb9J1Z!4v@Y;%q2dpWLhe##OAXrnH!5LRo!~@o>=2jveu%G1Bw7ubx$KCG6`@|zdRC-p2u|wxp>8c}R^64a5Lm!Ep-D)a`CM~7# zc3Ps19imIif!Qti+1b2aEj2MY)p5+Oh~p65YE{M<$7EVn!N(!Gv>cdS5zBecukYe% zKR>_Bv!+R_DdIRLSF0-cI7E{cRjNK=AID@`Rl&z0y0jdawTL5Io^+M8z9ua#q8XE` zl|?i|G-+uO&6rFp7u*cdrRBivis)<*O%ca2xmuNo;}G3yIWU=4Rqy~>4$Q8I&IZvG zaU7GYRTX?3VqC;=Or}*8d>o=n%Yj*mICLbIt}?V!J0@34M6^S6tL4CCT2k;pS`N&v zh|UI~W+JCLj>*-k3O)`oF5);Q)2a$S4$-CM!0d|XY!FQm$1%BDm5AdI-D){7nO0Ts z09p>rTEv+BmKM>B$<@jQH$x1FXvSn(x!`7qE-eRUS43xnXo@(F$Reg9hR0v-mNrrYL78u2GaR ztY#s{E*%qP467;1*riPxl_+CaT9lD}&8R3QEG>%Z(q=U%t&>GfSXva*rA-=H6cbic zl(9)x(QsfjMcGKhfz=enL8B67469klu}L;2${1F&kYkrtG#pq>Q5-ZXQO2;8D17v@ zaogj-QiXKVa9}A>0%(XRHmqhL$0k{2Q3qDDkQ-?@u$rPcXjGz%VKoamHp#|B8N+H8 za_rKIh6Af9ii1Wa${3aw#k5DLs^P%Wh4j&IU};hOG_oirtY#s{CfS%MV_3~Xj$K;O za9}k>anPtl8N+IdGB(MoYB;c(h1^KPfz=enL8B6-#Ht=Kq@%C%PZ-K~3jV$S_ua?g z@!@`d{CI!5r=M4+X-?k!(=xy+0^Qlo2qjy7FlrTHR{)&-qQhPi1%>NvNn{H2 zipj<+>{fHQksRoSyITM7 z19~&>h7<}08;qnGY%mf5WE*VoktVRgM~c7(8`1UI1RGpL^-m4(V|zy;1R{|E%r+u` z!UY>_q#10mk!G;LMg&mIgbhB@3^w>kGuU7x0?0NvXQbW!>Hg>Y-97!F2aM85V_K1? z1nAYqX+UobKHNNh-t7Opc|RU*<}cs8S_&r;6@^3TlN++qCLCoiV~RsQkwusi}Ahuy`7PJ&Q1T`NWy*NSw*nF`VO}PrjLYM!U|t zcPilYCEUAzTK@Lbc_a2}oj2l&H_p^~BlLQmuWsMzt?#GrO{Le(w)(TD)(-sDdONU3 zYJLvBXf|A()I~-<7E;6c@s{4WpT(J*Lhd%QA|0sR{3-pWDULMy8I)FS4DkEI^RKkl zPiw}!*h_bG@&!RZ9zUOsk08!4@a&4IyZ}?g;g-V1p(YDskhN&Bj8WLR3&u`51TL#a zl?`rb804&3JYy7g?ho{-i*|XFi%;9V`4g9S(TKxk)f$`Lz$GJ&Z`LfHF$#P3!?{o3 z9zVq5vL9km%k?4_IcpY+n;4_8#SioHM#SfVbHV{&3pd1ag#atPV)=rgXUQh2@e4&L zEDIJ2TDFBy;4D`ttZoXS(6bZNA?5Tjt~gQDi$v?hrT;;5`j(n>l3GOIlS(r#d!fL} z_?uBnGsGfiH^kyy-1H7lJSg9|R#TJP?Gc^bf+Uv=73SMf~*X zth+gQIefo<{G?A8(+TX_ZJXuE)4*nFO8|BR2)?ujC`9%OD1bH&D2TX-TuBz8DDz;s z!`z~AN&typcnHKE8$me;lmRBjubaiWcV5-=XtBOe_peSZsJkV}z+C!~#nK3!YXSft zTaHV{-8H#c4Dv~Tl{dnwLimB{W~#m$lD&{eTU?QW!(EY~RWW!+9mPA)6`65Sj7afR zA^gBUXj!q_;_&#%`;~69D3J3~Gc=`*=tZH|jbX$tcPa?(N=*oRY=ps?Q&hd%PkVkj z?{On?;gwL_g-=4Um4czjiM1a+-Jnm;0$tMsYEOTCoWC*8xA4=w;^}Fi`4i3i>!G-p zo}QYgha4%_>Cg0()9F(ZeOjXX^!#xtP069hDDvl~%S1|1#A|dn5uMQ`Vx^o4w_G5r z%S6kLIQ_Q91-yrhRC*5?XX<%{&ylgT_mGh|tL!~wwCufU@uxX;nZB2zAtIIDyF^B1 z-$Uk<-EfJlE)z|AZ(eEX@9B^Uxa5KYEu9An!|kCo9w=1GNIL-(h;zC~u0C2@qX!IBx(*nI6=K%`LuuCm1948-b--ZRb+hHvUHbI99`UXN1}a?#48sbs z>wuxO>wtkcr|ddlFzq_(a`d^COg;~At^);Hx(*bE6)If^3RSxf6o_-ut^);|&*TQ1 z>FHMor#n9zPE#%wi_LJEtS%QV8{#sVT{7P64KjJ^u;tzBR_6=xAwCv610QK>8K++a z!bi)#z{g6NT{7P39>rwByS9>)Vgr z7ohXyG74Uj+3J-XunbAY;Nbv*%U&a-We~+4=NsK=gt!c%GC=xgefx2B z`}k>oqmpHR;|Rdop6l`s`0^taC`8s46hMEt0tLxk0u(@h%mM`|`w=Jv=K(gZ?o8lZ z;1*upuQ;bT5}$>ZML`C#Iyl*5BV=)fxD28)K*z}4NvO*pDhG6ovNs8N8AN7*^^@*Q zl7I{%GXeODe!VZf7jfR+tnZzsHwK@6;j{HSIr*NK0#2(044m~nw<IY3>sx>s>bT@*HpDotJn((N6rap+k{AuNLk=1Bb941}@_ zBDf>*%U&a-We~+4=NsK=gt!c%GKlq0des7CZ=->O_cj`wH4W+R8hZNVUc57!9^qO% zXl#5gK*MJ(LENDMayJ06hXyG74Uj+3oVw!xbAY-iZvUa@R&F`cf@NUu#TuJM6+jty zk;{f7pfd2HXCVbr243vM9EqQsfl!t~1a~BU*=vNf45HZMT&X*a5SKwz2C@E`UV8%B z+$@lDs|jhs3v`whfnBy61b!A2fv#H)B0}8av=B90MkbLA?6YYB&W;g88MvWK!w9Mj z-0;~if++(xxVDSn%D^2Ub_rRCI#oXOnbSpp=tXJ)&-{Vrj9;cYImFsDz1YsOOpQNI zrGS$Fb3Z%RDK~<%ssQN8I0K)-H!{y!g1AEiNaQN z+{!IS7Fij{YP%7eMHN69c#+G7BcL+yqGurmQU+e^#2ksAn}JZ4K?HXse%Wh;v<#xy z<6NmbjS!bXR0gsB(dTA?+Yb$#8%;~70QjsW09-d4K!Ds00N_*Q<>M3} zK-q5q@XQ}*PTg^UIY4b1x3|$KZN%fQ~{Yt}gQETj;YK?HLoe(p9xSq2f@ zk@#h=5z;b&S(u5c2EGq(gTBRb-ZWi3LpeXj( z2yqL@qA=MqGKpkhpG^yJc8nm(zztm*Mo?wohR=LpNb*j8f{jXBJ2GH+pi{VxX0iK!0MiB_YU~7{eVJbZ9xHa8v+H%tq2rA zwed9#1#aQhZIN?|Bk@^iSvX}NE2xt_HbNFxh|3@<19XhsorJm! zqH;jTD0`ETmqBC}SU>5`BnikMGLv{;>C?0IkNBpkt-;eu0S&-9hpi4Ze{hJbFE9sq zklZ zbs0qEfR0i2CLu3_$Skma(w#{XkU?Z7@xEFdZUygc7U+Q8(;@;(9kxI$@TZf4fqNyU zE#@Ld78#iZK1%E&v?x-xl1wxiY@2b50PG||l|cx6X(U0GK}d+Ku;9udBtXgm9wdGl zS_nNUF>+7EGp7-YT_+E)aWTtXCl|N7iZAEpH#()?;u!ss?&-eC_v`tkQtZT=ed^1n zuf;umOE>+S>DwD`>2)deMUIEt8~U1V8cc7PnqElo`{CEO-@bgMf4&`leLUYCk9%gc zpKtzfoYIi!SlO58NSxJdPjs}>rRZ3k8JiUyPx~-=F}aUgd+P2`AL3#yeTa+1YL!03 zMXP;?i^Z8~AL8O^AD+y22^+#fEnNr;#8Q;C79%ebF zd=P;e?n8uOjTYWRglaku5r}iDd>0dJKBDCP*~xYEFKOSF+H2%iP>15Y=%e%{!ifVf z5sY)2e0Yg~m2Qzs>Ot|tSs!J;)(gT?BI28-9y4AF}@FV!F^ zD@3feUWiDX)uJ_q<)zOX(mCzCkkOYkQaaO2o^Z^kTFhYZy5|H3&@CYtoc#{7W65Z$8y0eEgh?Yityz&yalt$P$O7r2G9=-s{G_v^<``VOG2Pv2o9lwzYY zK*!J`qq4w9$x8QD?Jz3=0+cEP;Fs$iAcR&J5TIV~073M62MAEFcYq+v^$rk%^8g#S zUhjapz%876RBzP|>rsFJrPKiUWsd@c&;kPj)IACiME58_0G``WyDobaFb}YC>mCKn z1z=(Q0m6qss!3^$m`2>Z~yQ zqOrnaNMD6T;4H9et9j6Cs8RDE7=KZjVD!d32#!$AgJ67i5K`mDJOE{ZRa@1lyB?d) z@~Jb>H=r#Z?ab{05oA##M2N<;K!j0x^wyt9IRG&hY5*cM!~jI8`V=5WvsVG)H2M`F z3TK0NfSM=0&Ko&T05KL-0z_!c6F`(|o&aL7v(UOX<_RzxyaQCdx(k-+TyX39?g+A| z6Dp4jgi(6*Il6Fd$B&Zvmok7shL&JlyNRk@Ex) zV^JkQ1jG97s8uyjP?;>}2_Q~mo&chB=1EATUH`ta}b9 z2dGW6z}-EX>4do6r(q+M0;4iO$It?!vcN~lO80g^VO9bJC{+f)FV`|a2(2(6K)t2` zg6MS(5TIV$06~`P8z2Pd0XA;E&H;0QTR8Wq-fA7zqW}R)sR8iI9t8-Y1qKADdlVpu z?oogMJh!2CUG^wo9$@3vJqnl$+`?J(?q1Mzr()Bd+!3J^81wKmlQg79E z2%WkVG(f2`h(9=lRv456JcwTJKv}?pEY~|=9$;71>m4u`xVv%->8;vfJqi$@lo|j( zx+iy(NDB(XfxC-(RDL$CyB+l1pTh;1q{hKZkHUpe z_3tfqQRNTDUsOLBy;1z&2&bl(#SO-vGFU#>792zRDm)9U+G-y38fw%$2*zJjCK$aj z4}v39^B@?%IS+zkH0D8Y1kM7hwweds_1J%wpW6i@$f8Dw5RGYp2&44q_3lVH05KM7 z03tNR07R+!6d(rAaDgh!i<_RE5HBSIBn)3t@r!h|e zQ8*jC1Jpd}E?AC3uapWr$f8cDJT4GMm8#c)Bjo_ZSdFwOt5vKiqcF5lW>n8WW+P~s8QEj1+4R$0Ge6oW zvym)IHP15}K|3)cpvh>HJ&%R)3*2Q#+bLz`yW@g`{QnRQ|^)JU#0q7|Lc}~x+jlKhTT_o z-8FZkUFbPQ5O`AK1f96cHlGHeyraQSys5&`$Wm?qc2Vuw8`RIHKe5FQ#?G}41};k- z3|&-ubw1LDvhnFKXg(Uj3*M{7HbP6a6L2(m_DD3c&?bjOV;9w4^Pi{K!PvRh!N6sy zgQ1H`ug*ukPV<>m(0nw47ra-^XM~n&C*WxC?2%|>p-m2n#xAP8=08udgRyh1gMrIZ zJE0#jd({$F*d|x3PV$vaa0y(vV%2a4Sj9IIIQtSr8dk6sFatBnj{d|DTg5W)ntNK1ly9+7Yx@_>XZ!1Kq$s>Z7`aOmKyf?b_`cGwZw?5HD>*-;1NvV#uIwv~z8alxuw xcGwZQ?5HDh*-;1NvV#uEwQbU_%4LTgk;{%cB9|R?KrRj1*n*xr8u-62{}0d79&`Wz literal 98224 zcmcJY-EQ1Q5{2*c6h`)Tkw8By=w`3?Ar^~)BP*5=TQ)Rul)U{`lT)m&{vo@%PA!5& z>x3T<>zuCcX;PejeE84f%j4_Ex5rN(e*62wr=eJ+J|NYZX*}q=?@Wa19{P(91^uM3^-<+SnzMamWAHRM*-T!rSO@F&S zy}X=0etG`8!3DK^eII*PSc7b_CVNdr~_h)y8Y;1cNTrn_oL39pQ|U9UIt{fWNTs*9c0`ittganF8K4HePTTV_G9sup#v@ueG9prKyrxbtNToNw zc0@YD86`Vn8>esHPuqh4nREvas8VDm&;ggaI!V^vr)fm(v| zD$Rs!&Yi8S9YKZRCNz6eJAz8%^wUg8oziKj9YHF+#kC`nOlNiN2+9C8=w;e=#K?%C z+8B?h`hp#iYU4F^f{~^_XMXL7bc8cXc7%-6$@QNVkWIRiXEYP&0Iuu&DQyHAp`4$| zQ<@2MfK6jHJk10e&(5nf6Y@ECva)sr6=rKflP9$!s5DML&4kn`orc;Gq|#elJ0i(+ zR@aW83{Zn!r|tO|84*+);}KO~up?4!yrxbt()8!duN{$&a7O6>L3g9+hB)2PUT$uG z`<`x{KYx2UH*ROYJ-(d3y?*C3KocVw zP2#Y^yigVfuUHc1fp)8Y8)2mZ$`z}uG%%VOYg%c5vaqV^21Zwy2imprI>Jf=lq*(Q zX<&3K=7F-Xs_F<~9%yalWrD{xSUZ#}CRMjH2CTG0S(sGa&gcsBK)Y6Y+XY#Wb^goM z8lYUU%1Q&HTQLumg;iBY2=hR@R(jh-(@F!BD^^wAz!PTT8XxB<_yHKN&uhszNid9uNFvhJkKv`H-bpxX-%meLO z>1`KHD-BStSY@Sw(XE&V%EGFuBZPUNwUwatn6}am<%&tw?Ti5{?NAmbRkt&`!aUHf zmELyIw9)|Oid9w`7~P6_pe(GaIzpHS+O^W#E}B*vpj@%4>ITNRl?EsatEz5bbcNZV zPuu>hsy6>R((cMtj;+LaU~~b-z1AHe9#}edtERUi9$2%QyTg8`A|6fIe2U(!~<(u;}KF> zGYC`G@ZJ$*g=U1QY8sE-x5fxl))v6T=HtZ9u$NM+3+Oj~1HLR1y;z;rd8LOd{S4M2!kV}muTIXL3R ztQmwgt2uaVCBy@3TH_H?Su+T0S~ED}s)~4E&1%jH@xYqac!X5ewHZ6KZH8CIe{Q52)_UrS@JZhS%R#otE zL|2xBS(!MrC6O;6PapX4gb#gHR)pFLj*B)v5|UjuL-~d?;W^E#7y`@bwGr3w)a5G}S zL^G3RNx{vCt}F+$YofD3G))|5aaYR>^jrqh|EvnGGHMYAsxgA%*-GOT~WsKD<lhlu?;7#?q#Y ztZPP1F|o8MrbC<6pv$^M)Wp)Jm=0~q5K~O7rYU2CtdikiHBFh6;b1jQamc7l8DljI zIX1|~Oc`S}3psXZCBwmLn&OaAnKH&wrtsF!xMmOvzc*BZjo~wf}_S_ML)%o__oI_VD!h>2!X3ef~uMyt-@;o!%~E^7gM*01gRs zXIBWhT7EF$!^%3G|A|#w+X&ui+KRK8rDmo*ND=z>5hDU*BVi&(TERq)w1SBs z5kO+F7o_{s^QX_RA3xDQ=mDd&)0k%DJ0c*%wRx%`Tx?D=vENv4;=G|Yab?+Dy zvTDlQ^ZotRT^_nRpI&r<(1zPw-l?IDcwED)yH|S0dSi*?Num{cSVt>xq>>dg z8~A+jX+)2Fu%f~L`}g$DdMWriQf1_Ceg1)l{vWNezkasZP&C^iaj%bGex#XkX~^r; zxOJr|ZBXFT`ThC)1l0#dinB}EV8(iBxeZ=ciX4)sEkLNl5~c4T1_X}A9#F*~Z3M{^ z2TZ^_8DL%5tUpXOIUJ2WXsqC-lp%TI07@N}xVo_Sna)0`B53RZf|d7P1Wz15 z(2_bVar1PY7D(j!!_;7a!7@h?x_~%UeS|hBa5Z3&g=)G07?&S_%7-TaC&dTA^xy;l zS2uU}aqJBjf@4P^u-sD!P+W!Jw6754<}9V8v8^`cst=Vp>H`&S`e4dQAF6WEhufyQ z%O`{0#7Wz?H`h<^H~kwQjh%+$rGf#=&O)KN%TTa54#jHsq3A6M6e#+jL+tg3Vg+Sz zEQ!Tavm8K`MsQUMplXriAGgJR26XsN$~xfwWlw;`i*VgWz$i3T+@sr>?xvL zH;y?pi`lQ3{O-<)gEV!u)4Vg-U*Vzr6{5Q0G$&1`P?&yyg_T_~zz=n`)4boQUF#bL zGRWkoa=1M_z4r!i29zE?G~ULbSrLyd9z*66IZ&xG2&pv$JQ(!xib8>z>-$k3*9S%_ zQ&1YGq3H%QA(~r~8P_<$VXejJP;_Qoh+s^FWC9uoEQpnekEj15Y8zxjaoBl^i%ePYxW;spsUt zPoR?n$FqJ(4*V37oMpnK><9LobP5xTy@S2ny)j1 zrn7p<4EhX`nM)e;2RqV}%Am}k>3L?*Y_2+I2F=%*LDN~iWCnc(of(>v0yBf8mzhDb zx$1dlko-C`NII+6nL(am-98!XUYG5;#+Ka9=kL%Hu&aUkgHIvO23|w>Nmi37m#L-Y zcjg*1l_Nz#Ct!<%@XH{+Jc}N53Nbk71nC;WPjcDHT_RM@UFkP;*578EX+LW$5lqes zj+n;-M-{VhOdSFo_h{WR>wN_Gc&*}5i%-|c`oyC+B4@H>X#S$!d>cqpECk6u!GL)a zfI-C@7*e+Y7*<9q47pTvyY~ab=Hnj*Wv%d7>Cq3@8xXMLHNgGT`_S*eu6GwxY0*r! zX>SI!tHe{q9<0&`u5k9k#*k&$qV>j*^PUi)I3jDhGeoG4$;RPCZXF!*Az`OITA}P15ScdzAXZ!g zLX*}Iu#8`@+7NKw7XtJKgiLpZ0L`dW(65Z7w-uGIA?nST=}maM_YM?KTEbpmv@zK% z95HVQII0+gW9oK*dLD zQ_(ub9<0&`t_)SI)(CFi1Mq6`>C#;QuQ(zP?g?8D1!)70Ca96@6AYL)0T@)gfgzL5 z0IZBusL}vzKK|ia;jz-AAFdaYgyl^D$VET_l;Q8`eg%}B0wVJk0K|$%K&Wm5K(M>@ z7e0ykj_?;u_8v#ft;bQtc^p$4kK;l*b>Y)e)omLHt2HPCq!v$xO#+?zLUNW~e`A>h z4j;`>JZS@ay~@U9qj1E$4dAF^500tZ0FEnd0LRpA0LM*QgEA08)@=Y+2d-~K&JLFY zwjrCOf{~nno$1A{5e^k^pgIGv(gC1K1F-q{hiio=K#zX7UPu!39?LpVJ$*{^Wbosd zY!i;Sbkl*|;G?FZuU+iHDvjXEP{nGE;O0F5uNI#!-39Q9Bl6&$u=P-oHsEN28p%Gv zfO!*uLB$&wGU*J!%1DJO4Z!B(AFdT1D?R$*dLcW-3EYQck3^F67wD5FPQ8-j+k4Iql)u5rZyhOg>>q|r=_YRyTL~*>w*iK0tzat z#it8433L=ga%SBP9(cBYVn5hQWwZ7=pp}_710+{$LQ-`@K$4{;AgMgej+Yr$-}1H) zp*SLIx-~?oj!E6&GQnzIiUl0Xj_pj7GdcsX;u2VE2v|A;SZxS6?+XEX1JaWEsF{y)opxH$*6o$eQjB5vpUd z0hr$P;E?v_hdy4N7i8EpQE*vtMAm>4MMw2q>as7!<9Fkyr;|s&AaB`hzPZ+4kav4ZE$Z@J zSeI|eqd!@`LFSHryOX|5=Hc#^z9R2#;5`=%XLC}+3ORwKDg@8^*KmcLBF|XJ3HTZ~ zY9S{{S0N^+B&F<7ze_C-pCa%H$`pa|>_DC(@G0sPflt8Kpi=}siB8caI$0JIiLfA9Z3BvOE3Uq+5=cw|Ne=%bEvJl`zM#^FIMs{Fqu6vOj7{8et7@gHusd1cP z-NfQ0$|W=Oho}{b=JOTM>=<^AEk9vr;Hy|t6n2)%3-|mo;05>`4KKiF=y(A>OXY>P zRS#H!Jw?L_>tB%Qr^oQuX)Q=nDHq)z%nZ)WrL2NX5;48<(_1B$DiLQ(7ffFkSufTHqf z!zZll4`{^^S*!a4S~V(l`cVpUPoBu>E_jY^Qfp>PMg({9Ektb$S=FF7T{ugN!4}t# zfo|~={S1xG8X6kCbu=`XHHR3e^4n`?&irDq#pPq5jr?N3Remv8&O(4~jr?LY zhZv{|(Vg3l_YrB((ib=s8znN5VdSC+Fm5BrFt)03uf4lVi~<)GMqwKTMxm=-fdV(X z1q$Bi7brCA4zg0^x!2>|c}9VYBA~F1JfqN6o>5@VMTmioJfn37S*gm=?fS=iK(lk8 zm0XkosyK?RYTWB{)y3=O>WW-N5cF&4RH z7>m_A(ZHI9fW`e@`}LTQVdmarh^6%ysx}D2OiClTmo9sWK2B={ck8VoYVqkxdLukr zQX9mUKG2(2JbeJe%$>pz%RYdiYVR@3x({Hudc0zooKtkB%07Tqi%(bG2e67GaweX2 zS1!lnBFsx<_xy@KiuLQ4{Qs{^ZJsW0HExm`%55nOE*Qf&-b&kjhvG30WF zKr4>OL_JHORmWsQMsjaf^YsT5HTMj~99=hBE1;c1s*NG*QH<0ZLoWLRT5&|y>i&RM z9g_{|58c(y@z(IvO45qWU4v>30c*FQdV|oCgx<_%b!==UX{F|NA(cjOwN*&9F=XA( zk$PjuX>X$e~<@Ix(lAE`x9{P;v1a0bqP3G)wVZXI7^Gc7L~<78->MyFI`9c z!eAG@%5QgW zv+g4%$uZ=j2vEgQY*mR~vvijj1um+9!Ztbu3SIRI6qrX|h=GlMfz};lr7B0S$GP*2 z0vAO3QR87MH1y$}N%Jp-*f$Vyd?UY~R483itifWkKNj6zp=MuD4oM!_3-Mxj~v Rtd-3#mc6;IzF78u{|DuKKpg-8 diff --git a/dnn/src/CMakeLists.txt b/dnn/src/CMakeLists.txt index d9f28286d..962ce8740 100644 --- a/dnn/src/CMakeLists.txt +++ b/dnn/src/CMakeLists.txt @@ -183,6 +183,8 @@ if(MGE_WITH_CUDA) gen_cutlass_kimpl(conv2d tensorop8832 CUTLASS_SOURCES) gen_cutlass_kimpl(dwconv2d_fprop simt CUTLASS_SOURCES) gen_cutlass_kimpl(dwconv2d_fprop tensorop884 CUTLASS_SOURCES) + gen_cutlass_kimpl(dwconv2d_dgrad simt CUTLASS_SOURCES) + gen_cutlass_kimpl(dwconv2d_dgrad tensorop884 CUTLASS_SOURCES) list(APPEND SOURCES ${CUTLASS_SOURCES}) list(APPEND SOURCES ${CUSOURCES}) endif() diff --git a/dnn/src/cuda/conv_bias/algo.cpp b/dnn/src/cuda/conv_bias/algo.cpp index 086746ad5..bc5bcb294 100644 --- a/dnn/src/cuda/conv_bias/algo.cpp +++ b/dnn/src/cuda/conv_bias/algo.cpp @@ -304,12 +304,13 @@ void ConvBiasForwardImpl::AlgoPack::fill_imma_algos() { void ConvBiasForwardImpl::AlgoPack::fill_dwconv_algos() { using AlgoParam = AlgoCutlassConvolutionBase::AlgoParam; + /// preferred algo + f32_implicit_bmm.emplace_back(AlgoParam{64, 128, 8, 32, 64, 8, 1, 1, 1, 2}); f32_implicit_bmm.emplace_back(AlgoParam{128, 128, 8, 32, 64, 8, 1, 1, 1, 2}); f32_implicit_bmm.emplace_back(AlgoParam{128, 64, 8, 64, 32, 8, 1, 1, 1, 2}); f32_implicit_bmm.emplace_back(AlgoParam{128, 32, 8, 64, 32, 8, 1, 1, 1, 2}); f32_implicit_bmm.emplace_back(AlgoParam{32, 128, 8, 32, 64, 8, 1, 1, 1, 2}); - f32_implicit_bmm.emplace_back(AlgoParam{64, 128, 8, 64, 32, 8, 1, 1, 1, 2}); - f32_implicit_bmm.emplace_back(AlgoParam{64, 64, 8, 64, 32, 8, 1, 1, 1, 2}); + f32_implicit_bmm.emplace_back(AlgoParam{64, 64, 8, 32, 64, 8, 1, 1, 1, 2}); f32_implicit_bmm.emplace_back(AlgoParam{32, 64, 8, 32, 64, 8, 1, 1, 1, 2}); f32_implicit_bmm.emplace_back(AlgoParam{32, 32, 8, 32, 32, 8, 1, 1, 1, 2}); f32_implicit_bmm.emplace_back(AlgoParam{64, 32, 8, 64, 32, 8, 1, 1, 1, 2}); @@ -317,10 +318,11 @@ void ConvBiasForwardImpl::AlgoPack::fill_dwconv_algos() { all_algos.push_back(&algo); } #if CUDA_VERSION >= 10020 + /// preferred algo + f16_implicit_bmm.emplace_back(AlgoParam{64, 128, 32, 32, 32, 32, 8, 8, 4, 2}); f16_implicit_bmm.emplace_back(AlgoParam{128, 128, 32, 32, 32, 32, 8, 8, 4, 2}); f16_implicit_bmm.emplace_back(AlgoParam{128, 256, 32, 64, 64, 32, 8, 8, 4, 2}); f16_implicit_bmm.emplace_back(AlgoParam{128, 64, 32, 32, 32, 32, 8, 8, 4, 2}); - f16_implicit_bmm.emplace_back(AlgoParam{64, 128, 32, 32, 32, 32, 8, 8, 4, 2}); f16_implicit_bmm.emplace_back(AlgoParam{64, 64, 32, 32, 32, 32, 8, 8, 4, 2}); for (auto&& algo : f16_implicit_bmm) { all_algos.push_back(&algo); diff --git a/dnn/src/cuda/conv_bias/cutlass_convolution_base.cpp b/dnn/src/cuda/conv_bias/cutlass_convolution_base.cpp index ab5785467..fb89ec6e6 100644 --- a/dnn/src/cuda/conv_bias/cutlass_convolution_base.cpp +++ b/dnn/src/cuda/conv_bias/cutlass_convolution_base.cpp @@ -272,8 +272,10 @@ std::pair get_tensor_alignment( alignment_src /= src.dtype.size(1); }; + /// TODO: need a better way to check whether tensor core instruction is used if (format == Format::NCHW32 || format == Format::NCHW32_NCHW4 || - format == Format::NCHW64 || format == Format::NCHW64) { + format == Format::NCHW64 || format == Format::NCHW64 || + format == Format::NHWC) { get_tensor_alignment_tensor_op(); } else if ( format == Format::NCHW4 || format == Format::NCHW4_NCHW || diff --git a/dnn/src/cuda/conv_bias/implicit_batched_gemm_float16_nchw_hmma.cpp b/dnn/src/cuda/conv_bias/implicit_batched_gemm_float16_nchw_hmma.cpp index 74843e9da..6ef571448 100644 --- a/dnn/src/cuda/conv_bias/implicit_batched_gemm_float16_nchw_hmma.cpp +++ b/dnn/src/cuda/conv_bias/implicit_batched_gemm_float16_nchw_hmma.cpp @@ -23,6 +23,7 @@ bool ConvBiasForwardImpl::AlgoFloat16NCHWHMMAImplicitBatchedGemm::is_available( #define RETURN_IF_FALSE(stmt_) \ if (!(stmt_)) \ return false; + RETURN_IF_FALSE(is_compute_capability_required(7, 0)); RETURN_IF_FALSE( args.src_layout->is_contiguous() && args.dst_layout->is_contiguous()); using Param = param::ConvBias; diff --git a/dnn/src/cuda/convolution/backward_data/algo.cpp b/dnn/src/cuda/convolution/backward_data/algo.cpp index 1196ceb22..acd41aa69 100644 --- a/dnn/src/cuda/convolution/backward_data/algo.cpp +++ b/dnn/src/cuda/convolution/backward_data/algo.cpp @@ -41,6 +41,7 @@ ConvolutionBackwardDataImpl::AlgoPack::AlgoPack() { all_algos.push_back(&algo); int8_algos.push_back(&algo); } + fill_dwconv_algos(); int8_algos.push_back(&int8_nchw_dotprod); all_algos.push_back(&int8_nchw_dotprod); @@ -54,6 +55,39 @@ ConvolutionBackwardDataImpl::AlgoPack::AlgoPack() { } } +void ConvolutionBackwardDataImpl::AlgoPack::fill_dwconv_algos() { + { + using AlgoParam = AlgoFloat32NCHWFMAImplicitBatchedGemm::AlgoParam; + /// preferred algo + implbmm_nchw_fma.emplace_back(AlgoParam{64, 128, 8, 32, 64, 8, 2}); + implbmm_nchw_fma.emplace_back(AlgoParam{128, 128, 8, 32, 64, 8, 2}); + implbmm_nchw_fma.emplace_back(AlgoParam{128, 64, 8, 64, 32, 8, 2}); + implbmm_nchw_fma.emplace_back(AlgoParam{128, 32, 8, 64, 32, 8, 2}); + implbmm_nchw_fma.emplace_back(AlgoParam{32, 128, 8, 32, 64, 8, 2}); + implbmm_nchw_fma.emplace_back(AlgoParam{64, 64, 8, 32, 64, 8, 2}); + implbmm_nchw_fma.emplace_back(AlgoParam{32, 64, 8, 32, 64, 8, 2}); + implbmm_nchw_fma.emplace_back(AlgoParam{32, 32, 8, 32, 32, 8, 2}); + implbmm_nchw_fma.emplace_back(AlgoParam{64, 32, 8, 64, 32, 8, 2}); + for (auto&& algo : implbmm_nchw_fma) { + all_algos.push_back(&algo); + } + } +#if CUDA_VERSION >= 10020 + { + using AlgoParam = AlgoFloat16NCHWHMMAImplicitBatchedGemm::AlgoParam; + /// preferred algo + implbmm_nchw_hmma.emplace_back(AlgoParam{64, 128, 32, 32, 32, 32, 8, 8, 4, 2}); + implbmm_nchw_hmma.emplace_back(AlgoParam{128, 128, 32, 32, 32, 32, 8, 8, 4, 2}); + implbmm_nchw_hmma.emplace_back(AlgoParam{128, 256, 32, 64, 64, 32, 8, 8, 4, 2}); + implbmm_nchw_hmma.emplace_back(AlgoParam{128, 64, 32, 32, 32, 32, 8, 8, 4, 2}); + implbmm_nchw_hmma.emplace_back(AlgoParam{64, 64, 32, 32, 32, 32, 8, 8, 4, 2}); + for (auto&& algo : implbmm_nchw_hmma) { + all_algos.push_back(&algo); + } + } +#endif +} + MEGDNN_DEF_GET_ALGO_FROM_DESC(ConvolutionBackwardDataImpl) ConvolutionBackwardDataImpl::AlgoCUDNN* ConvolutionBackwardDataImpl::AlgoPack:: diff --git a/dnn/src/cuda/convolution/backward_data/algo.h b/dnn/src/cuda/convolution/backward_data/algo.h index d46a5652b..f0098a0d9 100644 --- a/dnn/src/cuda/convolution/backward_data/algo.h +++ b/dnn/src/cuda/convolution/backward_data/algo.h @@ -41,7 +41,9 @@ public: CUDA_GROUP_CONV_GENERAL, CUDA_IMPLICIT_GEMM_NCHW4_DOTPROD_INT8, CUDA_IMPLICIT_GEMM_NCHW_DOTPROD_INT8, - CUDA_IMPLICIT_GEMM_NHWC_IMMA_INT8 + CUDA_IMPLICIT_GEMM_NHWC_IMMA_INT8, + CUDA_IMPLICIT_BATCHED_GEMM_FMA_NCHW_F32, + CUDA_IMPLICIT_BATCHED_GEMM_HMMA_NCHW_F16, }; using Mapper = std::unordered_map; @@ -315,6 +317,82 @@ private: std::string m_name; }; +class ConvolutionBackwardDataImpl::AlgoFloat32NCHWFMAImplicitBatchedGemm final + : public AlgoBase { +public: + struct AlgoParam { + int threadblock_m; + int threadblock_n; + int threadblock_k; + int warp_m; + int warp_n; + int warp_k; + int stage; + std::string to_string() { + return ssprintf( + "_%dX%dX%d_%dX%dX%d_%dstage", threadblock_m, threadblock_n, + threadblock_k, warp_m, warp_n, warp_k, stage); + } + }; + AlgoFloat32NCHWFMAImplicitBatchedGemm(AlgoParam algo_param) + : m_algo_param{algo_param}, + m_name{ssprintf( + "FLOAT32_NCHW_FMA_IMPLICIT_BATCHED_GEMM%s", + m_algo_param.to_string().c_str())} {} + bool is_available(const SizeArgs& args) const override; + size_t get_workspace_in_bytes(const SizeArgs& args) const override { return 0; } + void exec(const ExecArgs& args) const override; + const char* name() const override { return m_name.c_str(); } + AlgoAttribute attribute() const override { return AlgoAttribute::REPRODUCIBLE; } + MEGDNN_DECL_ALGO_TYPE(CUDA_IMPLICIT_BATCHED_GEMM_FMA_NCHW_F32) + +private: + const void* get_available_op(const SizeArgs& args) const; + AlgoParam m_algo_param; + std::string m_name; +}; + +class ConvolutionBackwardDataImpl::AlgoFloat16NCHWHMMAImplicitBatchedGemm final + : public AlgoBase { +public: + /// add instruction shape as member of algo param, because f16 tensor core has 2 + /// different matrix shapes (i.e. mma.884 and mma.1688) + struct AlgoParam { + int threadblock_m; + int threadblock_n; + int threadblock_k; + int warp_m; + int warp_n; + int warp_k; + int instruction_m; + int instruction_n; + int instruction_k; + int stage; + std::string to_string() { + return ssprintf( + "_%dX%dX%d_%dX%dX%d_mma%dX%dX%d_%dstage", threadblock_m, + threadblock_n, threadblock_k, warp_m, warp_n, warp_k, instruction_m, + instruction_n, instruction_k, stage); + } + }; + AlgoFloat16NCHWHMMAImplicitBatchedGemm(AlgoParam algo_param) + : m_algo_param{algo_param}, + m_name{ssprintf( + "FLOAT16_NCHW_HMMA_IMPLICIT_BATCHED_GEMM%s", + m_algo_param.to_string().c_str())} {} + bool is_available(const SizeArgs& args) const override; + size_t get_workspace_in_bytes(const SizeArgs& args) const override { return 0; } + void exec(const ExecArgs& args) const override; + const char* name() const override { return m_name.c_str(); } + AlgoAttribute attribute() const override { return AlgoAttribute::REPRODUCIBLE; } + MEGDNN_DECL_ALGO_TYPE(CUDA_IMPLICIT_BATCHED_GEMM_HMMA_NCHW_F16) + +private: + const void* get_available_op(const SizeArgs& args) const; + AlgoParam m_algo_param; + std::string m_name; +}; + class ConvolutionBackwardDataImpl::AlgoPack : NonCopyableObj { // defined in cudnn.cpp void fill_cudnn_algos(); @@ -322,6 +400,7 @@ class ConvolutionBackwardDataImpl::AlgoPack : NonCopyableObj { void fill_int8_dp4a_algos(); // defined in implicit_gemm_int8_nhwc_imma.cpp void fill_int8_imma_algos(); + void fill_dwconv_algos(); AlgoBase::Mapper m_all_algos_map; @@ -337,6 +416,8 @@ public: std::vector int8_nchw4_dotprod; AlgoInt8NCHWDotProdImplicitGemm int8_nchw_dotprod; std::vector int8_nhwc_imma; + std::vector implbmm_nchw_fma; + std::vector implbmm_nchw_hmma; std::vector //! all algorithms diff --git a/dnn/src/cuda/convolution/backward_data/implicit_batched_gemm_float16_nchw_hmma.cpp b/dnn/src/cuda/convolution/backward_data/implicit_batched_gemm_float16_nchw_hmma.cpp new file mode 100644 index 000000000..f1ec797d7 --- /dev/null +++ b/dnn/src/cuda/convolution/backward_data/implicit_batched_gemm_float16_nchw_hmma.cpp @@ -0,0 +1,146 @@ +/** + * \file + * dnn/src/cuda/convolution/backward_data/implicit_batched_gemm_float16_nchw_hmma.cpp + * MegEngine is Licensed under the Apache License, Version 2.0 (the "License") + * + * Copyright (c) 2014-2021 Megvii Inc. All rights reserved. + * + * Unless required by applicable law or agreed to in writing, + * software distributed under the License is distributed on an + * "AS IS" BASIS, WITHOUT ARRANTIES OR CONDITIONS OF ANY KIND, either express or + * implied. + */ + +#include "src/cuda/convolution/backward_data/algo.h" +#include "src/cuda/cutlass/singleton.h" +#include "src/cuda/utils.h" + +using namespace megdnn; +using namespace cuda; +using namespace cutlass::library; + +const void* ConvolutionBackwardDataImpl::AlgoFloat16NCHWHMMAImplicitBatchedGemm:: + get_available_op(const SizeArgs& args) const { + int alignment_diff = 0; + int wo = args.diff_layout->dtype.size(args.diff_layout->operator[](3)); + for (int candidate : {16, 4, 2}) { + if (wo % candidate == 0) + alignment_diff = candidate; + } + alignment_diff /= args.diff_layout->dtype.size(1); + NumericTypeID accumulator_dtype = + args.opr->param().compute_mode == param::Convolution::ComputeMode::DEFAULT + ? NumericTypeID::kF16 + : NumericTypeID::kF32; + ConvolutionKey key{ + cutlass::conv::Operator::kDgrad, + NumericTypeID::kF16, + LayoutTypeID::kTensorNCHW, + NumericTypeID::kF16, + LayoutTypeID::kTensorNCHW, + NumericTypeID::kF16, + LayoutTypeID::kTensorNCHW, + NumericTypeID::kF16, + LayoutTypeID::kTensorNCHW, + accumulator_dtype, + cutlass::conv::ConvType::kDepthwiseConvolution, + m_algo_param.threadblock_m, + m_algo_param.threadblock_n, + m_algo_param.threadblock_k, + m_algo_param.warp_m, + m_algo_param.warp_n, + m_algo_param.warp_k, + m_algo_param.instruction_m, + m_algo_param.instruction_n, + m_algo_param.instruction_k, + cutlass::epilogue::EpilogueType::kBiasAddLinearCombination, + m_algo_param.stage, + cutlass::conv::SpecialOptimizeDesc::NONE, + alignment_diff, + 1, + false}; + return (void*)Singleton::get().operation_table.find_op(key); +} + +bool ConvolutionBackwardDataImpl::AlgoFloat16NCHWHMMAImplicitBatchedGemm::is_available( + const SizeArgs& args) const { +#define RETURN_IF_FALSE(stmt_) \ + if (!(stmt_)) \ + return false; + RETURN_IF_FALSE(is_compute_capability_required(7, 0)); + RETURN_IF_FALSE( + args.diff_layout->is_contiguous() && args.grad_layout->is_contiguous()); + using Param = param::Convolution; + using Format = Param::Format; + using Sparse = Param::Sparse; + using Mode = Param::Mode; + auto&& param = args.opr->param(); + auto&& fm = args.filter_meta; + RETURN_IF_FALSE( + param.format == Format::NCHW && + args.diff_layout->dtype.enumv() == DTypeEnum::Float16 && + args.filter_layout->dtype.enumv() == DTypeEnum::Float16 && + args.grad_layout->dtype.enumv() == DTypeEnum::Float16); + RETURN_IF_FALSE(param.sparse == Sparse::GROUP); + RETURN_IF_FALSE(param.mode == Mode::CROSS_CORRELATION); + // check if channelwise convolution + RETURN_IF_FALSE(fm.icpg == 1 && fm.ocpg == 1); + const auto* op = get_available_op(args); + RETURN_IF_FALSE(op != nullptr); + return true; +#undef RETURN_IF_FALSE +} + +void ConvolutionBackwardDataImpl::AlgoFloat16NCHWHMMAImplicitBatchedGemm::exec( + const ExecArgs& args) const { + auto&& param = args.opr->param(); + auto&& fm = args.filter_meta; + int n = args.diff_layout->operator[](0), ho = args.diff_layout->operator[](2), + wo = args.diff_layout->operator[](3); + int hi = args.grad_layout->operator[](2), wi = args.grad_layout->operator[](3); + int co = fm.group, ci = co, groups = co; + int fh = fm.spatial[0], fw = fm.spatial[1]; + int sh = fm.stride[0], sw = fm.stride[1]; + int ph = fm.padding[0], pw = fm.padding[1]; + int dh = param.dilate_h, dw = param.dilate_w; + + // check if channelwise convolution + megdnn_assert(fm.icpg == 1 && fm.ocpg == 1); + auto&& stream = cuda_stream(args.opr->handle()); + + float alpha = 1.f; + float beta = 0.f; + float gamma = 0.f; + float delta = 0.f; + + const Operation* op = (const Operation*)get_available_op(args); + + cutlass::conv::Conv2dProblemSize problem_size{ + n, hi, wi, ci, co, fh, fw, ho, + wo, ph, pw, sh, sw, dh, dw, cutlass::conv::Mode::kCrossCorrelation, + 1, // split k slices, always 1 + groups, // groups + }; + + cutlass::library::ConvolutionArguments conv_args{ + problem_size, + args.diff_tensor->raw_ptr(), + args.filter_tensor->raw_ptr(), + nullptr, + nullptr, + args.grad_tensor->raw_ptr(), + &alpha, + &beta, + &gamma, + &delta, + nullptr, + nullptr, + nullptr, + nullptr}; + + cutlass_check(op->run(&conv_args, nullptr, stream)); + + after_kernel_launch(); +} + +// vim: syntax=cpp.doxygen diff --git a/dnn/src/cuda/convolution/backward_data/implicit_batched_gemm_float32_nchw_fma.cpp b/dnn/src/cuda/convolution/backward_data/implicit_batched_gemm_float32_nchw_fma.cpp new file mode 100644 index 000000000..7487f1708 --- /dev/null +++ b/dnn/src/cuda/convolution/backward_data/implicit_batched_gemm_float32_nchw_fma.cpp @@ -0,0 +1,141 @@ +/** + * \file + * dnn/src/cuda/convolution/backward_data/implicit_batched_gemm_float32_nchw_fma.cpp + * MegEngine is Licensed under the Apache License, Version 2.0 (the "License") + * + * Copyright (c) 2014-2021 Megvii Inc. All rights reserved. + * + * Unless required by applicable law or agreed to in writing, + * software distributed under the License is distributed on an + * "AS IS" BASIS, WITHOUT ARRANTIES OR CONDITIONS OF ANY KIND, either express or + * implied. + */ + +#include "src/cuda/convolution/backward_data/algo.h" +#include "src/cuda/cutlass/singleton.h" +#include "src/cuda/utils.h" + +using namespace megdnn; +using namespace cuda; +using namespace cutlass::library; + +const void* ConvolutionBackwardDataImpl::AlgoFloat32NCHWFMAImplicitBatchedGemm:: + get_available_op(const SizeArgs& args) const { + int alignment_diff = 0; + int wo = args.diff_layout->dtype.size(args.diff_layout->operator[](3)); + for (int candidate : {16, 4}) { + if (wo % candidate == 0) + alignment_diff = candidate; + } + alignment_diff /= args.diff_layout->dtype.size(1); + ConvolutionKey key{ + cutlass::conv::Operator::kDgrad, + NumericTypeID::kF32, + LayoutTypeID::kTensorNCHW, + NumericTypeID::kF32, + LayoutTypeID::kTensorNCHW, + NumericTypeID::kF32, + LayoutTypeID::kTensorNCHW, + NumericTypeID::kF32, + LayoutTypeID::kTensorNCHW, + NumericTypeID::kF32, + cutlass::conv::ConvType::kDepthwiseConvolution, + m_algo_param.threadblock_m, + m_algo_param.threadblock_n, + m_algo_param.threadblock_k, + m_algo_param.warp_m, + m_algo_param.warp_n, + m_algo_param.warp_k, + 1, + 1, + 1, + cutlass::epilogue::EpilogueType::kBiasAddLinearCombination, + m_algo_param.stage, + cutlass::conv::SpecialOptimizeDesc::NONE, + alignment_diff, + 1, + false}; + return (void*)Singleton::get().operation_table.find_op(key); +} + +bool ConvolutionBackwardDataImpl::AlgoFloat32NCHWFMAImplicitBatchedGemm::is_available( + const SizeArgs& args) const { +#define RETURN_IF_FALSE(stmt_) \ + if (!(stmt_)) \ + return false; + RETURN_IF_FALSE( + args.diff_layout->is_contiguous() && args.grad_layout->is_contiguous()); + using Param = param::Convolution; + using Format = Param::Format; + using Sparse = Param::Sparse; + using Mode = Param::Mode; + auto&& param = args.opr->param(); + auto&& fm = args.filter_meta; + RETURN_IF_FALSE( + param.format == Format::NCHW && + args.diff_layout->dtype.enumv() == DTypeEnum::Float32 && + args.filter_layout->dtype.enumv() == DTypeEnum::Float32 && + args.grad_layout->dtype.enumv() == DTypeEnum::Float32); + RETURN_IF_FALSE(param.sparse == Sparse::GROUP); + RETURN_IF_FALSE(param.mode == Mode::CROSS_CORRELATION); + // check if channelwise convolution + RETURN_IF_FALSE(fm.icpg == 1 && fm.ocpg == 1); + const auto* op = get_available_op(args); + RETURN_IF_FALSE(op != nullptr); + return true; +#undef RETURN_IF_FALSE +} + +void ConvolutionBackwardDataImpl::AlgoFloat32NCHWFMAImplicitBatchedGemm::exec( + const ExecArgs& args) const { + auto&& param = args.opr->param(); + auto&& fm = args.filter_meta; + int n = args.diff_layout->operator[](0), ho = args.diff_layout->operator[](2), + wo = args.diff_layout->operator[](3); + int hi = args.grad_layout->operator[](2), wi = args.grad_layout->operator[](3); + int co = fm.group, ci = co, groups = co; + int fh = fm.spatial[0], fw = fm.spatial[1]; + int sh = fm.stride[0], sw = fm.stride[1]; + int ph = fm.padding[0], pw = fm.padding[1]; + int dh = param.dilate_h, dw = param.dilate_w; + + // check if channelwise convolution + megdnn_assert(fm.icpg == 1 && fm.ocpg == 1); + auto&& stream = cuda_stream(args.opr->handle()); + + float alpha = 1.f; + float beta = 0.f; + float gamma = 0.f; + float delta = 0.f; + + const Operation* op = (const Operation*)get_available_op(args); + + cutlass::conv::Conv2dProblemSize problem_size{ + n, hi, wi, ci, co, fh, fw, ho, + wo, ph, pw, sh, sw, dh, dw, cutlass::conv::Mode::kCrossCorrelation, + 1, // split k slices, always 1 + groups, // groups + }; + + cutlass::library::ConvolutionArguments conv_args{ + problem_size, + args.diff_tensor->raw_ptr(), + args.filter_tensor->raw_ptr(), + nullptr, + nullptr, + args.grad_tensor->raw_ptr(), + &alpha, + &beta, + &gamma, + &delta, + nullptr, + nullptr, + nullptr, + nullptr}; + + cutlass_check(op->run(&conv_args, nullptr, stream)); + + after_kernel_launch(); +} + +// vim: syntax=cpp.doxygen diff --git a/dnn/src/cuda/convolution/backward_data/implicit_gemm_int8_nchw4_dp4a.cpp b/dnn/src/cuda/convolution/backward_data/implicit_gemm_int8_nchw4_dp4a.cpp index 1f01859c1..d499556fa 100644 --- a/dnn/src/cuda/convolution/backward_data/implicit_gemm_int8_nchw4_dp4a.cpp +++ b/dnn/src/cuda/convolution/backward_data/implicit_gemm_int8_nchw4_dp4a.cpp @@ -54,7 +54,7 @@ const void* ConvolutionBackwardDataImpl::AlgoInt8NCHW4DotProdImplicitGemm:: m_algo_param.stage, special_optimization, 4, - 16, + 4, false}; return (void*)Singleton::get().operation_table.find_op(key); } diff --git a/dnn/src/cuda/convolution/opr_impl.h b/dnn/src/cuda/convolution/opr_impl.h index 4d22ae36c..0a3c22b74 100644 --- a/dnn/src/cuda/convolution/opr_impl.h +++ b/dnn/src/cuda/convolution/opr_impl.h @@ -102,6 +102,8 @@ public: class AlgoInt8NCHW4DotProdImplicitGemm; class AlgoInt8NCHWDotProdImplicitGemm; class AlgoInt8NHWCIMMAImplicitGemm; + class AlgoFloat32NCHWFMAImplicitBatchedGemm; + class AlgoFloat16NCHWHMMAImplicitBatchedGemm; class AlgoPack; diff --git a/dnn/src/cuda/cutlass/initialize_all.cu b/dnn/src/cuda/cutlass/initialize_all.cu index 3a43f8de6..44d6faf38 100644 --- a/dnn/src/cuda/cutlass/initialize_all.cu +++ b/dnn/src/cuda/cutlass/initialize_all.cu @@ -55,6 +55,7 @@ void initialize_all_gemm_simt_operations(Manifest& manifest); void initialize_all_conv2d_simt_operations(Manifest& manifest); void initialize_all_deconv_simt_operations(Manifest& manifest); void initialize_all_dwconv2d_fprop_simt_operations(Manifest& manifest); +void initialize_all_dwconv2d_dgrad_simt_operations(Manifest& manifest); #if defined(CUTLASS_ARCH_MMA_SM75_SUPPORTED) && CUTLASS_ARCH_MMA_SM75_SUPPORTED void initialize_all_gemm_tensorop884_operations(Manifest& manifest); void initialize_all_gemm_tensorop1688_operations(Manifest& manifest); @@ -62,6 +63,7 @@ void initialize_all_conv2d_tensorop8816_operations(Manifest& manifest); void initialize_all_conv2d_tensorop8832_operations(Manifest& manifest); void initialize_all_deconv_tensorop8816_operations(Manifest& manifest); void initialize_all_dwconv2d_fprop_tensorop884_operations(Manifest& manifest); +void initialize_all_dwconv2d_dgrad_tensorop884_operations(Manifest& manifest); #endif void initialize_all(Manifest& manifest) { @@ -69,6 +71,7 @@ void initialize_all(Manifest& manifest) { initialize_all_conv2d_simt_operations(manifest); initialize_all_deconv_simt_operations(manifest); initialize_all_dwconv2d_fprop_simt_operations(manifest); + initialize_all_dwconv2d_dgrad_simt_operations(manifest); #if defined(CUTLASS_ARCH_MMA_SM75_SUPPORTED) && CUTLASS_ARCH_MMA_SM75_SUPPORTED initialize_all_gemm_tensorop884_operations(manifest); initialize_all_gemm_tensorop1688_operations(manifest); @@ -76,6 +79,7 @@ void initialize_all(Manifest& manifest) { initialize_all_conv2d_tensorop8832_operations(manifest); initialize_all_deconv_tensorop8816_operations(manifest); initialize_all_dwconv2d_fprop_tensorop884_operations(manifest); + initialize_all_dwconv2d_dgrad_tensorop884_operations(manifest); #endif } diff --git a/dnn/test/common/checker.h b/dnn/test/common/checker.h index a3fbee1d5..688de2f0e 100644 --- a/dnn/test/common/checker.h +++ b/dnn/test/common/checker.h @@ -569,6 +569,7 @@ public: }); return ret; } + megdnn_assert(false, "Expected algo not found: %s\n", policy_name.name.c_str()); return ret; } diff --git a/dnn/test/cuda/chanwise_convolution.cpp b/dnn/test/cuda/chanwise_convolution.cpp index 8ad2160e1..2202ffd67 100644 --- a/dnn/test/cuda/chanwise_convolution.cpp +++ b/dnn/test/cuda/chanwise_convolution.cpp @@ -497,15 +497,15 @@ void check_chanwise(DType io_type, DType comp_type, Handle* handle, const char* } } // namespace -#define MEGDNN_FOREACH_CUTLASS_CHANWISE_CONV_FWD_FMA_KERNEL(cb) \ - cb(1, 128, 128, 8, 32, 64, 8); \ - cb(2, 128, 64, 8, 64, 32, 8); \ - cb(3, 128, 32, 8, 64, 32, 8); \ - cb(4, 64, 128, 8, 64, 32, 8); \ - cb(5, 32, 128, 8, 32, 64, 8); \ - cb(6, 64, 64, 8, 64, 32, 8); \ - cb(7, 32, 64, 8, 32, 64, 8); \ - cb(8, 32, 32, 8, 32, 32, 8); \ +#define MEGDNN_FOREACH_CUTLASS_CHANWISE_CONV_FMA_KERNEL(cb) \ + cb(1, 128, 128, 8, 32, 64, 8); \ + cb(2, 128, 64, 8, 64, 32, 8); \ + cb(3, 128, 32, 8, 64, 32, 8); \ + cb(4, 64, 128, 8, 32, 64, 8); \ + cb(5, 32, 128, 8, 32, 64, 8); \ + cb(6, 64, 64, 8, 32, 64, 8); \ + cb(7, 32, 64, 8, 32, 64, 8); \ + cb(8, 32, 32, 8, 32, 32, 8); \ cb(9, 64, 32, 8, 64, 32, 8); #define cb(tag, tbm, tbn, tbk, wm, wn, wk) \ @@ -516,16 +516,29 @@ void check_chanwise(DType io_type, DType comp_type, Handle* handle, const char* "_" #wm "X" #wn "X" #wk "_2stage"); \ } -MEGDNN_FOREACH_CUTLASS_CHANWISE_CONV_FWD_FMA_KERNEL(cb) +MEGDNN_FOREACH_CUTLASS_CHANWISE_CONV_FMA_KERNEL(cb) #undef cb -#undef MEGDNN_FOREACH_CUTLASS_CHANWISE_CONV_FWD_FMA_KERNEL -#define MEGDNN_FOREACH_CUTLASS_CHANWISE_CONV_FWD_HMMA_KERNEL(cb) \ - cb(1, 128, 128, 32, 32, 32, 32); \ - cb(2, 128, 256, 32, 64, 64, 32); \ - cb(3, 128, 64, 32, 32, 32, 32); \ - cb(4, 64, 128, 32, 32, 32, 32); \ +#define cb(tag, tbm, tbn, tbk, wm, wn, wk) \ + TEST_F(CUDA, CHANWISE_CONVOLUTION_BACKWARD_DATA_CUTLASS_FMA_##tag) { \ + check_chanwise( \ + dtype::Float32(), dtype::Float32(), handle_cuda(), \ + "FLOAT32_NCHW_FMA_IMPLICIT_BATCHED_GEMM_" #tbm "X" #tbn "X" #tbk \ + "_" #wm "X" #wn "X" #wk "_2stage"); \ + } + +MEGDNN_FOREACH_CUTLASS_CHANWISE_CONV_FMA_KERNEL(cb) + +#undef cb + +#undef MEGDNN_FOREACH_CUTLASS_CHANWISE_CONV_FMA_KERNEL + +#define MEGDNN_FOREACH_CUTLASS_CHANWISE_CONV_HMMA_KERNEL(cb) \ + cb(1, 128, 128, 32, 32, 32, 32); \ + cb(2, 128, 256, 32, 64, 64, 32); \ + cb(3, 128, 64, 32, 32, 32, 32); \ + cb(4, 64, 128, 32, 32, 32, 32); \ cb(5, 64, 64, 32, 32, 32, 32); // check both ioc16 and io16xc32 @@ -541,9 +554,26 @@ MEGDNN_FOREACH_CUTLASS_CHANWISE_CONV_FWD_FMA_KERNEL(cb) "_" #wm "X" #wn "X" #wk "_2stage"); \ } -MEGDNN_FOREACH_CUTLASS_CHANWISE_CONV_FWD_HMMA_KERNEL(cb) +MEGDNN_FOREACH_CUTLASS_CHANWISE_CONV_HMMA_KERNEL(cb) + +#undef cb + +#define cb(tag, tbm, tbn, tbk, wm, wn, wk) \ + TEST_F(CUDA, CHANWISE_CONVOLUTION_BACKWARD_DATA_CUTLASS_HMMA_##tag) { \ + check_chanwise( \ + dtype::Float16(), dtype::Float16(), handle_cuda(), \ + "FLOAT16_NCHW_HMMA_IMPLICIT_BATCHED_GEMM_" #tbm "X" #tbn "X" #tbk \ + "_" #wm "X" #wn "X" #wk "_mma8X8X4_2stage"); \ + check_chanwise( \ + dtype::Float16(), dtype::Float32(), handle_cuda(), \ + "FLOAT16_NCHW_HMMA_IMPLICIT_BATCHED_GEMM_" #tbm "X" #tbn "X" #tbk \ + "_" #wm "X" #wn "X" #wk "_mma8X8X4_2stage"); \ + } + +MEGDNN_FOREACH_CUTLASS_CHANWISE_CONV_HMMA_KERNEL(cb) #undef cb + #undef MEGDNN_FOREACH_CUTLASS_CHANWISE_CONV_FWD_HMMA_KERNEL #if MEGDNN_WITH_BENCHMARK @@ -1324,6 +1354,81 @@ TEST_F(CUDA, BENCHMARK_CHANWISE_CONV_FORWARD_LARGE_KERNEL) { // clang-format on } +TEST_F(CUDA, BENCHMARK_CHANWISE_CONV_BACKWARD_DATA_LARGE_KERNEL) { + CUBenchmarker bencher(handle_cuda()); + size_t RUNS = 100; + bencher.set_display(false).set_times(RUNS); + std::unique_ptr> proxy{ + new OprProxy{true}}; + bencher.set_proxy(proxy); + + Convolution::Param param; + param.format = ConvBias::Param::Format::NCHW; + param.sparse = Convolution::Param::Sparse::GROUP; + NormalRNG rng; + + auto run = [&](size_t batch, size_t c, size_t ih, size_t iw, size_t f, size_t s) { + param.pad_h = f / 2; + param.pad_w = f / 2; + param.stride_h = s; + param.stride_w = s; + param.compute_mode = param::Convolution::ComputeMode::DEFAULT; + + TensorShape src = {batch, c, ih, iw}, filter = {c, 1, 1, f, f}; + + TensorLayout dst_layout; + auto opr = handle_cuda()->create_operator(); + opr->param() = param; + opr->deduce_layout( + {src, dtype::Float32()}, {filter, dtype::Float32()}, dst_layout); + float bandwith = static_cast( + src.total_nr_elems() + filter.total_nr_elems() + + dst_layout.total_nr_elems()) / + (1024 * 1024 * 1024) * 1e3; + + bencher.set_param(param) + .set_dtype(0, dtype::Float32()) + .set_dtype(1, dtype::Float32()) + .set_dtype(2, dtype::Float32()) + .set_rng(0, &rng) + .set_rng(1, &rng); + bencher.proxy()->target_execution_policy = {}; + auto time_in_ms_fp32 = bencher.execs({filter, src, src}) / RUNS; + + bencher.set_param(param) + .set_dtype(0, dtype::Float16()) + .set_dtype(1, dtype::Float16()) + .set_dtype(2, dtype::Float16()) + .set_rng(0, &rng) + .set_rng(1, &rng); + bencher.proxy()->target_execution_policy = {}; + auto time_in_ms_fp16 = bencher.execs({filter, src, src}) / RUNS; + + bencher.proxy()->target_execution_policy.algo.reset(); + param.compute_mode = param::Convolution::ComputeMode::FLOAT32; + bencher.set_param(param); + auto time_in_ms_pseudo_fp16 = bencher.execs({src, filter, {}}) / RUNS; + + printf("stride=%zu src=%s, filter=%s, float32: %.2fms %.2fGB/s " + "float16: %.2fms %.2fGB/s " + "pseudo float16: %.2fms %.2fGB/s " + "speedup: " + "%0.2f (fp16/fp32) %.2f (fp16/pseudo fp16)\n", + s, src.to_string().c_str(), filter.to_string().c_str(), time_in_ms_fp32, + bandwith * 4 / time_in_ms_fp32, time_in_ms_fp16, + bandwith * 2 / time_in_ms_fp16, time_in_ms_pseudo_fp16, + bandwith * 2 / time_in_ms_pseudo_fp16, time_in_ms_fp32 / time_in_ms_fp16, + time_in_ms_pseudo_fp16 / time_in_ms_fp16); + }; + + // clang-format off + for (size_t b : {32, 64}) + for (size_t f : {3, 5, 7, 9, 11, 13, 15, 17, 19, 21, 23, 25, 27, 29, 31}) { + run(b, 384, 32, 32, f, 1); + run(b, 384, 64, 64, f, 1); + } + // clang-format on +} #endif // vim: syntax=cpp.doxygen -- GitLab