From 8dcb8408629654ba29b15147b03fdf1a5293556a Mon Sep 17 00:00:00 2001 From: trm <956303669@qq.com> Date: Wed, 20 Aug 2025 21:29:33 +0800 Subject: [PATCH 1/3] finish step1,2 --- .vscode/settings.json | 54 +++++++++++++++++ build.sh | 36 ++++++++++++ gemv_quantv1 | Bin 0 -> 57712 bytes gemv_quantv1.cpp | 130 +++++++++++++++++++++++++---------------- kernel.cl | 67 +++++++++++++++++++++ quant_gemv_test_origin | 1 + 6 files changed, 239 insertions(+), 49 deletions(-) create mode 100644 .vscode/settings.json create mode 100755 build.sh create mode 100755 gemv_quantv1 create mode 100644 kernel.cl create mode 160000 quant_gemv_test_origin diff --git a/.vscode/settings.json b/.vscode/settings.json new file mode 100644 index 0000000..e7850a6 --- /dev/null +++ b/.vscode/settings.json @@ -0,0 +1,54 @@ +{ + "files.associations": { + "*.json": "jsonc", + "array": "cpp", + "atomic": "cpp", + "bit": "cpp", + "cctype": "cpp", + "cfenv": "cpp", + "clocale": "cpp", + "cmath": "cpp", + "compare": "cpp", + "concepts": "cpp", + "cstdarg": "cpp", + "cstddef": "cpp", + "cstdint": "cpp", + "cstdio": "cpp", + "cstdlib": "cpp", + "cstring": "cpp", + "ctime": "cpp", + "cwchar": "cpp", + "cwctype": "cpp", + "deque": "cpp", + "string": "cpp", + "unordered_map": "cpp", + "vector": "cpp", + "exception": "cpp", + "algorithm": "cpp", + "functional": "cpp", + "iterator": "cpp", + "memory": "cpp", + "memory_resource": "cpp", + "numeric": "cpp", + "optional": "cpp", + "random": "cpp", + "string_view": "cpp", + "system_error": "cpp", + "tuple": "cpp", + "type_traits": "cpp", + "utility": "cpp", + "fstream": "cpp", + "initializer_list": "cpp", + "iosfwd": "cpp", + "iostream": "cpp", + "istream": "cpp", + "limits": "cpp", + "new": "cpp", + "numbers": "cpp", + "ostream": "cpp", + "sstream": "cpp", + "stdexcept": "cpp", + "streambuf": "cpp", + "typeinfo": "cpp", + }, +} \ No newline at end of file diff --git a/build.sh b/build.sh new file mode 100755 index 0000000..4bdeef0 --- /dev/null +++ b/build.sh @@ -0,0 +1,36 @@ +#!/usr/bin/env bash +# Simple Linux build script for gemv_quantv1.cpp using OpenCL +# Usage: ./build.sh [output_name] +set -euo pipefail + +CXX=${CXX:-g++} +SRC=gemv_quantv1.cpp +OUT=${1:-gemv_quantv1} +CXXFLAGS_DEFAULT="-O3 -std=c++17 -DNDEBUG" +CXXFLAGS=${CXXFLAGS:-$CXXFLAGS_DEFAULT} + +# Gather OpenCL flags via pkg-config when available; otherwise fall back to -lOpenCL +CL_CFLAGS="" +CL_LDFLAGS="" +if command -v pkg-config >/dev/null 2>&1 && pkg-config --exists OpenCL; then + CL_CFLAGS="$(pkg-config --cflags OpenCL)" + CL_LDFLAGS="$(pkg-config --libs OpenCL)" +else + # Fallback: common system defaults + CL_LDFLAGS="-lOpenCL" +fi + +# Optional manual overrides similar to the Windows .bat +# e.g. OPENCL_HEADERS=/usr/include OPENCL_LIB=/usr/lib/x86_64-linux-gnu ./build.sh +if [[ -n "${OPENCL_HEADERS:-}" ]]; then + CL_CFLAGS+=" -I${OPENCL_HEADERS}" +fi +if [[ -n "${OPENCL_LIB:-}" ]]; then + CL_LDFLAGS+=" -L${OPENCL_LIB}" +fi + +set -x +"${CXX}" ${CXXFLAGS} ${CL_CFLAGS} -o "${OUT}" "${SRC}" ${CL_LDFLAGS} +set +x + +echo "Build finished: ${OUT}" diff --git a/gemv_quantv1 b/gemv_quantv1 new file mode 100755 index 0000000000000000000000000000000000000000..5ab1c830ebbfa522c16512beeade77d5e7693492 GIT binary patch literal 57712 zcmeFadwdk-+3jbhes|!XBO0BvZ!Um!_nQRbS4}&4r zbzP)ZEp63j>)ZCRtyZ)is1=-m0$OW>T5D~q>#6RlRYa?3t$DxK%xrd8v+ra3{64?G z82#=w*L9zdbI(1K;R=6nZgEjjiE@t;^)*$*xmN2Kbk|<>C-(%QtC?yvzeg!e4MQ)r zZH~L{2wa`xx#E3gxoejMCER#7<$QUV>n|%MPjZNqaD|RdUv^xwQdKy*ge#NNDo4+k znNuBmR(iwP>22TX2E3^IVgqG%IGp3!;vi$g#=C<@ zo?^!=`O-l8)Vu{t6wP9Z&82K5-l8x{Lp@-#dERofjYf_0vy0;g73Mdt&~o zCw{k9{7AkfjQEj5;&YpPP_K7wcGk^`;LP&T#j1r+$Bix$HwpQ5=#M#w{xThFsQ6bDsQ=yq zUBLgj1?uOr0{+i0P_8Qq^oO?#wELX}`pMV=`9HEiK3fXp^U?zG{Gx#UfCBN~TA;k& zD&T)(f&3iH1V2>&`B{N}JFY;z9aA9xHy4QikplUAwSfJa0`^%2(p_30pZ657KVG2S z&nXc983p{$EzmA!7pRAy7bw^J1?+l(@qSgGe$mi1q`qklNM9a?%`f&y9bFmLK z9v>}G-Ww^`(L==lW#;Fj)OeNpnd7LrevT)fVpl(R45n)zLA#7qhbi-J$50^^$?-nv zivOhRpOE%ZO}*Ln+Z|nVgjz>_B%ZJv58mwKZrV}&n{GUsYmd3{Z_3NhYt+B^x7_$m z*ZS&C!+ho$Vcw&S-N)B5||3Fw(Gmc6U=#q|=EHKPx<4#4Rw=8!fv(H^Sp4$I+%e)k*MQT+quGsFt=*rIai$g7uwiVHp zp-5+EduRRWjkqpeOn?(Zp~ls#tE;PPmNj%WH-_4~qMekK-$+Aipm9;Oy1H>?LuV-3 zNrQC-8vTnJ`0@J*?+rc03GDv!2?g4_gecw7P8;Jir~J)Lg)7djvYDh1!gr)@ zirq%bTk2QYMOgfAMN>Vsv%4+Y+!}FObWSxD>aBy%*&zR6JD%7X>56o&iulPn{ih+c zqPwAUd8ncBqVDESsgwna+NG_V+wc|*wEF!E7qu^To4b0dQ(26u$g=LHLee3pAqH`Xxf*$u&df0&hGg#=J&QxV?(rYrS!xA)v<5Cz`7_p zt&vf0*Fbwqm%skss589I$W9>38kUn%S2UXsr^!2=OMz3rgME9B-M8)RHnhl;6s4c8 zZj5w9o7>wQzn`ZqrZKm-$rzV%*~9J(46W_UrHe7SRcFUii2aMshR#dqaLslGSFBzg z;=0h)-qz649K9s8sw$i6KEuXs?8VZLK1Z*#Q z&O4g_y{p4u1?R4mpL1%pr-S9qT^-qpOa@|mchnuaT)hSdE)83g9cC@s5@iB% z^3W_p@|zKfq$^^_)!1=~B8-gfIYHTUtQM6}jmI7<>I_j^2hUbC68oUHM#nyH&YaNH z$yI7zFfe;gXv*X%lc%dQ>irAm1W~H8r`1^>OPAt%+T^KIoB*y`U9ILUU0ORk5UQR$ zb#l!| zT8iZW+vjjwAE8Dfv$2fgY_zMDsxkQV5w_foTl`5X75s{IIO$6)CHxj~CsF2VMQCEp zhLhZtgM=-zf4L#y5P#2euT;Ug2KC2WGE(kH_F$7k-1|N5-Ct_vj%Bp^c+kCA-Q&7f z;eTU8FLyx0)G%U^I5(~oy;O}wm!pXM^|}5{^x1pAhkJ!O%=O=3>ppd?t6#F~F3v`& zlUzM^#U%V6pys*y(|29So!fA=#MLjd!Q4E~?ea_cqK8O*6ssS)@n`q#v-g{7wxpue2KW9k zYqx40d-+CJIznajiah(TE^<8Rzi_177R#QwRswb;LrH6pY3*J3}{W6$nCi2Y2DJ-h!P_JtmMc7H(ZXM60~{Rgox_Sm!g24a`` z&mP(Pda=v1h3t{tmk|3|StByL?;-X@9{c(`9YyvTr0%lE29KSAF?b|Bc4wqJPPcpP zY+(!@8$EVOJ9}*M*rg9;56ffE?tO^cPLDm?*TvrNu`^%>kLNsghUMUq^4Q&>WxJ(4 z_ECermD=sGkM`L2dhBC7cJ)g~k-Ymn_HvIsJBGx$!ebxn@ju>U_wJWz9{U$O{wI3u z@_a3ORC(+NWsS(09{a%_yY8`n(PK9}_Us;+IM;jZhj{!i_1M4Uv4=hOLp}DD9((p# zw>Wor?6TLGJyv_{M`Vpi(_=r@z*~ znI60DvDbL)hR1%A$6oKTpX{+O_1I7G*ux(C*F5%>9=rPrhV9znvFjfHt3CEwkKOdx z=XmT_d+c={`+ASv@3C+2*ynldNsry|*l+jP10MTEkNq@{eUrz2y2oyL>9R4DIXb7y zhEV3foIY9f;+%eh=zBgHEWd30WM0ncvf-0?DyPdvPiAXQmkplGJvm)Ab~3l+blK3! zT%XfrBPVlZPL~awOlMA)jhoD}oGu$SnT0uBHfl1nbGmHMWTxcwS)w1E(@zooz??1{ zG@0U@E*mqMJ^vgmzb^XAIbAklGEe1nc@UA=n$u-NCUZ|tmyMXrZ8=>wU^3U|^tqy6 znbTz>Cet~n4@@}tTS_(gjwBwLe+DLOU%jusB~QOUPyc0}{?k1D);#^DJbhiBepQ}+ zd7i#1PjAc9oAUJY^YpXw^aXkP+&ukjdHRezeR7_DOrCy3o_=tiJ|<5emZyJu@xJZx zL7x70p8mHy{e?XJ**yKpJbinfz9mn;KTrQD%-4EqVI=dHOH&^q=PGx8~_L<>~7N^p79HsZHT zL{vX!BsS}u1Y$dkL~p&g#5ONwUJS;6qKTI1I>?CiiWei6Fc8~qUn(7uq&5+$9H$z` z_KKs?vv-{Ds&;a1#O-Jvl!b_qnBF{tgr^YiXM@x{=ezQfTnHofd@%O6^f2)ozt7G} zY#=?{DWs_BZ*dtl;@(28T}t@$7qCm^{RSC`&rKWg>8eAiVC+6iLd5PZ6IuO%5xZ}v z7~-4jC6Pqg$E{LDy%uSte@Qjs5|)JI`kj~1J@|R^Mz?TVbe1vcS(H2TEy-V^Y-hWa zcJq3>y6hTB|4z>1n>8n?yPTw?uF~(>E^{QGo7yA?w@5VyW3Nb7?HZ47Ug-o|>IB;; z!A3bbwA?!pah7f% zVCEXuk`l|k;xG1MF!o;h5SC9;>7<-)buP%l#+@N(#QZy}tzc}ko$PIlN+%-x*qK_7#Yps~2x}y62~%)4 zjpq#pSghAB!4FwX1`~HGsk7Mgwg)NXTFgzpcdUok*Jj50#`aE>YWqFaW+c{&iQpP0 z@j@Pr*o{?al&AVJ-&kezSV#ND>Tfs2j8Bc4{_Z#MEfJ6z2~@kuAcuk2z^|p;4?5-E z&QO;)e_v1WYfmGqCyqG9l*e)S5(P^t`^8a5`RsiTmv7Dt2SNB)CJD z%237*b1(79F==n&Y6)9y^<_%lGk!tVC2G6v_oNe(>?)I(?9AEY|8*lq?t1MGF^hYr zVC>~!>{Du~*7idkJXX6&()qO%uO`;G?HGt{ z?;Y^`*kN+h@pEKs9H@p2e^JRSH zCLm{jqx$cG*sk-9*vm%GtGnwL+da%koYH{TfrGxwH2Z!o1op?|@{lN#`PZn~8$F1v zmCNOvt(5}?QobWi>0R4giU050{-<9rV$!EZY`4+(_NhkS-V&o|yRqxj=r|&n??zBQ zkaF&Ovf-sR%u^oa%Bi}KU1Ic{B9}**CZnT__$kK_EdAIXnszyP-ClMJ{))~c>|pv5 zK65dSZhq>Q8i|93PnPj7+d{Rmy+m`OTqtAjZjpVR^f?R(XSA-T+co(v-yYi?+tZh> zSo1Q`TrokJO08iUzv7CEJ8d~{)d=&6OSd`swkwu5Q3p+hT-1o+)19l$OFVXJY)|6{p_j=Sj&|S2xKD~<}WAA^-6{&Vf z?cybis50G%Ev~Gni`7-m^fQUA`!#v!eayG+Zt3HVyJEr0S%JpBKSVpz#kz(NmIixzirR_Ehhm zee2>>JVRRr8Xu{PEw3Ce_F%lFvOJiWHrt5TRaR6#6)c&!zSj3Z-9d`qX$KoUTZ`=r z<+~H48mT>gjlOqA7=4)$J+GhI^U|q)-vgDG)W&CRsw+y>mdv`6aHWAJU)zqN?X~gJ zt7}V2??Fqy=llAj@A;PY2Wvj=zO^n+@>rhsO@88=!+dM6Cuf8k8R*$oTwC-PBlh^3 zmzZ-QruRV{_lWJNjXZm_?qj{Ne^>3V)o&2`>CO%Ev&qzP= zP2B2cDE$!&7dw?;&0&?%A5ojv*$dBI+5J+xXT;}~2jjj#Z0r0 zEf_zd`Z2js(7bms?wAlJ9xr10j}Tl_4*LEAX>@l}ezRC>@=e zw;RFedz0$_`-q-bPVIT^)H}+%Kst6Xp5$eCKZS zEvC*weZGl1gEhN4Z}v^>OIK1TDS39DksOMBRGwYtkA3?8Qt5Ux>0ouePx;n;M6ZuM zDgFN)&H{~=pKg>mWjjUlb2}5;Y9x-Yku_xE@Own`y<93q{^@+qi~5b&$3~H5(4P&X zu}}QhRaPbD4)ml_J(h;T)W|q9;!}boIDQAclp3WWM8;M!Nf}CiR`c2Uu2E+G==5^oX z*QAdRlgXhkJ)$Qie*7w4+Jm3?tnU@=h+SGairzJ!`g)c5M|$#Xnm_!R90r{n$|$gN zNS5jXJj;BF@j7uT99!wo)bM#)F;*|1H(i*WT0qkO|Y`g z*W2ifo~7Q=BlC05-Zht9HD^01o7tp&yF3YJXdx>JioEvl2k%{G%Rl(Rx>Ai0a z3_Ms){0y?G>@CX;_YH5$QtI{e9k_8vi-`(OUyhZl?-wQBxSkHie_XPgaPxa#^{p*s z%+;EgR?bwuwMTGeto3b)R?d;QR=-IoRxaA#MZP4NEhW^JVlr#63|lSfbET2=A7Zat z!i*!O__Rtn{S(fC*ds>VPIG?zj#4Mhzx&p%mvqd_B*|AOBuRG9kz5aRQ~V`CNzOTB(J(pIN7SWOfd(k-LDROq8xb|!<=_R_da#4K# z-RTx7Q@pIK2|d3?_DGi88oAWTo>L>iL`!9bRKJ!!_YJ3RN~N9?b(Q7Sk5xalrOK%p zgKSE>sPqyk5OZ=l(?#r^AXlV{^jJITU}c4rF!qjgxpF&X`YTc>x{Xw5oQDZgi`z&* zyo{G5*E6zgFSo<_RW{sSp63jP`1QHYNop^eL42@A^c<)@tY}{XMj8F z?b@>wwMVj~ialwss*Kn_WMMV4_c0^Bs@xaollfbd#frh?Skbq)tTy&2*9TUp#-v?F zO`q?|nT!*IE9{l>`5Wl~TA=Y!BffygoHJNJ7?buGyFNB*w)(D=3zXFVl@=%d-Vc21 zHgHng^LEj1G-)Mzn6Dms}=s-MmkzuSwr+{_3jroUtNy}HVY zxhYF}O9ht%D=6KL$c(;JSs*b^%})fC-n5=e^ntRF3dZ_to$I)J9boBn*gO_VSNV&c zpHotIBTIm`oxa77bD`eJRVGlgtD7bDJFogqc*1wp2jr0XzHG+PjGgB9rBRrr4Q6w? z|JbHj=}Cjvv)4=e-WU;xeZta-Bu1Uh(&?r;)=u%dURBi3HFeaRN6}yHD}>bX?5j(; z8az@Ita&BcWS54L%#$lnFlN=RNjXa~`_d(A6FbXwMO$m#EL-kn=$W*08G6D$cp19l zhjKCaA@_u8_V!ZEz8;WD_YVv%MdOQ&o^6ysYOIDD_4Ue{iuKsceXc0d3T1Uiaj7Wd zQeh-Y=Suspooe?>se)K-d1iuWMO%%UXMJnmVT2g5CFPmF$;qUCsUv)CPj}JNGyk-W zJ7spkWbs@zvqgqfu07rLPVD!x7E0=bv5*8KXcYP`2w8u>c&5_lJZyo*;ytfN)N4z{a&_B>nioygr7{J0+n1*Q9>1M_GC%fwAeQN|is-mAzD!rAV@jx{Qt3-6M7@zXUz;E6m$g-u+^`hw zw%2yziRIH;85VY*=2n8`YrKL{a4227IFJYy>+br>_W%tk_X*qG_0-!Vdj2Mh_Dj1} zZG77OMSE*YrpbC`du~N#)I8q3LsnEP*s;y7sJ{EIy`t)iZIu<3eFwmbiio0@5z(wI zMNhJ(dcd=$nj$0HSyPEGKfX$@DtfmzKI*!nKH8-;5HAYG&Sh4oJmZ^Y$Hp;9r{>50 zVy_ojRu>)zG$$7-N<7E+z^r9PY_D@8nM;i&oSoXo+~LSN&lv@!=_g-vM#1*kxlwS7 zjDlPC83lowzjWUmj7#2R00d%h$zZToUfgtOfgY@$Y2aao%|(5W7z{{7Jfsq|tOcyqys?S$u1kjTTQk0Hi$^o{l34x8-98$8#r z+iLok@XD^=1{pn{ecG|3SDto=n;oj?Iax@a{is{^BF`$W;C zrLK4d7w17qF!pv@lVEYXUhLJ8tR+5?+KAt1KVvC5CoQWzsj>%0lK&6Vx7rDnohk(| z=(cQFmI`XqH|ZxD;sbjFz0B9UatsAa+$WDcs4nLr)Q7Cf z1B<0OdgW;l8I@-w>Fp^wzeT)OKgDDfKm1%tlIm-U-MF6e(G508G2?bHwA)21;qeQb zwDGbT_%?R*UX4+zim06AhNq8{DoB+5gfz+ceUd@DRbstvzD&E%$o9Ut`6k^bA$^nV z=VXYmqjezC6;5sZND|9F+_}$wa29i(pCx*e_#+?uvVxUT?xhnW)?*=q8xWGg?Y2ji zUT4?7^T>_{y5dFYJiU?_HCADZ&B_T{?gnKj$%7i=Oglo{x7(h%525LN^kC{qG9Ki9 zI-Xc5XEj#PH^-uFgTA`0QpbTwJJY91@`Y8 zRuHum(q8ZmNV5$-9sABgyTno<=|3<7f2Z4kHCvQNkoWv4h{HcFUxUyx|a zj>fA8~0drXJ@b_%G&I>dZovwm)YJo&y)p z46;9;xbb$W(pax8;=R&_XyS+VH4%ZBg~1*J*HfAC&0@q-kB~1(}7`%3c#P6$; z6zFU0))?_ITUcJOTeFpo#q^zIZ?I6CBx=&V+{tCR9qM`8yZt!CPTPo=r!V$~vfNNB zoKOmt<fkJDdu93)!DfqeqoVXUM&e`4G}6K^|uuf+pP z6t$!G2KOxXWUO}s=R10nuElrP?%od)XVEs##d`Xq> zu063Ef6Te>uE*`${@8u%(d=htvBb@sQVI#0zJ_kWxrB4}2e@Y2`Mbz*kbF4~!(4}@ zj)TPKI8ak|h*KR0$-miq}4w4_oL2iX5#QT4-b0r0L9JnCa z4u5tWq-c(V+~-S(t+vDbiH?;^zHPn3u~s=&u1L1^CdWF{vC7tq#9%s6%X-6ckU3l& zRyYnFgATG75{Cui5Z}BS2deNnglz8ETS+^5Ma3V_DTycj-RH$?Vkt(s0Es`jk#s|E zDU%(K9nRYsn7Q)kpU^nx5sFd&Yl-r~jU-E|T*4*6AWy0h1QLO}O1Il-*a?%AMDUyC z*m=ggSA5p=8NPWtH_1ksd(K>8$Kkq2_cOSy_SNk)8vPTi+4hMRF|reOa2`Z`$Q|{w zjC+|6-TeU`3Nf(4#6LgwZuMgqDErwg%MI>!E>IV!^y5-TJ&$O)S?Y2|0IumI+3k6= zsQWh$+pl8$NOGT8&0^7A8*C3IHi?V(6+w?T)brBOP$J1U|`jbJn-V+PS&p8YX8RUfBzuU2hl% zY?T~rC1oqqEw{Gpd|teB-4tDj1r`6tq2S;|NM76rY*;_n7~?w4NtR(<$@%4x7&$&v3B5Fq(#|xk(WaT=Tm$9JlG?}%ibTpnEtKZ1MD`T z@#P**wxI8nwQ2e~lA*ub>(sa%nz>qTp`01ojw*ePM36ZA%k|9Qa{VDAXOj`xP(H!V z=8ew#Z0>t*FQyL>t7l(&t-UWT@A%0$y@x2w)iYJ}65sv$95#m~qJy87AUo1qo}q!~ zd@42WDZH=uO)({=zfJn-uiA_DeVm_goL>;=eK0FOG;}gRUkf`&ht(C z9>2b7lf8>w^SE#A3g#4hlbBW3z*Cu}GR5qhpKlOjdJR!z_kBO2%-QpoJm&||wX*g) z&-wFo%dCAn=b0Bhfv$8t^3WtR!4s&G0hJ9jYVdqV=5Wrl`&Z)bKEI6pp>}aCn|X_j z*gxem;=8}>m`0`6(g@EoHvGQ3UJEq7WyI3#^PJ-Al|5#AhTSoTp2!3`w#0sx@B{p^ ze7;j`@^a3DwAF&dtSTc`!a6*d2$u$9->PsfWJb@wilT?julXc8CXhJm6IpV;o{l{w z*Va$6`~QvQ^K0Jnt^Jhfxbr=ZH&5B~Jbbb2U?h&eWtp7`oX!?WOqK7 z9KI*lA0CMFQ2guLd?#%6_1sKkJotTtM}4L5)v`zNiJis|B#l_VaY{R7^L;xhYNEb~ zt)xe>AGEPsocoxUN5&pAs(0GwFAZPwl6lr~w+!mOH7{|)`%YHR9LSk37L`sB>l?f( zWsC#*BwzBNexJNuxg4LhJN$|+^EJtzY9Me$SZh;+gfkl?bdH%dToQ#P%qI?GUC1<-G4{f*D%kf2DZqLyu~5) zeLP0{xs3gqPWlL=dbiPRMPkES64S_B^mi;Gcr3r>qYi4Y`+1_dILgOdP^2ro(@)R6 zmzjjgvXNFwpO3+_xtx8zDQ}ZVZ?vCl(hWZhGJC!!9kG>suwndHqmc*y{dB~y`Ffwk z=5$2A+Y$Y&>Ynng9gjmWQF)+rs%g9A0{3$D2d5==o z4c5HrTl+(tc~ao?sf#-T@o|p?6FNJuWql#1xAmn*;$FSgE>LRt1gAH-s&9hRnH)7u zXWHtE$%KRNo|oMvPTDx2KikPI=MTZM<6FvpwcOdo%g!V{`zf=%PTb*C);Ac>?h79F z`$2S;vQkQtsU{71ZpY>6bf@+Ds6y!$LyLu6+><*X~!$`RcyDsOy%sqJM-J_+QrfrmMrSN!08F|?4Dr6 z{^c7x-01xvI@3rT#*?oX)5rZ@s_uDangmZ)oPH3DMJm!?AT3VHWq57NB`R+x?U4ID z*4|E}gPzFctiy30`*N<&2A2-0v_ZZO=SP>~{V=Q4f~K_^kgM|4sjq z|EGT*>#ZFrBgcJFi^pFw%fH7gKRy`0m>xfmo%U0MlXe?5yL~;Iq_mr8e>tw0Kj~A~ z;mYHQ<1`-i8L`qzgZEJ6Nl1C1<^$iAkM0xXjKpy>0yU2Yee=??sN6&aon~JF9^F2i z3F4ncvaQC`YlHoHc3n3-Ayv-{vwLJVd!C{Mt};;ccChI0 z^LfLjJiuFByq@vM`26{BPwn<$Dv-FG_gLO2>OL|ckE1^8FZ*5($sa=UfyAsM-sHt5 z71-<)d*BKqURoKfNe6v%?LxeVFN3!P%9%MHr4p*P%NvLDsapH}$@y&btX)FbY}PY6 z3SHzJI!}7lZ#lP1oLDU-{weyK?s((xPwAj%58XEHZrhv~oRl_p?WK7-ryE6ordH(< z*Ym-mH`C`3SAgd@`Gs`m_iQ-_+w;_5EEAmct`U2ZcPkBPasHy=x8CMEHJB)!YVd@p zoV(_rG_`Lo50|&*q6pMj0pDE9sCj|9Fnn{~FpB;zP4|1dm=Z8C?ohT|Wg~A^d3w|1 zT+nOY4oZc_J`ToqM-LwyzrF{^w`D&gYl@xrDZQrqEh;rS-l@mfhmwKdqz}?P#4Rz* zkIky&LEF4Q(K`X@e)HuScwYNb?XR#(OAkCY5G?&q=J8?iUiFB7?}0E#k34BuRwot{XI*D-#2 zKj`iuOTj;_kcp`*Ol{@q1K|lSYhJ;dcZR$4yN4Nc8cOjd&Li=q@BT zC5Q^6cOZI$q{W-EPI8$G2%MX*veBgH*x{uZJB*@t$ceAFlOCB_#1H*E{pCUToyM^x z#-vAmYp3C4htC}1^xcO`sl*3&dgHg(i`3pH_A_6Zx#DM2ZCma43HJD|yg`Z-KbvI# zXiVB+>|z(@!>-eK!ytOFT`#@|<}(9W#m=k$Tf5((Ak|hb|Ki>h8#Nn0MMlkb+!Ljp zc0=0}^LiO8N3tK=!BB{Pp^10;yzz;Dwry&*Mt9}HIPHO%D@6B&*$aV`-wvFZHooVT zV)}RV3oQ1?7~OtP^o!Z`V{8P^Z%d_bV(#7@h|j9DyJmoX8Q{q<6A|yb(k{NW=TK|k zpgdebGS5(Jbi>PJ+hp|)sq)E{ zVEXu@3u>po`)(ut9f{nCuNm*AF@elqVXS$X_A}WrVWDb#_?*G+$<*(xW;0^KiyOxL7GrksZPb_hxBZNlZFVG|{VC#GtHk)QQlHM=eBxEpIX?kT#dxZ$!2Jssp0V&G zZ9c*6xl% zuJSJtEXOhbpY+B;7=iLYBP zk2E!Ow?sp3wYp8+*wEJ29@U(0e{0m+NhdjBw3dd>6_HLYy0W27Yi(E^>f#gB6DMVR z$UaS_4SqSBhSZum+gr7^?v@s+~Cm&P(-~4Ci$CO_dhbG8? zGG$OP^_f8I$^<@?tW1zk2b=O}g^9-mJ|nG6+3IC2ZGvLuZ_3`vNF?yFW+m?{@R4X` z%3}_;WlWIIteW$XCjWz4g9LRY5>yWo)LJB{UL>fikf6Se1a&nMR168~8YHMV64bRw zPzfZcbx2UxA$K4-%&GiLQ%VODl3+m!6mv=iXix<@)I%6Lz=W$|J#2s^+zuOI6Iiel z`r$c9K^k_$UQldDmqIy=f(jT1_a!wql~Y=9rY&2TFu;Wqd&{1k47pTiySOV|ka zz`bxkY=Q^jA=m;IY=uW)JM4tV;R$#W`r#>f2A+lI;Cc89yZ|Y98D53IK^oqGx8QBq z4SV2y_yG38$M6Y!3W|psMNkYSPzq%*42DBF8~`I=B#eU5Fa~^30b}7n_yUZBgWzEJ zB8-Pa;7f2Q90rHO5pX1Ea1rU?NO{ufnl#92^gmK_19dK{ZT)sW1(u z!wfhPW)Rjn3JxEY%k)V2!psqrK`Zf~O)ksh=B&chU zpyEhS*CIhBkf7EfL0yLgwH^uTJ4jI1BSC!^3F-zUs2h=>ZbE|k9um|BB&hErLHz&; z>W4^BHzPsaf&_Id64ZYoK_!u(euM;d8xqw2L4x`*64XzSpni%3_1{QPwJz zNKn5(g1Q3<>P{r6Um`)>g#@(`3F>YnsC$s0euV^eFA~&!NKp49K|O#3wFwF8*GNzg zB0>EI3F;vvsLe=FTach0MuM`Cp!$%Ywjx1oLxOq)3F=WKsO?BlJCLAuB0)Wd1ob!) z)NhfXoRBYHKO#XrhXnN}B&g?+ zp#F>m^%o?lzal}sfCTj-5>yHa>Lnzomyw`eL4tY}3F_$sDB_qeSie@ArjPHB&d&&pgu-| z`X>_9CrD8LLW24f32FcdOprQgkrE0#$|%CRm_` zp@Rt)sNv{ff(5D^9Zax59e@rdSfECrg9#R>k?3H81!@#Jm|%e#jSePQpvIts2^J_H zI+$RAsz3)5EKp<7!2}D`f#_g@1?mgvV1fl|96Fd_fjS5sOt3&5j1DGPpuUI>CRm`x zqk{<+s6)`f1Pj!c(7^-?)S>8Lf(7a@bTGjJbvQbhV1YUU9Zax59f=MmSfDg?Fu?+K z6grq-fjSx;Ot3(G868ZpKvklH2^Odc=wN~c>KJq|!2uc3np7N}Fv!2}CPT zFgloEfoebp6D&~6(7^-?R3kc=V1ZhW4klQjBIsa(1*!=hOt3($KnD{nP%F{F1PfF% zI+$RAx&R$aus~gi4klQjTF}7+3sfsQm|%fwLkANqQ0?eof(5Dr9Zax5U4#xMSfD!5 z!2}Cb7dn_=fr_Go2^OesbTGjJwF(_fus~gm4klQjR-=QwBVvKN1RYGUKz$P(Ot3(G z3mr_bKwXLsCRm^@LkIbDNDI{E=wN~c>I!r)!2%{ArGXA%Fd+#Rq(Bwn4myOvgd|vy z0#%GV=nw`Il3+m!R0-~&Ll{g*f(0p%^@9dFgu#R)Sdao)PiUY+7)(fl1u2mAg$6o= z!Gt7OkOEn6XrMzFOh|$SDUkJt20Dbngd|vy0$Gn}phFl;NP-0^koAcMI)uT5Bv_CF zS+8iILl{g*f(0p%^@|2Ngu#R)Sdao)&uE}S7)(fl1u2mAjRrb|!Gt7OkOEonXrMzF zOh|$SDUkJ#20Dbngd|vy0$C4fphFl;NP-0^koA!UI)uT5Bv_CFSubgzLl{g*f(0p% z^^*oVgu#R)Sdao)Pide-7)(fl1u2mAl?FP5!Gt7OkOEn6X`n+GOh|$SDUkJ-20Dbn zgd|vy0$Gn~phFl;NP-0^koB1cI)uT5Bv_CFS+8lJLl{g*f(0p%^_vDdgu#R)Sdao) z&uO4T7)(fl1u2mAod!CD!Gt7OkOEonX`n+GOh|$SDUkJ_20Dbngd|vy0(BJbphFl; zNP-0^koBPkI)uT5Bv_CFSubj!Ll{g*f(0p%^`izlgu#R)Sdao)Pimk;7)(fl1u2mA zr3N~L!Gt7OkOEn6YM?_HOh|$SDUkK220Dbngd|vy0$Go0phFl;NP-0^koBnsI)uT5 zBv_CFS+8oKLl{g*f(0p%^{WOtgu#R)Sdao)&uXAU7)(fl1u2mAtp+-T!Gt7OkOEon zYM?_HOh|$SDUkKA20Dbngd|vy0$C4hphFl;NP-0^koBx4O#>amU_ufsNP*nH zX`n+GOh|$SDUka)4Ri>D2}!UZ1#*9;*aB*B6dFzG1`bO?h9 zNw6RVssVSylQ3D;qU_ufsNP*l> zYM?_HOh|$SDUkb14Ri>D2}!UZ1#-Wsfev9XAqf_wK<+;^&>;*aB*B6d$o;4WI)uT5 zBv_CFxj)rFhcK9s1Pf9i_p2J{5C#*HU_lDx{#64V!eBxYEJ%Ue&uXAU7)(fl1u2mG zTMcvwg9%BnAO&*2tAP$-Fd+#Rq(JU}HP9gpCM3av6v+Ls20Dbngd|vy0=YldK!-4x zkOT`-Aot4}=nw`Il3+m!10BL(LJ}-Uf!xn)phFl;NP-0^ko$WLbO?h9Nw6RVa=)*E4q-4M2^OTT zm?>{fu?OHiOl$)fY?*b>N_jy-j;xf|szlGq@7gae+b>aw**?s*lWXl0{W`tq(2|qo z8%;+*#}5YvswAT1z`#;u1+oLV{>K9Y<9YL96S8U;J>_5z4&3n`295@7#>6LT(g0e;_lFcy<23Sy zzhwy$uf^Z1ESC5<#yLV2eRH8IT3u0e=&0f4>q$fGd_6|h|6*X^81Yj#+^)$=PUR(x znLprTmBgvWRLmVS{xsjni_1-Q>X&N1I<@j>takWwk;z+~PZ@|!4v9Ak8?c)a-aRC8 zc`2lWm*S~f`xwt}5*{KCRbr_fQ&BsneEt|bi~on9AGepqjYDkxa0l(d%PcH~##9)H zSK)SShe@5!9aC{#@$50<*Okm3qa{jbkC}LFS?!prtA^E%nb|YEZp`YE{-P0MW@4%x zGZ8O%ojs;v_L%b8ky38yAD7aef8t!?y0*A>%=oKHYR70jrN#G+7&D$c;@KEEKYPlG z3a={(S9#aKz_QPihxJ7zV~U5 VoSN7_Y>qHop*?N%H9!yi#shi2PO$4%O)l799y z{h+jqiYBj#;xopKzgF^kRcYOri9KZt#_TRGxr&q~Vxt^RLF=+%K#m0nwtC}tq5S(K%p_y51mxe7>K_;YnJXRt0#A9?oRd9B1P?X;S(mGtYQBrJ#6 zuEnP3*^=1K_1L8E+>32BHW_Q`j3L^yZm>P)jj1p08?q+rMxHZNE7xt?{>1N|{Jk-z zzGwJV!>%n$l&&kej`Sqm=X0GR2BR4 zdc@q33-c_-$YLE6!3Mj{86)e6vYu{ooXk&d{MQ!ONuR(YSCjqNfUV>T9*Y%El0;*vU?M!9cbp4~}DEB!8+ z^N*>xs(8-mVY9C*S-P$?QFiUHp5Y|}2Taz^nXHPhnXFwrS(QvEnyeM~qW_-}OpA8k z)riX^E|dA|A@WxGB5{oxRzYtmzIGU2)GfWPMEdel=Go=c!-Y~09Fm6?Y$@JdzH+ck zxp^p8CVL*5H%2eHLFOTfnKuvBjhsbXn~1BDH`T|wajla%Eg^m4+EM~sCH2!YtRzGn z(n^D~nAER~-8$;`eXeg&Wa*r2eas&{tnRv!b)^fWKF***b`?*qT1o{(vlXBbR_5(l zgl*@%bna)%k6H08ilgT|B$ZyeoWJ+^P1*PG_sX!ZaQ>6w9wy1!r_UXxVP1c@ zm^U7NCg(po;-cSb>h2@waQ@+u2Ip?O4g7FI(V>$GdHM1?O4VCMWuKO+|1PcN{G-yc zwPjgvQj(Jcc>6}#t>rwED7(H~eXpqOs&aKlQQ4E_>dxY_t>tQ6N!jz|>aLP!itaC0 zYloHnyj9+gR5trn&WZftpB#UYn|KaB2x=ee zZ+GOOu3qlSFS&Aq>z>_%%)0;MQ+qo=$-c83*?ZZvo3Ojxtj4o$Mp*x9|~FS_n~-Hc}KNw*-`d@pq)l0TG`Bb)vajw*jl zX&-#<%dRQf!8Y-?%(h-u z*p(fwG+nvgl}T4_bfx9WepjYkx!aYh+DX5{m6|K7T&cUV-j!ijcDT}X<$6~pUAfVf zmMi;RnR4ZBSE?y)`mWSmS>;OImG!O+yRyTTrYqOGGU>{VuC!d)@5+=bce_%_U#xOe zxKeXvl`C~u*1Iz7$_`hWu3YcRq$@YN(sE_LD^srA?MnNhjHE9QDL6D&R=HAlWxXrI zuIzB7>B{x4OuEvE?ce_`-{^$-)kXV;oi}IBN!r9E%evd5-CFh3$x|m+9Y3SnR;OG( zW!mJbX_H*@zTxsUnGZ|VyYgIcpF_rR5iZxd6NcuxegQd@|4Qt|JbkEe*9qC~e6nlT zUHi*!JpU(auW<|$UHcQReWh!cyMcAI-w#XeWw&ekWod zDxW9Y_6lCfbiB*+8#(j>{+Ac9uPb2xSpoZ}1?-12K@FA9D(qv64^`RgfRjt!9xdR% zp@98c1?=C)F7x&z5vqaPWYqTcz#mA|C0smZ;QQn?BKfGjYhp+!2cK)(nFQ&s|D=y3fRAnUGi_O za8mrXn{eE!cck?rW zdEh|e(cJ6TF|NJO^}jyP|3?MlAIbg7Q2Epf*ySUsL;07_NeyMczJUFgw!K2F4>}pQ zpY?P4Q33zsxPg;;NV(XT>qEbc31I;4F&AK!ah_# ze?t5hli4ayB%D5lUFL!O>&OcQ;y-E3(Dimf0lS6$N0dyDy6I|;JcI%cRsUbL?G?P9 z?f6^fUXQv9_}`iJ@7A0B4iQ1+A-|n&mlF?*68kt8`%vkc*vGP7(lpy(A<3BqueO$o*k@SDLUguv2kF)KEsq~+mc(T9M_ISx^Ik^D4w6B>r9y{H5viZ-3 zxE#CmL-ljVV!96Z6^MT{l`~ZR%qw8;Dq!Dg+behl-$^iAyfW6&L&d+WfPGZ~`**NQ zJyf+i@!RjSaJmb-w7b>m*iA>MiQ|S&cXk2$GTY90k2n$7?{RSYE!SRkhGVzif8lgX zfp}glVE+mmC_~l5*RX3&)cgFNQ^5a41?-O$u)kv4skenr^zy!q9QPbNbbkJTeW?ES zQ33xN8#|gEwAJHm%374^^LUy8e@{ z|Lp$JXig-aO?l&DDfXe_>BcVk%)bt*MU9=&>dEaY6lz%39EvuqP%9#>t3nrbH?&1p zRZnj0=un|^W;U*jG+yZM>};+-y>W6!LuZsv*@ezoP+hgMp`|I*)Y9G%ov}O;<+FCx z(`{{v?a~yN9cXZ;*`G{Yjt*O1fdTOY#g-^(>YHo~#nwL**Y;RdkhN`Qa@3}3G zbVV1QzF<-G#89Yl_3G;C>M7n2#045jx4Lm-dX2V&s-5{i*zbN}>M;=OM#a&-j0a&pEd?G`Bf9p`Ua`yCq+0d89M4qPZ&?q5fK_koGn{u^0+1 zZx5|#X%2mL-`0bHMcc~8ag{0E(t~2qMethrp|`e zNN9O?YwINh@n|7@^1VQh*HCEg!rJ-N^Mbli2+#8tED8Ay*J#u&RH3@BFQ}a#n4?1T zf@jRG4TjE`J9m+PacFVv?4Vyl*d41W)VT6OeO@pydroKyTSn8QZ!c<}%9kSjjgEck z(oohOoiepG(%BVhYm0=U7dzDuXnx84tXPh`~ z=F}P0(^0C$GHpuDw3-WrExGiOYz5-(z%RyAeHiBqdioGwvT&$NH4 zr%anVb;|Uqem@;`euy3r>ZB|BSJ_z}?1*lMqr2vRXwl!ww+k0WS|SZy5&tS`TtB1U zzhF)l8X?@mY(aHO??T7@FZ8=?(P z9am3tE1{`yO#b{x2>L(HV&Z4Wk?y2BlUfgvbY#wyeKY7d|GQIN4PC?e_NJdnz13TN zQsTj`;C{w)-@d@qI+&fh{On}EjX6Xo3Y@|0zr49E)ZO*juZ?~tiR`r6=rm0AG_p)r zXbClRbv3VO^S3TMy^)M&KR0Sm%4f=)yziGdv+>IJSLJ%p5L(9c)fQ@Oh&Ha&=PX@X zJ3B!4pXxd?A@PCGZ$_LNVM3TPdAg#hTOw`ym}F{ngEX`-?=wAzBC8uC9Z|kIx=*y{ zEQ(I43WcI8JKHY~5zUI|O8J~(QU{a%Q)&Q&(Jw zriHwTIgR|8&%1ilY>l*bMWXu@nmN3yJ;XOym$yWm>^iM{c6;ZA^SHcs$mQ7^hqp*| zUHiJY)6o1#>lw=~;LEW8x4Y}-aifUhE73s3N8ltvfugL=KIf9_k{m`lMFHz<2uic^ z*?YDXuh-dKU(OAps*Qice-`T2x1w*z>&SFZzMn$Q(}6q8!(WF{f-axCj?eph6EU1LdR1 zpq_s;cb^vNI!+X0i)~Y1A?_48+^)=BbhTNX%TN2d&yr+1IT=8F;r%7T`<5^ES`j#9 zFkjl5hPI&A99^-0P-YP0ZWG4%$p>(yZV#&EJcTcI7w*LK$uOO@MaCf)hF(Q`$YqLK zHnA2h`Dn8!Jgn$mas-X~Au6(bXvE#vc`6nC-DPhXL=L%0SIkRo++g}X7Q7t$9rQkS zMfv57lV@U;Wm0x|Bb($(KURD8+*A6f8BRNAnZxd}=2om!f%j?nFtX;l72BrV`cu36 zqAxezki)>*Vy}%_sNqjBqAD_ZkUYC9^0(J?-R~>5?1I)O&?H z^i+NQy3?=5SQRgl0ukM`(hBr`G4f~-K^r_3$25_zh1Q9g9F{$Hufif}6aXw7dIB=OLENke~%bbA~IfkfujkUKJU>YKt3i zp|&w#jV5}vU-NEfp1IWQza|QmxT>k6PKiRVi8b4e79zT8QkH$hz$dV6>G>qEd{6F} zp!g+N&JPIpVXqL|ZMhV?Sc>7KoUeGhdL3F4w__JK&`p@YUV?p9aVDsFRZEzj&k+jF zDzk*jY=Qf<-*eI*!s9Bd&{6MJtwFvpYwdx~*`rq88$>zY zY_=06;h!R%Py2(kon>`Czv`D-;HsFbX5w&h8w^myQz-hyQfUHjzvI2;cHDobNa0kS zfpT{xH0L*W<&*)x>Tv(0%%^}jsLG2CrL=onc^;zuJR=9YwA~3(NFQ9IEI;$33GJtW zooZr`4gHI{V5r-Wxx6B2Zsw<_@ct>w$?3YIk{Vyw@gYFaYEUcZOL$-u>dV-06AB)S z4HqMBda*?68!$X`P{f9VC|0G8HI9rtfCuZpno-HOt>Jjk$uRmGN!vV3Nu9!QsLc>Q*7D7Vln zHr^4_XU&o|eZ3wC+vJSN$xRwIaKp;zZ^&La5=*e8dimVa)#I*$2i_if@e4r=nI0mnop}J7DptC zaT$+tb5W*)D3&(}QLIAlxX9H3Vw{f6iS~Gi-gt*$1_;DY3k~>a( z=8F^lVDLyk+Yx?>=aP%gd*-8CXa*Geky>ZzH{HsQ@E}k0l78mL6EZ)Zu`G@(BgcmnF^k~PwjeZI6nP*G*PtZOb|Mw983ypAL;3t+bbUpcp-To2azb*`c zLU>0O$0_0Xf8!Zah?ERVGv*-^hU4*`@HO<06d{Yxyk|nTlfI?!p72fJ3yDT;+I(K- z^}ethqtGqBzQebGfGkTyeC9I~?izfNe<>W7eDoda-0}JUYY0DIGd}xyC!pWvA6=;7 zq8$zFU*fUiCh(<>^fSL`$Kbzsmn}8C=ijRDF$w*;W tyoG #include #include "half.hpp" +#include +#include using namespace std; using half_float::half; @@ -135,6 +137,28 @@ inline int8_t clamp_int8(int v) { return static_cast(std::max(-128, std::min(127, v))); } + +// 激活 B 的对称 per-32 量化：输出 qB(int8) 与 sB(half per-block) +void quantB_q8_per32(const vector &B, vector &qB, vector &sB, int k) +{ + int blocks = k / Block_size; + qB.resize(k); + sB.resize(blocks); + for (int j = 0, jj = 0; j < k; j += Block_size, ++jj) + { + float max_abs = 0.f; + for (int l = j; l < j + Block_size; ++l) + max_abs = max(max_abs, fabsf((float)B[l])); + half sd = (half)((max_abs > 0.f) ? (max_abs / QMAX) : EPS); + sB[jj] = sd; + for (int l = j; l < j + Block_size; ++l) + { + float v = (float)B[l] / (float)sd; + int iv = (int)roundf(v); + qB[l] = clamp_int8(iv); + } + } +} // 量化最后一维 量化块存储 // M: m x k ; produce BlockQ8_0: m x (k / 32) void quantv1(const vector &M, vector &q_M, int m, int k, int blocks) @@ -192,42 +216,20 @@ void quantv2(const vector &M, vector &q_M, vector &q_d_M, int } } -// kernel 示例 -const char *kernelSource = R"CLC( - #define CL_TARGET_OPENCL_VERSION 200 - #pragma OPENCL EXTENSION cl_khr_fp16 : enable - typedef struct { - half d; - char qs[32]; - } BlockQ8_0; - __kernel void gemv_q8_base(__global const BlockQ8_0 *A, __global const half *B, __global half *C, - int as, int ars, int acs, int bs, int brs, int bcs, - int cs, int crs, int ccs, - int M, int N, int K, float alpha, float beta) { - - int row_id = get_global_id(0); - BlockQ8_0 valueA; - half valueB = 0.0h; - half sum = 0.0h; - - for (int i = 0; i < K / 32; i++) { - valueA = *(A + row_id * ars + i * acs); - half value = 0.0h; - for (int j = 0; j < 32; j++) { - valueB = *(B + (i * 32 + j) * brs); - value += (half)(int)valueA.qs[j] * valueB; - } - sum += value * valueA.d; - } - - __global half *p = C + row_id * crs; - if (beta != 0) - *p = (half)(beta * (*p) + alpha * (float)sum); - else - *p = (half)(alpha * (float)sum); +static std::string loadTextFile(const std::string &path) +{ + std::ifstream ifs(path, std::ios::in | std::ios::binary); + if (!ifs) + { + throw std::runtime_error("Failed to open file: " + path); } + std::ostringstream oss; + oss << ifs.rdbuf(); + return oss.str(); +} -)CLC"; +const char *kernelPath = "./kernel.cl"; +// kernel 示例 // ---------- Host 辅助：编译、运行、衡量函数 ---------- void printDeviceInfo(cl_device_id device) @@ -251,19 +253,23 @@ void printDeviceInfo(cl_device_id device) } double KernelTest(const string &kernelName, - const char *kernelSrc, + // const char *kernelSrc, cl_context context, cl_device_id device, cl_command_queue queue, const vector &BlockA, // m x (k / 32) - const vector &B, // n * k + const vector &qB, // n * k (int8) + const vector &sB, // k/32 per-block scale for B (n==1) vector &C_gpu, const vector &C_ref, int m, int n, int k, float alpha, float beta) { + auto src = loadTextFile(kernelPath); + const char *kernelSource = src.c_str(); + const size_t src_len = src.size(); cl_int err; - cl_program program = clCreateProgramWithSource(context, 1, &kernelSrc, NULL, &err); + cl_program program = clCreateProgramWithSource(context, 1, &kernelSource, &src_len, &err); checkErr(err, "clCreateProgramWithSource"); const char *buildOptions = "-cl-std=CL2.0"; err = clBuildProgram(program, 1, &device, buildOptions, NULL, NULL); @@ -283,25 +289,30 @@ double KernelTest(const string &kernelName, int blocks = k / Block_size; size_t sizeBlockA = (size_t)(m * blocks) * sizeof(BlockQ8_0); - size_t sizeB = (size_t)n * (size_t)k * sizeof(half); + size_t sizeQB = (size_t)n * (size_t)k * sizeof(int8_t); + size_t sizeSB = (size_t)(k / Block_size) * sizeof(half); size_t sizeC = (size_t)n * (size_t)m * sizeof(half); cl_mem bufA = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeBlockA, (void *)BlockA.data(), &err); checkErr(err, "clCreateBuffer A"); - cl_mem bufB = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeB, (void *)B.data(), &err); - checkErr(err, "clCreateBuffer B"); + cl_mem bufQB = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeQB, (void *)qB.data(), &err); + checkErr(err, "clCreateBuffer qB"); + cl_mem bufSB = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeSB, (void *)sB.data(), &err); + checkErr(err, "clCreateBuffer sB"); cl_mem bufC = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeC, (void *)C_gpu.data(), &err); checkErr(err, "clCreateBuffer C"); // set kernel args int argIdx = 0; clSetKernelArg(kernel, argIdx++, sizeof(cl_mem), &bufA); - clSetKernelArg(kernel, argIdx++, sizeof(cl_mem), &bufB); + clSetKernelArg(kernel, argIdx++, sizeof(cl_mem), &bufQB); + clSetKernelArg(kernel, argIdx++, sizeof(cl_mem), &bufSB); clSetKernelArg(kernel, argIdx++, sizeof(cl_mem), &bufC); // strides // A:m*k B:n*k C:n*m // A * B^T = C^T int as = m * k / Block_size, ars = k / Block_size, acs = 1; + // qB layout matches B: n*k contiguous; for n==1, row-stride 1, col-stride k int bs = n * k, brs = 1, bcs = k; int cs = n * m, crs = 1, ccs = n; clSetKernelArg(kernel, argIdx++, sizeof(int), &as); @@ -319,9 +330,24 @@ double KernelTest(const string &kernelName, clSetKernelArg(kernel, argIdx++, sizeof(float), &alpha); clSetKernelArg(kernel, argIdx++, sizeof(float), &beta); - // NDRange: 1D: (m) - size_t gws[1] = {(size_t)m}; - size_t lws[1] = {256}; // 可以动态调整 + // NDRange: 1D: (m) — 选择安全的 LWS，并将 GWS 向上取整到 LWS 的倍数 + size_t devMaxWG = 0; + clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(devMaxWG), &devMaxWG, NULL); + size_t devMaxItems[3] = {0, 0, 0}; + clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_ITEM_SIZES, sizeof(devMaxItems), &devMaxItems, NULL); + size_t kernelMaxWG = 0; + clGetKernelWorkGroupInfo(kernel, device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(kernelMaxWG), &kernelMaxWG, NULL); + + size_t lwsCandidate = 256; // 目标 LWS + size_t lws0 = lwsCandidate; + lws0 = std::min(lws0, devMaxWG); + lws0 = std::min(lws0, devMaxItems[0]); + lws0 = std::min(lws0, kernelMaxWG); + if (lws0 == 0) + lws0 = 1; // 兜底 + + size_t lws[1] = {lws0}; + size_t gws[1] = {((size_t)m + lws[0] - 1) / lws[0] * lws[0]}; // warmup for (int i = 0; i < NUM_WARMUP; ++i) @@ -362,7 +388,8 @@ double KernelTest(const string &kernelName, // cleanup clReleaseMemObject(bufA); - clReleaseMemObject(bufB); + clReleaseMemObject(bufQB); + clReleaseMemObject(bufSB); clReleaseMemObject(bufC); clReleaseKernel(kernel); clReleaseProgram(program); @@ -375,7 +402,7 @@ int main() // A:m*k B:n*k C:n*m // A * B^T = C^T int m = 122753; - int n = 1; // gemv, n值固定为1 + int n = 1; // gemv, n值固定为1 int k = 2304; float alpha = 0.5f; float beta = 0.0f; @@ -402,6 +429,11 @@ int main() vector BlockA(m * blocks); quantv1(A, BlockA, m, k, blocks); + // 量化激活 B -> qB(int8) + sB(half per-32) + vector qB(k); + vector sB(blocks); + quantB_q8_per32(B, qB, sB, k); + // CPU 计算 gemv(C_ref, A, B, m, n, k, alpha, beta); // gemv1(C_ref, BlockA, B, m, n, k, alpha, beta); @@ -426,8 +458,8 @@ int main() checkErr(err, "clCreateCommandQueueWithProperties"); // Kernel test - KernelTest("gemv_q8_base", kernelSource, context, device, queue, - BlockA, B, C_gpu, C_ref, m, n, k, alpha, beta); + KernelTest("gemv_q8_base", context, device, queue, + BlockA, qB, sB, C_gpu, C_ref, m, n, k, alpha, beta); // cleanup clReleaseCommandQueue(queue); clReleaseContext(context); diff --git a/kernel.cl b/kernel.cl new file mode 100644 index 0000000..96a6c50 --- /dev/null +++ b/kernel.cl @@ -0,0 +1,67 @@ +#define CL_TARGET_OPENCL_VERSION 200 +#pragma OPENCL EXTENSION cl_khr_fp16 : enable +typedef struct +{ + half d; + char qs[32]; +} BlockQ8_0; +// 激活值量化到低位：B 拆分为 qB(int8) + sB(half, 每32元素一块) +__kernel void gemv_q8_base(__global const BlockQ8_0 *A, + __global const char *qB, + __global const half *sB, + __global half *C, + int as, int ars, int acs, int bs, int brs, int bcs, + int cs, int crs, int ccs, + int M, int N, int K, float alpha, float beta) +{ + + int row_id = get_global_id(0); + // 允许 GWS 向上取整：越界线程直接返回 + if (row_id >= M) + return; + + // 本地缓存当前块的 qB[32]，供同一工作组复用 + __local char l_qB[32]; + + BlockQ8_0 valueA; + float sum = 0.0f; + + for (int i = 0; i < K / 32; i++) + { + // 组内协作：加载 qB 的第 i 个 32 元素块到本地内存 + int lid = get_local_id(0); + if (lid < 32) + { + l_qB[lid] = qB[(i * 32 + lid) * brs]; + } + barrier(CLK_LOCAL_MEM_FENCE); + + // 读取 A 的对应块 + valueA = *(A + row_id * ars + i * acs); + + // 使用 char4 向量分段做乘加，减少循环开销 + int acci = 0; + #pragma unroll + for (int t = 0; t < 8; ++t) + { + char4 aa = vload4(t, valueA.qs); + char4 bb = vload4(t, l_qB); + acci += (int)aa.s0 * (int)bb.s0; + acci += (int)aa.s1 * (int)bb.s1; + acci += (int)aa.s2 * (int)bb.s2; + acci += (int)aa.s3 * (int)bb.s3; + } + + // 缩放：权重块 d 与激活块 sB[i] + half sb = sB[i]; + sum += (float)acci * (float)valueA.d * (float)sb; + + barrier(CLK_LOCAL_MEM_FENCE); + } + + __global half *p = C + row_id * crs; + if (beta != 0) + *p = (half)(beta * (*p) + alpha * sum); + else + *p = (half)(alpha * sum); +} diff --git a/quant_gemv_test_origin b/quant_gemv_test_origin new file mode 160000 index 0000000..d043d9f --- /dev/null +++ b/quant_gemv_test_origin @@ -0,0 +1 @@ +Subproject commit d043d9fded6701759727e2f3c6e79aeb5a1f5b77 From 625fcdec6744c48a2aeebf39ba8b2d441c78a600 Mon Sep 17 00:00:00 2001 From: trm <956303669@qq.com> Date: Wed, 20 Aug 2025 21:48:16 +0800 Subject: [PATCH 2/3] =?UTF-8?q?=E4=BF=AE=E6=94=B9=E5=86=85=E5=AD=98?= =?UTF-8?q?=E5=B8=83=E5=B1=80,=E4=BD=BF=E7=94=A8image,=E9=80=9F=E5=BA=A6?= =?UTF-8?q?=E6=8F=90=E5=8D=87=E5=88=B00.6ms/3ms?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- .vscode/settings.json | 2 + gemv_quantv1 | Bin 57712 -> 67032 bytes gemv_quantv1.cpp | 118 ++++++++++++++++++++++++++++++++++++++++-- kernel.cl | 52 +++++++++++++++++++ 4 files changed, 169 insertions(+), 3 deletions(-) diff --git a/.vscode/settings.json b/.vscode/settings.json index e7850a6..d5ee39f 100644 --- a/.vscode/settings.json +++ b/.vscode/settings.json @@ -50,5 +50,7 @@ "stdexcept": "cpp", "streambuf": "cpp", "typeinfo": "cpp", + }, + } \ No newline at end of file diff --git a/gemv_quantv1 b/gemv_quantv1 index 5ab1c830ebbfa522c16512beeade77d5e7693492..856e43ac4f2f4ffc209db35edf1282c1f4351459 100755 GIT binary patch delta 23537 zcmaKU3w#q*_WvYJ+EAcLp{1>W1PCkzftID%Do6tjOfbdrQdFK6YDM0*h%S<18@59j zERMylid%Pmu&xhWKmjSFfV8lR1r_lH^#xNw_5ELN*V3DlE$crU?3eO5p{Ig_(P0BOo zJL0b~q+B+yOlZvbZSC%to)7-z1L4;>y@uU!u;*(n%G#)YeI}r!8Gp==5{EC-*3m4? zhsv8enuvEN{_asei@M8nuaa&_GHq3kL?V2q46)zi%xWRSr&C_3-kNNh%gS$Oc9z4AnbmeYtDBR4mE>(? zo!aIsv%Y|7_&3ua#w06g+bqh*(FvyE$|_qjllV?qA3NA|E0a6U^gqJNPh_5bt@La^ z)ASMPw*E6#5PB(ZnLGQ_(SbdMWTEjOIt!PR{R+aZ=v2amLmVj4@MMh0c%hq6a)bj8 zZ3LNtlYF%%kCN;fY143RY&kXjId}j)eqolT2y*P&6U7UZXo{Q~{ zLz;mSP5wCWUYJVR%Q>J;7U7Ck;cQmnn*>3dwZd}EaD~f{g- zz$cxXXA>PQ^w4R>w+Q`tc|X+02p?(+N>E@_%hM> zC4Va9+X=HY`lB}=MCZf^3pIKPD{mF3TI5hs>XYc*D4`s*Ao!_tjsX=09MhonsxshMN?NaK15Vojt^R_r$-p4v`nB1B z(?HU%?FL+5&4?E=;Q9b2yxxFgdgzy5XfQA^3tO+_23(s-jMr$uJG4p*f@;7!8t`TV z9%sN?7>@dr_jUtAILryc^{GUN&49;q5%%ja;8-_XuVe$>xfK-zrvXoBZGQ;#GcXbj z1=0+77XzMYz>^GkwgK1ge8d(Fc#0;^?B^L6Hy9L6HQ-$h_$&k7&43pg@a_hDu>rr) zfG{$~U3FyKQCc(MV{FyKy(``PpxW?=L)6u8-dry1}}1DmPvcPvqXP?0&@AeQd8?PCSP851`OLb%W5ja&Mst+DtEsRB*+{T>b% zR9vv%csblB`dL9Ju=hX%ar%v_KhiHp(9cKE8zSg+5%g*!?T5jYMga7(2zp@zT@XRf zh@kI`pvOnhV+^#PYJ8gkAbMB?Jt%_i8$sU~L3fFu<09zjHrmgc{OZ$xY~qU%^iL7= zsR;T+1br-mJ{&$G1RWGi(*qnAa{ z3nS=)2zo{YeP;wcK7t;j(U`8h#W2>NgYeSp#awt2ogLg4KP`i%(sg=S@h zz_JK>VJq#E&Z*_fNME9IPgbly`D!?<{`m@i;=%3-T*m4)QdFqN)kz&IVNf4G`Z$8*4QWh&MRXCdw$7 zepHm+6XhJ&ED?Haxxqh)@Ir5Hb zxRbWdqMT8Q_)J}R`5JM*n04I#IEBtkT$DcdNvG5?Bp57m*+lf!-CF%@>K=5{wXOTG z7)-2uidh-Y>UvIHjUqHNHJ$s^0F{KU5vXI+I@Y(mll*PH)IpR-e1RxWy&Kk5%XMwV zL*^Zl$sn0VXj8wy57|1aFpJW@roqtflP;<8OogYu&BS2p&C3MKHEC!v^{Rga)x0K; zdiWb?!kBB5FFK;G1V*|JAU%5wl9?pAok`YIpqlDyk)?`9kd;bqrIEz{Mn>eC5anx- zKf6y?w1|&Cqs@J4d#;^Xx`Xmi_b2XWKdA6YzpKkRX(Nptc-}$z5YMwoRfasxwI8s3 zqbhud6n4`{?l{9vs|{%kG}tdrA#>A(*1aVIw&U`;m!p*I8tc*gX|l}=>c+evT- z7d${^>jDnoP$VEr1Bzv(O6nK*L0w``WonQ(yFx&%M^Mb%CtYw~B(StaX&BMVTr*Yq zWki<$vN{A?blA@ZPwO$}BiaG(LUBJRR%JU?;1n7K`VP9eM6i;Y*1ertoOyb|?y$a_ zhZz%x^9oVkmIrqF(LKah@4XZbO93B{K*%1Liv`q^9`gC6-*ysvw*@~iPbJj1E*eD@ zh{8WAi4y4tebQyGRLi)k6R)YJBY%$DtYqHF83V{&u2S<=oow*V zZ?(6%S4g7`KGKMSk2In*-`rCoyt{K!cpDoS%Ux~y$-^r29qdts)M@b6Ct-R01sEp! zG@66K8di?xEApesmYTx}GN`ujd@|*CAy4zr#nMY^H_l5Vo`YehRa>TF&mJ208LA=G zssBZnsT+P(CS)i09jLncZ`Oy#HRlNKP`e zuoP)|eWLJb(bTt%)6_}Ji?}Pek@x7g2#H_X({ACBc_~jbgi=1 z-PJ6wRsQAf-rxEJgo+1D5lcth0zL-!_eCAV;D`bcYRb=O>N&7}!1~oXrP5Nj86I*a70(EI( zRzndMsB_^4`|jqmyVPpPQJV+C#io6E>THY>W#s64yL3X2=1AY?N+*1B$3QNv^wz7n zTQ-j#XfmI>qMXc`W0I5{k2kpocxa_61lVY$cvKYgH?iFUfY(@X?kLU&Q_-is8~eq$jmUFIuQV-}n1$0)I5-5n(y z5}+`2{HPWyQ^!iNxBf`ZI}-QJQ_hb4#ym2hd?@z$^YJs`@J<`6c19Eq9H^a&exO-( zT3rhsmSjLMs?Gq4?mu0o{K-2oc4`yNxH?rhtIY6r={c3u*twduh$BeNFK4mh4Ruj= z#~G#4n;v^g<%QOrQ5w8X@0ZY_7jkNamZMPV3YWzc%yC)Ba>`M()HkD{vUG1$AcVnq zoCv$8S@l5{9(_`(WR1HycDSbeyC&t)aoxLQkb|`8OFxPU?vv)Y9O}7zYw*6E@+5$-ufd9g`tXi{#w zeNZPS%6o$rmp+6l`YS7M@7qa5DX(EfnWQ0hJyw-7Z^( z`Q(h5Im#Wm{+Q@&%%0yaD4*miv0G2^p)h=ovUL1JbHYgF!1&RLy>V#S^gG>hQtL73 zac@=^rz$tzk>bD8D_x{bl1HA^V~Qw!GAT#ehgr9rZRW+`ZIaKGW7sa`xC+xl9`sdk zjKgF1HoEO^){Chv?%FF>Uw)G>OSM1#J>7DHx{kh1%d)c?SYIa>XHwA z(iz-x#(ftK`$ehNXKM0F<83*J0&=i8$dg?K@UZ|9L_v;p6&8LG^DnA3sDcMK)^$Um zEbxVWU4K%Vzb`lcJ4_u<`WGzZK*P9lqz}_C5+%x`!@1JuUa6^8rB+-sil%R$D7PS# zhD<(r;3MvCyj_>v2VJ#|Rw^}>2EbIrGTBPEyc?B~dAGz2FLDQqwh2l}UiXPH*sSWa z??N;78g){ww72$L?5eNBs9ovlk#74tLsz(iD{2JOVRzK|TFLeUfZ#|Ks!zmMwMHJc+uy+tP4fxi5!>B z7j%razcYNPD7CCQO-9>PCUBLe%kC)4(=Z&mV4s26rmp;!4~yv8RNWXFP2V66!l>-p%C7eg zi1wzwXD{3D2dn8dij)3O!D0#ob1)v%R~j*eC&Erah5IhAbjf|EdkV%^-b0wKlW`l9 zey7L-J!P0mJglSvU&f9EeDYT9Yg}c7#_s*{G36oD|2lEMpxU>#8d2Of_QW=!FV{q@R^!v*}sAYz$zuf zcc1*Cs}T}@Q6ATig&VDR8H?!rJ?lCW-QN*BDBKRouU&7|X zUiKv>E*g3YnhQ(*v87}uV=R&a5d*fsYhQQ*o9}x3HQ2j!+<-dz3+(ehvZwl7Y4+5~ zF6YHES2J{2?SXDoU7S7Or~&hWS~7&%pcX_s)Q3)>(}N={s0ud!to6`p|07dUnX4K# zqc+=wR1ErDHr9|Jyis?s2IRUv6zlwDt}5Vk@xu}+UOi5=lb>=S4xm_Jj>}1%Bg!2I zL%k?{O2eJ6hC^Z3JeD^p|_e+>RF_89Vlbv1T}{>IZjikKW_etYR@fUDRhW+qiHt)4N(V1 zJA_>7i+omhxa7m$18snH|D-P~k*m)cNR*th%CjI)CA{!L9U5^aKqQCqH`n0Hd0Yo(RH86F+|{gl2aB)Oz%E z>Pt36V*KE+;fI+*4wr`diReH&NW%!4z{-`ek*ZE+MRvJF@YK7Q%e!33AbglMG?vwg zG^5yoUE+q#L#R3CmJB8v-6|jA|CcB~fA#BBssn6#YtpSy#h; zYB0?4Km`-v$0I=xI#Vv>2!**9a7geXPi2ZO4Qm~pviPJ0)w#4TFG6gK$WD|l_~hbh zuk^k;@e8sWsN%ZlcmjumcKcu7#CS!{I27sr8 zP53|_imT;_cEsRg)xNBAMIGVwX10^NkM$e`Wm$z3Kjy&>RIY5~Hdzw%l*0=W5aqyj zBwi`R+(P@S0%lXNhEPD|NKlo)FjR$%%dE}VJ|u)Cpddp#N>b;cW)jmy2N%6aJQ!fj z9IRy34c3q@F_@;>Ki3U-r30i~A?f_wOflbIoQ8QB^z=huO{D?{m+UeKBZLziD+B?a`tYbMj&sGAMTWI;CDSmJ-m>RMBV9w`i6U zD#Z=#?#X8g4i3nFKN(zVENat={)LLlGP%t{9oidQQtKcJl)Pdn>u+Mb!5L(jGq5mV zEM8J8IfK-42Hi7>v4X7nbNdB7nIP)|`=JvPMyLzCujQn}QlJ6(y1;P_vA@|IxG3L= z9y%r+K>q|nQspl2#1tsWFR(30ejkwrCj zUG-!9{rVXrs34QvZPh&dk~7FL&cN{*bA2CYkV4KtTQSC~oIw@f47x2)W!G~C)rK?B zoy72~%Q=TU;v5R6smM&upsI65dyO%MGpKT$K~Wi%7|0m8{W#iDLz6g~#?d$pHF4CR z$zh5@NyXXYtSxA<;S9Pl65}XmRJStXHO3BN1Zy^dffl8+P#g!co;Dnf?zTXWdi7YV zDBaOX&Yui5C^rk`c3UBYLW_|HiW9&9!D>B|hv<-)za{Y*k?7BBsbWojyk zYYTJgG0b-Y=~a&=T$ zv{p9Xc+{0-_7IfK(-Qp$3_y?Iv0+1cf-+n zIES@iGrjT)G~06U(H69_q0=uje~q0UPwf}nMvE^d$W*r0c)kZAC&s0d;d2LRU%n54 zdLI^R${Ba23Bn#)ImOzh*wU{@+TTeSXKFx!=_n9PSc5y@q#<}$H(`}_vqK>IjTl7W zdC&VeZrV;o)t*P0L0Y-hm&mJNO$+7yf8&He4J7kCIj^olp5}N8FIY~(kD*^tkUDxd zjYoAj10E|8Wlsg|LTL+7!L|Us7{_pyi_-|2Uh>@ItbX@uA*z4o^nEmS>{z4_1$Z6T3?xg5{+qBb zV70eGDL?BR!2BBVqAr~W@khD%gjVrSA#}C@b!mI#SG)qj=gdK<^Q){uj7QP4ajrHA zIuOm(oKzOiO2|15Rq8)Zl4C|4sSx~-Sbt?KqmDSR99+jdFx8BY*tvKTNsnXF8>tI$ z8br@kT30$Yuy>YX{ljX&1~$&yZebp7Ba)8nF40y{wNifuJ=It=HTjcR3Rw}pY^qUI zOzysRK2}t^Bfukn2**I!3#prUr8kMeV)V>OwTfuew*js-t%hTNsHi#y} zC$3@JIlF||3$A6;C`IX|1F*4E$AhnL!M#`z_58H|Pk!++rCLnzx9z=Y-z`tyjJ-ES zDR7kb-t;tHl(sox$rsG%I*&DAh(K(;u_+hH1+3K^af*_h#*s#jAU>DtnU26_B(xdN zfdEoD8|fUs_=jhnm>_63_WVnfrekA%(O&jmdv@2#MKezOkXfTih-H!Q#wX*RPrGq@ z;23ri*o=3P$=8iMxcB%pRYLYr1rl0Us<`$%$-?Z&CGT6g0QrP6L3OTM5! zjA}^%Dn{G!^CTY(u$B+jIDq;jtW#;K^RTWRl_$;f(K5m5*ohaX(QN^D@&-slBkcri zS~l>7JUR+pcpS}-Q&qcnANB?zFl$@PVwE2o^Jci;Kz&qu%4`pKvAgo1-Pheh>rnph zbf`HPS=@Sk(o5Bl^Q0fe?4pCg*Ecbx|EvAUERQ!y+pd6=4!NcOO7e5amVR@^U3=Cn zTPn#Xi9Nsb)_!Y|0?j1YfU`}%8O8;8BU`qS5O&J4Xp!nukCZk>VeC-%N$k<0B=gr! z*rPmbCA`6rqMVREPPW99{vNfgn@{$oAg3*jF6-sJZIc!||VK?}q^g5b`db3`L@@Zyhq=FNZ!F zHdE2VVm^%fQ)6`WDg%HToTIrU-AAcI^Lg;G>v24mz>RDt=gg-qvEk;7V4!Q-WK z>JYAlR!tsNz;C7{g)_2S89ZJ(rwYUf)`(!xvOu8%u~4gOKrEcV`apU7e-!Y#z$TVw z*~Qw((gO(Bp;l#3FbiDLN0A2$IaJFbW&OR0elmuDf)uGeSrKN8#R^&s@lqZN<3*LJ zNVHG_{Y@I%C@Qaq8PxD%;Gk&Fsi%>OH@f5^6J~-;dPE*+=&^+#(~>=uWv}O8Mp0;y zD>ZA;#ZN4>SVQf?LJNJVeE~g!mP#Y9eejyX3p^NU3olgSjH3EPMi)-=Sa+VXng1u4X0oXJZ6U<GN zq)HDh$SEFLoU21IIk1`R)G-rgh|u*lsVm zu@{j{sA&mxE}a4QL!5yvRMQ&FGwpGI>F@3BtG9yI*7J0Fryz$O9WCmXX^-!A3m*MA zyji~&|D@eH(1-%;jN}Ag zqT5uC^inqz_SSxFzA(zZ(qbzZWj|(#4YI}N;$D-z?AUK~tGnu!J`PrrC~h;ajFJ=n zfwdENkPeQV@(x+Z`1}r2PPs43M0tmqvS$0+IgY6EoUEv(C!vH#$2GK=r5{8T^%C^} zg-vWVW@7`vRG&Q_3Z`;{b_%Ax#kyPtF>Smcm|9Lf7OdF}Er4VnL>8<_D?s7PzW^f-$D4)vn#2co-Gp1#DpAQCg85Rd1kt>Q(h6#=2gF6=WO=FZ=`e z>vi_PuQWd+Lo7sCutrUeD2wjXh_dYUpd?QjE@XXrYkU+e4yssc9ov!Acr*P4D3c(OUD& zwrFj*lblZ6i2r%w(OPe`50BQotd1l08l-!2=|!#5zaEI8XiYuw9*dYlV5k?cU&a2_ z5WdlyL|2G){2CU$Npyi&HP^771?vp4{&5Yf1o2=8_@ds0p7rtlEH;pSJw6lB3O<=F z6)fKJNtN3P!aG*W*jU&`q}DO?R`2v`bp#sT7^Bg#eUPcZh#-aPtyEO*zJEqrjMmZ} z&4-Yzk!kf$X^qhu2taUQL#)|?=9 zS!b$#o%IdC*lGy7>1PMKv&#cFE>_a#ru8^oDhP{dQAe>wA@spk8vF8yzu+wh^~4^f zWNshx{nM0J=DN*WL&}-C-8%K12G06Dkdmh37HHkDNJ%c}-fb-2EVK`{=1@|42!<}$ zr&(_SrFQ>7nNZN*oW4kTtl<83RkZU9T5A_7KNkFHL=pBfP3fo&VhTL*ZMuhhGNG1N zU>vx3H}OdAm`88(E)^n1*&z_tYqVzrPRO`1Q$)9PFuk z=?}?TXkZi1_P5~B0b zHur%Xgg5m3-`4tqL-7Qx`N{mDAKvWB#wIE>qv+^a7di>O%1oT1nRHU^jONC3SBj;B zWM?#xI=BFvmcAN8WvT~MQ}AKqEnfB#KgcQKbgrznmd zTjIqcyv`uab*TT`rL-^nbIbuWwwk(2@f7wlPpMHJC>&(&!;@ckG3q@^2Ag&%2MeDy ztM4iW58UXi-%0DiRiAX8UY)xECcN~@lrI%fSl#)q^5O$C&EB1gb)L)I`(0(gykSXa z6Jh7&9pv+5dIG=Pgh%_wb|}l|jWe&=q3oY0nZrAjJ046ltJ{@=2S=N?>`?yw;J}!o z1e6=RT{-z+7bRoCezUS&`F+9k?40e;B##Jm23LKXh6-;;n!k9;p7C)8_1kZF^95A* zW1x`sIlz z=xOCcqtci(|Ff z9hY1hl@sS&5F3{a4LNZ(5l2X*;0pd4AWYA)gmG~W5s3p$_CBy~O3$ZVO6`Uo(bIrb zKcSr3FeLgZWW@^Q#;5OgRwfBT*RCC6Y&1x+QIP5`q7Oep8(Z?N%F|D0%_u8D69cz+ z&~9)8` zraLYTZ2?0t$XmdTzGI(wp}FxQ=Ef7G*N+;5|Ag72J7)2cPLy9A?IxZ6adFkQsE?wc zH0K(nqe!U}e`}$%97FRC=(9Y89~3H$SslIBDqGe?Jz=I1LN9#n+87QWMIp<>knzM} zl#T8fJ?4q1br#t=erMt z#}YFIdCOx^)*;#fgXXo-vUQ#1i70ABCu(@pzi^`45i+w42EE)Mv)2fFGob!uZtoJA zgJE+mH%Bl23WZRV2V)PAF^cvgjM3Pck0r%{CF4QrYg|<05@0KE2Y4KFqO+3uY+CgD zKtkJs%e%_?XGOoqG<=WAddeo8FJq;?PKaTXQUv!ZJu^CN_Zy9T4NDDh_ddp2x}~pzl&cg=r~-I*l#zwh{l<&7la=T30#KxZqUUpiX>tWN;&MgnCICO84FN~HV? ziQ|wj?~-{SNjTYMH1c>q!Eys+kxu^2WPPKnaM*m`Ke`H^$9x>Up{uaXe%ZRStMFPE zhLW$yniA^9`3`h0mg_J7R^aQE!<)iWSaBTVIbA>SCiFs zd*MynRTBu(_!uj(vQ{NS{zv{X40h z&b_Ft|96ieD@sU05U_b-m;UuLY{=PV(^Bjo*;Q4qeDm)?J@K6+c45G=D;Z;$F3kfW zq@=$#IDYX59L1Ps*KuXZYdv~85&h7m-%uJg-MXo4<*nC7;_GBTzt+PPQsQ5~-<@d4NkRC(w#Sif$Ut zahmLOd9c!$@S!j96sPq@DA4d(8eWe59$f(quRh9YF7FqFmo-L6W9S9m15So3HN!NK z>G}}3Mff?)skYzGHF?MAHjHv#YIt@8evWWI4Co(u(u-NqR8+GH%^`i<#WOsKe9>11B&qHdTHSkU8ebrsK7@y83UcWmf+oL1bGZQBL+PqQ#7}tI2u{{6A<9@a5_bmI3W}n~VA3S`N@+mUbo_KRtL0=hoNkMpqN$%D@*9XZA`X(BAIco)% znXL`Yq1ZZH?;CdufOjJAyRRh)*wDZCg=(;ClP+QXP|u|_AxBoyBHF^N)NW5Q4WQqY z+dDwBIbP6LVNIPdfyq3`e*elOUxeiKTDFF`=^tv`>_@nUziNt1Xjdu~`E|Z-%3C{< zOe5H@Ous3wXp#+@2EFG#CazZnms&>uJXxNmKN_QljlWSG=&Rv3<7EVr*FRgP%P&&g zJCjVS*zX4R`!4%EOutqecKp1ZU6h00q48W{r8a&S=@PkXkRuEF2hg0F<b;EB6@iNa2kz~^Y#(4 zV)}%Btya7P!s~l4axb99jZbN6ojl#DOQgIOFAj>H8dOWJ6nVd*}bMMHGEb?8?V=N>i+ARHv*@z zAspgJiN^St*}6$N0*m-`D1M)$-|k;0^46;u%LPBvP^GYRy&Zxs1)ok+%(x zaq^_qcGo6pkPyhPJB7NIHEHe1hStfCxrIG1rDdew}sXEc1)1WvQ)4B4;rlGdww zE@ccVdcEFStMuH1PmqG^WaE1e(WYYzO-<2!=`Kav! zp-Wf;$T1t%r}SLz8(OKz6^S-`*ZZ*PMc~wD`u7Q)%nLyngrUS*Eh7&(lF+|Yh>Edm zsfO!cDlF0P8ogxf=g6_1dPEtuXG_eaqM>sa&Y7<)Ig+S6n9xl*yXVh#U;2>wi)PQ6 zJa>7~mm_cSBKC<=6ezOfPx|vrWmo9NuHKm* z&&1m&<;?WXo;2G#+vAy(J9FZs+h-!#aw=rDDxD7alI|Hlsc6`;x%owlmU{CcmhbW8 zqg(6clV(bD(2X^33a*7~)yBWy3)^s<)Aumo+ng{>-_{ z<}NIPCzI0QKl!qy=Ld((%1!&*DWSwaIYy^@Q9+L<@AiCD4pqG^(WCTlxGVV&r-j*f zh!ciiT`SFCp}Qm)2?HciC+zhT~We?SClAgzPh@V%K8&Ml(Z9$P8Dx+c|9$PQ6`=k6mQwhVLipilux9!tJh@fl|v_P z)+vN*u_teKXgvLEq}9rVlY?XRO)sLWH7AF)25b7Bw?a8`GOdGlYhcSc^Vr$gvz-Py w830zdw{=avaff1D%+yGG-6%1ThdZJL8jDlTqG&C(Lcgv;4G{H2Np@(DM#xd-q z&?|k|tDY??MjH8o21~#t!&b9Ws%KuFrntOB^EvZ_7f>%UELH;+&WJr(RKEg~^fc7aLRnuNzt-GLoN z;!VQe{p{P2`^2#<%^W4Z!H$GRvFYY$@ln>>5+!zF>w%WB@n1xV5vn3xVuSu8rQKLaz;s7t>QEnxh%Y8vUvAp(@xQ~~!f$eM`@7qKO>d&f@HqHe5C)-X2^Zc{fg%kLN56~|dI&{3RlrVt9M2_-DqpV2Lz6#gRt?vB zm_x(Yg>yS@VY+4ra{SpDqAC<=h8!9$YXwvVIdl>Oj2!iA1r%xW2Y~m%fU20I0@{!f z&T19T;uRh!2-IKcc|_D!AaGOTU?j*>!6YkXLFLpeN;D-s5GC*$LVmPM5EORSa}4fh6OPc5tHz;R)9Vq>U~BSNu>xv znbrq%-0u{4`GU1V70?+vZ4bgL7pihuZ>uSF?)&jaPSBHK!1F=tc;gPvMx)&5PlO!Z zp_b1XeN-9#ZZatJH{gl^A8o)}47kgHw;J%V20U7;n8@4Tz$n*3ga+law9e9p58+f} z{ZWUBfEE>m2XzpsK1v9mY{2z-gz)kB)9R1hTdEW>Ca78d*2g*F_Ze_~A|ZUD0oTVr z;nW2Blcm*|@F_Y-b(I}CV<0q!0iV7F9tl?fDbm{4g;R9jN! z20X)nXBqHJ11=fx+YERfaNhr@#g6JVZbaD{gFNp3HM0dMEyLreXZUCGP4Wh33!-}k(cM(q-8SqZf&@Z? z=k zF8Q2N#-140l|7vq?oPZO2uR-QqgMj~_{VJfQ%R_Pfb4e$_GDsJs!Tm1x13~-O^?4gaNT>M%60@2oat0jCoXkULGMDweR_jdy_&2zlsgPH zZ(c?o`Ya7kP<^XQc*J;2lus$|t#+to?NH02Xek#tXB0VFHJ@DAM3HvMKa;OK=-%qN zs#?CPRzYf=JE%VL_V-^w;bQ}A#=}wW;m5HRQe5ED5Y~K@m z$2}kZD(Tv}<*w}tc&qJ&Q!YgR7}Uw)>36tY`Qzrt(CDzwd)@csAvL!fzS+s0yTMg9 zBB#uNX7h8{s3)G`~x(=C)tO$k~INAE5>^yyOxf7#~g0m3*6~qjojVoDk7o$0Q6K*~8s@ z8CMuC-!=hUFt~urjjF5yKHGNA|51svInoyCNu=O>hTurP5`v)b07N2iQj#~Mg0jU4 zY5Q%Ff>iA9V~fl>J6o2nOPP&@|A90afH$P%tfCE8t{nJ){N1BERfjQ56@h3E9@Jv1 z<^c)80|}bSxGU`KJNlS*m9fw6$aG&&d`Y~HS>SPFd>o#VWLWd?W*Qbj!T8A~Uq=Hu zxeR>X@2sLIeLzKDwAn42v8h*T&17p=37km#{3GHk$*qBa>{$!MS#sMghtwamwaKYe1$hnT0mI?vtG{3mFqkKo2L-svu#gUp2zJw z=8rO{x#TO9)5iBw1pYdsP5oYxa8l^Y%Z^lf>fVVHFL(Ihh1I|1s=t1?Z7Lxf&7?D z;eSM&OQCX33hbg@bC2>Y&hY3bm^wd6RR*9sG?c(mP9X1+FDs@0rKUWgxTzCrL*@%L z&s(A_B_r z%gn#9i#rZSG38YXiE+9uJc1cLjvD$|&YDCEP8Yu56z&G=zX%izJp!gaU%ZZ?|7ZCv z8jjX$9%snk#q+~3iD9ytWKcHbF$So&YW=ghFQjHC+u~MH3k$PLSkuVZ@iPDhx-!gT zmu~PR-|+o3-TEvv!AY5|1zjbtWhLctIA|AWvc(mVjgPey@-K1u^cUtc*Qi+2*!gVU zsGj}zVQ~=(2TYQRhtmpzesi**gXA6F55mg+3yAqcD80_IX#uMkm5>ycj3cjKDn4U% z%B>I^N=t=&zJ~76ln2lo)sa&2Uv_=eg$@)gPoPlzI8Qmk9?O~9?NvlITRu5fKJM~$ z?3zNO!m=)r3OkoQWyL76l$`+B}hN$pQ&vhv1Q)*vk~wI;IL za-AJNK$n6tZ2YR6WNUKe@CVOPJKMY7<6^OQ{bYJ#4h!6s@R!yz*a=uELca9B;68)8 z^NVMxd!JI?0S^m4AQY;s1d6tZn#1lJGbnt?X=>hj#qld!H>O+fC1l2?n)wjdC1j=x z6ug7uDf{giwr@;Yc&Mhd{mct4NA5LHx=x;^Q}$~hP`o_eEPAu!%~WvA=WwhOaf|nc zc>L%XRw8UUEy`n%g+rFvY@Taq_~ien4w=uew_H8DO`@W_1@YFj&oKgB@?GOwJ0mWQeyDZM{4yMwsYKD@sVgUXYE;w%Qt*AszU91 zeKerrI&qB8JjCToKbOGVce}%e;C5e`cZpf@SooKU+7k}tuvhaYn6i7b-||LvU4(nR z=HF>smDC`~=dipe{unlHe2iQE%z3YKQjWaOIWbqhG)a=by2t6ukHc_TBFW3+tu9|d z{A_d>*%dz#Yk*$2l#3r<_^QKxN3L+$StY1VmV8tx0T0hlw}^Q&!LDss8A_DQq8!x+d(`%Z$aB>Bv-+45(!ZOo3( za{6)th(mt78i%}WL8w&_nL%}^`&ju{uH0Pg7ZH*nxEPdbn=GW2PfeMRTI8{3CiJrK z8H)XV!fj#83k9cl@NU*Lq347-*h=byiJiw1STEPgduz{yKl5WCL=cKkjd0rDFIWME zVKw4mXUMSkXWQnVvL&B3bqkyXJ3c!J7xyaE-&G4c~fQMW;IHyf0V2?7SDd6FT3xos*ms zX*+1OP)u?#D62EDWMCJzgeM+>5^qklOZKFyh@>J8m+Z5uNDJR^M|uiCUbIE}79chB z&zuHKAXBEG0C^W3D^PYS$_AC?<_5ScSbA)$W_Bd0`ds-C7DVOPPbB4O;X)tdY-O`O zTZB0UInm~6!7Sjk?e4Z_1Gnbgkj3RZbg;uNm5w@fW)sFaKi8-mRes))H}QkRwI5M` zmp6U*eOQ+G_7js}KDJG>c4novMh( zwZx+|b>H`I@+Q3%cCk5bq>-;GV#1V)-F1mXQfs+?21`AWx#Y8ogY>*SGJG1P9nLk1 zhj#-%%!NFl(9BIi+tVHu_)Xc&OR|ih8c67%*aaTNUz4p{&-U=!)-9m}B}tp?h?6bp zbvlo^f8(2ai_J4P0#))6_q|54A-RCKH zLlB_d10O|z++Q@wZOSjiUr6Fc^o>U8z)O)O? zz-{tFB5W$EY2w#$VPUMtjpnnVA0t%cx(nM5uWv5pGW{;wC_mcQWy{`6Va`ozP)1Ps zKFb>XQiUN*LD(LKLHo>aExN2e1*{< zdHGTNZ!oX7i558hIPS9ypFm9B9Lo9QWY^LM^3_?wg$DGeHpWE_KzgK2* z@#+lf?f0rRaofn8ui6R>N?gnZPqf>Mo!qiI&a}%4?GgjUzBitG!iVOKp{GotQcy># zv)9At1T74ydKj*v-M~XjxZ%IA0%6pSx*#`frKrjto_KjXf+88TaS@VSnffCi2`?a? z-f9AY%o6y+K`l({Kte@|bM_j>rW>)k+h;XWCY zys65TAM_G(<$YxR8;wjJ2df-1up<{|O!N~IFmi3hv7x51ybT2zrKl(SOx%1Tl=&BvoF^J5+_OY&Nk zr;J*DO>2&-C4j_+<4?4lMy&>}<+gb?qgv~pQ%0!_Dp+M;!!w3DuhL)TkX@C7)t@sWso1fJ-kc1G z)qB=L0DBFg-t(fClMly))q=dM1XK;r`d&d69+Rr_9NHXt(&K_PSev;k3snZWpfW-=M!w1*dwa|@r`8}W zTxEpHAzxJ$I)zYu`>70aTxFo;IHRM=Acs^2ZMI3{()Zl0@-_yR9nSbpWl-Hn!!6X% zzMoX~a?aR9)ezQpj=rO!sVa(fhoc)+G($z{ltUJZRNJ((P#H8&6Jw6bSkT6xeGf6l zxryVgUIGrHcoYflO?aDPS|$>LLX8=Q6n2b6efa=snIXAoFg5KO3uK^^cTLs7`jUmL zG6h-~)*nzP_f()<6pN*r3UuL;YyeVdHEAU&Pw+x`$ta5~{OciTx_rA@e`hPFMVp#_ zXJym+x?4$4ZjmmZmy*85IHeg({S}Um%orEbAuv`xpSJ%If!`Won}GAcBf=wsa^QRD z7JuTOAlDqg+t5lywmBRFO zQK~P>0L-a+o+P#U){tbHzqUE#qV%@59k9h^VMBJt8sM9QKl4v>UFr7+C#|%_F6L8; z)r|>YE^3&YSzlP8-F$zfKtP#xv~9jb#k!G>&!bSZ>bVAahwb9&VK;6#s5LK0w*LML zwyBnsp@MM1maolRw_m`9gbGsb`-Xz{vDKCNd%;R+l=pV+Mvg*I^H*F;DffN{3686C zLNP$}sTR_{jq}va=>YD|(&yM+^Sd+u42zq#Ws!7m<*m+waPwwN8Tw@kVx`S>beVVR zg3Z2-+|e2p4i8$09}6B4TwW{oaVt+Fd}IS5Y2;$c9`3iXv9mCZfiH2f>Xxf+_;>Ob zbRX;kKmG_zc-?hy7bJ3R0zcRI(sy+RS=$_5{No7Q`(5r58zHhDBHk{~VYPa;EB-BC zmuIwf3Fn;yk{4%)2k2&`iVP|zzv9)aLYc}u^2l55r~DM;-5?r>7_TzOaU}`PBiL7%1tTHkpW-NzIkrU zw6svbx%=P=T^*(HFlC2tmfS6UU)0yX-1kA~#C$QLci5DJ#s zc4v*U)wGbIZeNf}y;Acf8gtZj3hqsv>og!VX8Kt;p!DZE=$rV5NX3_aoY?LE$lj*1 z8;N~}Hl8SI73!7i4MR_PnxDJ#QmgotR^|cQsyP9S1$E601U_%h2Pt{oHV_S8dTG4A6h+lSbkJfV{ z9p`<6HqGP*$o+J+mIDZBg0lmAxZplhMIoyx*!cKXzO%GM7HUjvJ}}91g*4>JRszyy z`teArAB7$)2&Z>yh5OQ-GiW)yto(?^rk>E5TBY7{2jG#AtfAKjm_^8$K{HrwVQSn{ z1^8qGQyrgqW7v;RrNmWDy-|V2!>B+n-W7*u zqjxDI4zrXc2`qc*z}SJGk;lidb@t(ER>sw+T+N5Iean-gNO0YrlL*)B~uwviKm}z$KU!`x6oO8E+h9`OEJzJ#dh{wp=!S zc95m4=xVAuz(%haWhy?%yekHU{fyD01opF!R&;aYA}d$^5?5URy+@LNryYsp9f1?A zT}o<|GFzUUDS5|OCHa7qbXCf{u;?$iI>NQN=X0B9A2qxDnI!*5`Ds6eqE-@{B^lEA zH&W^?8*bRLA>{e3aFLYxSpjbAm)*=Q!VJX)Y-?WP=J31K0sKD63z6}wWxdNcFi|Nc z``+6n?;^WPeiZJ6AhAoHK*5Vtn3Qx~EwA|zDBsx5=w_5rT%YPuFNx>iip}#Rx=8aq zf!ldUQTnzCG)MzNq@;beRf8eG4K&Bl*zgVi2%B_ePs5D@*x?liO?~&Xmsd{7`Vu%~ zR?x<4Nj({*CZ|5z>niRX(!WfgVBrVFR!}{^28A@?xS9_7C809IJy(3U#3QnA*C(_1 z7Y4FtU&v&6FJv(P3%9cVER*?7*2-Dz*WPY1Hsp5t0*U^K>VlU8fm9lQ<$(bGTLwGQ z1xQO?3k0I!OBK>o9Cr%-i4?u11?eTw3*G@Aa?`g30;Ksf?Zv?maF(So!(5coIJJoMYSD4+?$`8_&kVd9|Q z%>`! z{?K8j9Dg68lvF~T5q4*Ub!-HLN&Z9pHG=mO@$eAVh`)CbRXTwQcSS_lC8Ynx-&TM< z*r2L-_UZd^q10H*x3dfH-x@j>+4vo-f7ShtEhsTIwnLb;1Xvc7s6O$C`4FG=9cz`(Vh(CCJ7j`n>_uL(Z)QmKB7R z1DkBXC^+u{Thhpmeeh)HK$Kj5giYAm5Rdq%n^L#PF$3yahc_$rK!u=`9S zBcj)aWJfqwo3kSli!FCXq?&TVBNCkv4rfI4$O!w$2Mol7wWjQd=+$Jum<{;Xpj*qp|Mz==z-2f~squlOig_kO~t!@_tu$HkLVo8f7s2e-Yic0g*8?6p|BRUu_!IF?8hfA^Gdf zKFiwB;;@iFyJW`$$%5&*WXDs@ayi|05jD6~aSKJO;3-o^`7p9a@EK`~fzzddN6x%4}5c z5{7XBQu8R?y|Eobv!4%HXC6oH-3|Yan37ZT;ZK3?j{~~Yv4=tT3FOVA+1YBz{h}y7 zCtfmHzHtbzgjmiwgttS6{_YU|Zn3QGEmVhE%6bcjLM^ZL7CsEKRP+|ESS=s+7Wap* z4*$p@?CbiL`Dcf4tUEk>Io1NTcVbECt5^%T#XT%wmiGX&nrDypZ~!jpNpMBa@yKuK z$m?guE7mz{?3@Y%)|=KJ<=RI3rrFn}vUxoycD?TUJ?gF#)$mx939a zLCa>VP$XLLu~eyOS#1^G5iN(U!Zwp-uT@wVVmW3Nc7%K>?zRf6LM?Aug;zrVLqxz5 z5X}!;D$nC-TJG_UxlUq+ffEi=;eDVmX1n0zp<)Atl_ zJo?89C0p5^gK^?kcI;rXN$J4C5A_oN$&wDGL@wE;N}AL}@UwY`;(E8@*G78u`KKKJ z@QEI)nh5#q%|j!^Jl1?DPAp;VKAmHd-ex;Z-7437nq^{_`(}je-(*!TWStt{aND)2 zrfLT0-;bLx6J%-lPa6IY9oPOvI!!P19Ks*T&((11-t?3sP9i>FlnSI zSc6kCJxv;;oQ@E9bOjLx$&g>eb$JJHDmZ?aDz61r=nvd1oDZnzjd!{xpXa=R@qafz zRAUs~#K6p~;r<|eI&gBRL~{s!^XD;+N40$&rx(LCJA4X zUvn;)uX6Qy=zxYd%~f&Q2+(the-s-*pTEW=VMd`ZMDMSPG=e{wwV?2`2{N0m29u7}MPmeUk;^Rqn+gJU> z8T^alW%N#a2`qA?=n>W23@z;60ROX(aqaZ&a!d+N6TM+j|2PrC{29Z)Y<`lW+zYdZ!jHPMi6EJ zk3_c%{@QguaI%&n-Gg6E8G7V3?TgCZq7TFrEdxF1jVUa&XS?<7{vV%`c4-hRbYh1Jqe>x&VTKeb6MH-&Z)*p`wF9er-$eYXVYN})-eoSJoUO7}* z`tx-&Yf`$f)XQD)|E)OEu`hdH>CdhzgDP8^JA_nr{dH|fWnD{yiG6u-BP+htpZ)!6 z=gO0p_JmZPY`tj4`Pou=^ncSuv5{p4QX&-$KYH}^pKK}663sWobJt zJ$-pd=^jx`4%esY@1H2W0G-GMss(+CtB5bX%_Jsw(zgqGdQB*On&TU^616lxv~&;f zNb4&quBSswub9Ntc3V}nWNYcIA>vS-EZbW8BuKl))wcp$OFsw^Q^NHfiS%~qg%EK_ z+ZI>fk*q1rG>fSnv}uD+0jxH>bLkqh*t?yEIx0%vHH#zkUF-msb}*{cZV`KT(bcr+ m6~Chiv86c{aX@(9N2;^0MwdQq5r>3d_(r8OMRxUnz5Wk`Hd_?{ diff --git a/gemv_quantv1.cpp b/gemv_quantv1.cpp index 05c1c3c..01bda7e 100644 --- a/gemv_quantv1.cpp +++ b/gemv_quantv1.cpp @@ -383,7 +383,7 @@ double KernelTest(const string &kernelName, half max_abs = computeAbsoluteError(C_ref, C_gpu); - printf("Kernel: %s | avg_time = %.3f ms | max_abs_err = %f", + printf("Kernel: %s | avg_time = %.3f ms | max_abs_err = %f\n", kernelName.c_str(), avg_ms, (float)max_abs); // cleanup @@ -397,6 +397,111 @@ double KernelTest(const string &kernelName, return avg_ms; } +// 将 Aq(int8, M*K) 打包成 RGBA 像素，宽=K/4，高=M +static void packAqToRGBA(const vector& Aq, int m, int k, vector& rgba) +{ + rgba.resize((size_t)m * (size_t)(k/4) * 4); + size_t idx = 0; + for (int r = 0; r < m; ++r) + { + const char* row = Aq.data() + (size_t)r * (size_t)k; + for (int x = 0; x < k; x += 4) + { + rgba[idx++] = row[x+0]; + rgba[idx++] = row[x+1]; + rgba[idx++] = row[x+2]; + rgba[idx++] = row[x+3]; + } + } +} + +double KernelTestImage(const string &kernelName, + cl_context context, + cl_device_id device, + cl_command_queue queue, + const vector& Aq, + const vector& Ad, + const vector& qB, + const vector& sB, + vector& C_gpu, + const vector& C_ref, + int m, int k, float alpha, float beta) +{ + auto src = loadTextFile(kernelPath); + const char *kernelSource = src.c_str(); + const size_t src_len = src.size(); + cl_int err; + cl_program program = clCreateProgramWithSource(context, 1, &kernelSource, &src_len, &err); + checkErr(err, "clCreateProgramWithSource"); + const char *buildOptions = "-cl-std=CL2.0"; + err = clBuildProgram(program, 1, &device, buildOptions, NULL, NULL); + if (err != CL_SUCCESS) { + size_t log_size = 0; clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size); + string log(log_size, '\0'); clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, log_size, &log[0], NULL); + cerr << "Build failed:\n" << log << endl; exit(1); + } + cl_kernel kernel = clCreateKernel(program, kernelName.c_str(), &err); + checkErr(err, "clCreateKernel img"); + + // 创建 image2D — 使用 CL_R/CL_RGBA 与 INT 格式 + int width = k / 4, height = m; + vector rgba; packAqToRGBA(Aq, m, k, rgba); + cl_image_format fmt; fmt.image_channel_order = CL_RGBA; fmt.image_channel_data_type = CL_SIGNED_INT8; + cl_image_desc desc = {}; + desc.image_type = CL_MEM_OBJECT_IMAGE2D; + desc.image_width = width; + desc.image_height = height; + desc.image_row_pitch = 0; // let runtime choose + cl_mem imgA = clCreateImage(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, &fmt, &desc, (void*)rgba.data(), &err); + checkErr(err, "clCreateImage A"); + + size_t sizeAd = (size_t)m * (size_t)(k/Block_size) * sizeof(half); + size_t sizeQB = (size_t)k * sizeof(int8_t); + size_t sizeSB = (size_t)(k/Block_size) * sizeof(half); + size_t sizeC = (size_t)m * sizeof(half); + cl_mem bufAd = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeAd, (void*)Ad.data(), &err); + checkErr(err, "clCreateBuffer Ad"); + cl_mem bufQB = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeQB, (void*)qB.data(), &err); + checkErr(err, "clCreateBuffer qB"); + cl_mem bufSB = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeSB, (void*)sB.data(), &err); + checkErr(err, "clCreateBuffer sB"); + cl_mem bufC = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeC, (void*)C_gpu.data(), &err); + checkErr(err, "clCreateBuffer C"); + + int arg = 0; + clSetKernelArg(kernel, arg++, sizeof(cl_mem), &imgA); + clSetKernelArg(kernel, arg++, sizeof(cl_mem), &bufAd); + clSetKernelArg(kernel, arg++, sizeof(cl_mem), &bufQB); + clSetKernelArg(kernel, arg++, sizeof(cl_mem), &bufSB); + clSetKernelArg(kernel, arg++, sizeof(cl_mem), &bufC); + clSetKernelArg(kernel, arg++, sizeof(int), &m); + clSetKernelArg(kernel, arg++, sizeof(int), &k); + clSetKernelArg(kernel, arg++, sizeof(float), &alpha); + clSetKernelArg(kernel, arg++, sizeof(float), &beta); + + size_t devMaxWG = 0, kernelMaxWG = 0; size_t devMaxItems[3] = {0,0,0}; + clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(devMaxWG), &devMaxWG, NULL); + clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_ITEM_SIZES, sizeof(devMaxItems), &devMaxItems, NULL); + clGetKernelWorkGroupInfo(kernel, device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(kernelMaxWG), &kernelMaxWG, NULL); + size_t lws0 = std::min({(size_t)256, devMaxWG, devMaxItems[0], kernelMaxWG}); if (lws0==0) lws0=1; + size_t lws[1] = {lws0}; size_t gws[1] = {((size_t)m + lws[0]-1)/lws[0]*lws[0]}; + + for (int i=0;i evs(NUM_ITER); + for (int i=0;i C_gpu(n * m); vector C_ref(n * m); - // 量化 A -> qchars + scales + // 量化 A -> qchars + scales（v1:AoS） int blocks = k / Block_size; vector BlockA(m * blocks); quantv1(A, BlockA, m, k, blocks); + // 量化 A -> SoA（v2） + vector Aq; vector Ad; + quantv2(A, Aq, Ad, m, k); // 量化激活 B -> qB(int8) + sB(half per-32) vector qB(k); @@ -457,9 +565,13 @@ int main() cl_command_queue queue = clCreateCommandQueueWithProperties(context, device, props, &err); checkErr(err, "clCreateCommandQueueWithProperties"); - // Kernel test + // Kernel test: AoS buffer KernelTest("gemv_q8_base", context, device, queue, BlockA, qB, sB, C_gpu, C_ref, m, n, k, alpha, beta); + // Kernel test: image2D weights (SoA) + std::fill(C_gpu.begin(), C_gpu.end(), half(0)); + KernelTestImage("gemv_q8_soa_img", context, device, queue, + Aq, Ad, qB, sB, C_gpu, C_ref, m, k, alpha, beta); // cleanup clReleaseCommandQueue(queue); clReleaseContext(context); diff --git a/kernel.cl b/kernel.cl index 96a6c50..3c590d7 100644 --- a/kernel.cl +++ b/kernel.cl @@ -65,3 +65,55 @@ __kernel void gemv_q8_base(__global const BlockQ8_0 *A, else *p = (half)(alpha * sum); } + +// 使用 image2D 存储权重的 SoA 版：A 的 int8 权重放入 RGBA 像素（每像素4个有符号int8） +__kernel void gemv_q8_soa_img(read_only image2d_t AqImg, + __global const half *Ad, + __global const char *qB, + __global const half *sB, + __global half *C, + int M, int K, float alpha, float beta) +{ + int row = get_global_id(0); + if (row >= M) return; + + int blocks = K >> 5; // K/32 + int baseX = 0; // 每个块占 8 个像素（8*4 = 32） + __local char l_qB[32]; + float sum = 0.0f; + + for (int bi = 0; bi < blocks; ++bi) + { + int lid = get_local_id(0); + if (lid < 32) + l_qB[lid] = qB[bi * 32 + lid]; + barrier(CLK_LOCAL_MEM_FENCE); + + int x0 = baseX + bi * 8; + int acci = 0; + #pragma unroll + for (int t = 0; t < 8; ++t) + { + int2 coord = (int2)(x0 + t, row); + int4 px = read_imagei(AqImg, coord); + char qb0 = l_qB[4 * t + 0]; + char qb1 = l_qB[4 * t + 1]; + char qb2 = l_qB[4 * t + 2]; + char qb3 = l_qB[4 * t + 3]; + acci += px.x * (int)qb0; + acci += px.y * (int)qb1; + acci += px.z * (int)qb2; + acci += px.w * (int)qb3; + } + + float d = (float)Ad[row * blocks + bi]; + float sb = (float)sB[bi]; + sum += (float)acci * d * sb; + + barrier(CLK_LOCAL_MEM_FENCE); + } + + half oldv = C[row]; + C[row] = (beta != 0.0f) ? (half)(beta * (float)oldv + alpha * sum) + : (half)(alpha * sum); +} From f905945fc76e15e6d4a226fb451e7f5022ad0a91 Mon Sep 17 00:00:00 2001 From: trm <956303669@qq.com> Date: Wed, 20 Aug 2025 21:50:29 +0800 Subject: [PATCH 3/3] =?UTF-8?q?=E4=BF=AE=E6=94=B9=E5=86=85=E5=AD=98?= =?UTF-8?q?=E5=B8=83=E5=B1=80,=E4=BD=BF=E7=94=A8image,=E9=80=9F=E5=BA=A6?= =?UTF-8?q?=E6=8F=90=E5=8D=87=E5=88=B00.6ms/3ms?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- readme.md | 2 ++ 1 file changed, 2 insertions(+) create mode 100644 readme.md diff --git a/readme.md b/readme.md new file mode 100644 index 0000000..f069f9d --- /dev/null +++ b/readme.md @@ -0,0 +1,2 @@ +* build.sh:linux 环境编译文件 +* kernel.cl:整理代码结构，内核代码单独放置