From a906c350f2261dd6b4801870e5642ef56da66268 Mon Sep 17 00:00:00 2001 From: Andrzej Janik Date: Mon, 22 Feb 2021 01:29:03 +0100 Subject: [PATCH] Make misc fixes (#41) * Update ze_loader.lib to the newest version * Export _ptsz/_ptds for which we have a legacy stream implementations * Stop producing build logs if we are not looking at them anyway --- level_zero-sys/lib/ze_loader.def | Bin 0 -> 15760 bytes level_zero-sys/lib/ze_loader.lib | Bin 75668 -> 77220 bytes level_zero/src/ze.rs | 46 ++++++++++++++++-- ptx/src/test/spirv_run/mod.rs | 4 +- zluda/src/cuda.rs | 80 ++++++++++++++++++++++++++++++- zluda/src/impl/module.rs | 2 +- 6 files changed, 124 insertions(+), 8 deletions(-) create mode 100644 level_zero-sys/lib/ze_loader.def diff --git a/level_zero-sys/lib/ze_loader.def b/level_zero-sys/lib/ze_loader.def new file mode 100644 index 0000000000000000000000000000000000000000..71bc4df4bb6534e42585f407e4fb5f256b561449 GIT binary patch literal 15760 zcmb_j=~5d<5T4(u%FE;pz!;l~9V+naIg%9C8EW$ke7=8$Uhh=yb?!z5EmthmG!@IB!yRZ*O;Q^mz{M#{}ZSh&- z_iYE>>+mH~*n{H|RF|OJ&h;$Hv;r0-d@}rZr?-Lh2r0urj@e z*0)i|2Yg=!E8{L4&O=V&2r=*_#=x?}X7&X8z171`{z_^aXuAyGx-90=Jj`b;rBhhu zw#%C=_AH+K4RqN@uWgSTOXWOfD&~b?qv0@~e0x_j`rLp<#|d|(dj z70&7s_R-AOuH)~(T*ERSKtu6+B={W~k-{>(feiQIW0WP=sw~s1h|6qUgJT>s+w?#( ztJB8UkhBH&1hA|}%-*#MK>sdP898;H{Hy4RCH@NAB7s7`Mw~LcJx06uwRP%DmC78t zN@l0t8`z4K+Mb8zgehGVVp;=*p~tObR@S=f^HWG0l|JknNWDj|!FbxkCPU4Y`+Ge9 z3|QqYXSs%-um>5|iX6{M(uPS3)34YK_*sNj6PwUq*c~0gK3Dkb4xbsm+iu80i`K~^ z`4duO)n0-M&sXX;t2Swu@NB21-ts0=c?6w9=si67Y9B8BoYvdIF3foL1o|lHnfKvY zj>D7bv8kof%l{tU$zG)Q9B1JtoKzpTK4>rEazC6y`j5f>-FhEu3H$dgs)_B~Pw3$5(lldUMAJ2v{PP*^!Z!p&F6U) z%3*A;kmB8gDpdi6pMmo|{7@^m*ywG{TG!D-b1sUtg=WKAnA5IlpQwqi^zi=>^WO%! zpXXYZS#n;zZ_Bk+OfNQ0VFA`#zcbLMz7_O&@_};wR`^z_a-R}A`~8J@`xWZdH9lKR z`C4_oU=55CsoRxF!u+Ho|ZTXAw9J_M7sq;5$)(N-R|LvO_ySIW9vb~bM1 zA{DV6TC3nO>IvsVZjtUW3mrnE6=#P2i5(n!R_#ApvFE_ye8-sBHp_VdZd+vd$1aX~ zYF>TOr+#kH=MuA`q2Dob_rduuMY%nFkGFm9bGHyFvvW(RI~?12d3*)ymuPRzKRMl! zUODs}{`H8nV8%Xz#+(n=$Z;}tX%}NL=YZv2a|_RWj58>6*F`)_&V-)yf}{PqY-4V~ zYgK^Pew=mQPHwq8<9$`lp6uCH_((6(%E{tz*7`QFF;sfgl-ou1CdTw7&Ylr7_B6>E zM~Elx9j*pFWwPrSCe_NR=W|3j_s-lE3}IUO9O1?-8t1ejzW8LSjxx{bR>_A=nw9)i zc&0d+>ShG0-u1kPCDVK{&a8g4mM}yr_CAc1jqzWL{*$?xNX=)nWv?- zcjxqVMhofqo+z2PRDXt^EAXNj>H1V>u+{o0pUIYdN?r@uE``PFktM-SOCxi;^mhE( zrf;en+V#?EdB#fgYMM@C(ldCXLk&Fmb9l|>IY$oRbt|IVj`RixeN0bvm@#D%lCJ9u z^xK=Rh7o16D`CredU;js?afQOFb$vj#RcSH)TsQ^IETuXe5j4f0y%>|#YyTDWaFe+ zmQJ$Kv;{U#K3RvTYnW(F>9l#ujqDe0kcs3C3aJWX{Sz!}8t*c%fy=YtPh|{+(~YBt zF|3uR=rx$fKMm|TglWlc(y!W@Dp*J80a;hcMYfN$l@XuHyIDn4zvg3et74|*c|g!?ED$enfnxbKUAi1 z23fA)N8CZ!_?hPFdC)ZUBK+A!{Q@z~w~1u7=+Q^w`(-u}8a*}Znbh}Li{xF*e zuIKmSd(1toHQd=rOqC{~RAK$<+k>w)vK`A-z;-OJM)hHDIxI7uIeIva!c&!0tF@a4=UiUnRKn!I{gNUy5N zbIBC!d$>2S6ODcD;ahJKc{r8O7GIHBu93%W7m-7krO)9#8S1_%lKFjID|T7ZIu*6? z1|6Thui(2&DI!EtW;;%nD|FqITK4|KS+j&AJqhQTPw`B3zm)T0htJ*caPk>buW6s; z7cH&)ZL$d7lq`(>7H^+H;{@q>ie!b7H&}?W&K7c-SLxLAoRcKuR)Q6Qw>^1hF@;x_ zcM2=@tMoGE_A*+#*vqo7NYeMxQ@e3nBe#g$?9xx!xA10h^-Nv+)?CZvk>Akg+yaX? z)ZFAG$?CYih1&6QYgk-cJLLR5$#STgMJBDI;mumj*r9;aylkonJ^@DTX zHon=bRbFW7dqccEVkf~VeGe>aHE*|XBa*nEAu;R0EpV^nEULROwTDKr-0yIkt-XS! zNl1G0)N}FIb3oaxtawVvCoPXN$V^?Cclq9~?vgANdGN-)tWScn9_Hzt;Kq{c4rI)m z)s5%qQPcl5sPR?E8s{_}iiP#W%2w6Fomgda;^JawkhiAg~n%nWlQnv6SPSUlV zlfAue$MU1;F#N-?G*7Y1W6%F`I<>d*(y3bQ$GNFT3Vmy#9n0S2@}1}Oc7EAI>?w<9 z3rAK)W;kXJ-py%DS{((W^h~Jv^Bemytyr3Sxi!2+KIGOo&(Zn@mCYnRrLugoUw>%m zy@!&jhj};VQ?QJ}9b%na^_`9>Xp#IR>VG+5eU&45dXGE_r%x{C)LlF^|J27`ukTK^ z!Ke3>+MrX7@$95Alx&&$>fERE*zqly!%bN)hm?xVeka`v*uG~Pe(LRr{ovgViFn%; IjNg&?55;PF1poj5 literal 0 HcmV?d00001 diff --git a/level_zero-sys/lib/ze_loader.lib b/level_zero-sys/lib/ze_loader.lib index 661240caddaa48c16b837aba6239073a33de3943..dfb3f84a4b9e74af7c267ba8b3a658794dd5aaf6 100644 GIT binary patch literal 77220 zcmeHQd%Rp#b>4?ZfDl3mAwYm|39kSFl8}T)Ddu@4kU&f>NGT;u?mbBc=FZ&Cy%R#( zBBhkt)>3MzMWl+Th>EC)sHljDh>D1ch>EC)iin5^sEF8aKhN6FnUk6Hujcpra_;QC z*8cW;@3SB49CUJ}Ik@qZ@yGVr|K=>1bKHsZ=gvEBw(YKe&N*@ZyamGRg7J)PH_X^u zPhsqBBaB^tGh=Tbr*Ijb9Ir^cV>4p{ui^>Z06&g*?x1iLp1_S06yA&{@U9sOZ^sjO z_cVnY@dVy;h{6qc0yp(3+=M6a-eVMQ!4tUoNQDpM3A_*C3tWRI@cu~(SK!Ied2k-=LU!iaxp1{WzDtrP@;Eq`e*W(F%9O)C-f+xqb5}&9sCh#Plz$ed9 zxED|0&bbQ6^E-i0O;&g#p1`M(9)S@&Ii8UC48jpW+&=@{b(#W%?*cwMQ{gRm0(VbP zxD-zgU^AW^FGzgp0>%VTHeUk1yjbC8Jb|xFRX`YD;dn*jfeRTEK;3x&_$u;6 z0D1CN;K6={JMjd*hWr+IKc2utvlXt#6ZksfD}XZmI>*xz-#CvkfqU@;zIlSet#|?t zPgl4aPvBeQ6p+{7;&?{lk#iXncmPk}+lZgQU3dbI&Qn0#9tFNLT>)W!2l(y|3QasY zP#-wHw_4#LJb}kfQMdz7;QKQbZo(6I9Qi5mW;}r(pj-q-@Z><496vl;;mddeKRQj} z9z203PEtTzp8$T`r+~ctG4LevQ{YNGfuD?5xExQ8=Ouo6A!7nR#1r@#@=4%PJb|aq zP(a>41^gWOB!IO19C#Y}B!IL$4g6x9!ew}JJS*|cM#cml#}oMFfWrNF0?(eTa2uY$ zuaIv7H{c08H%;LxJb_>DtZ*rw94Kdw=g(I_89xvFW}yPg;y1txvlNg=F90u2Q9#^Z z1YSZu2_VduIG&gIZH+O3r|<-Rw^8A7Jb~Y@R=6Ke;N^aW`|t$*fb*;B!x@y?8BdNUB@R88Wde`k2^_XU;XXWp!%tGU z6;EL142A3P1dc#@1+Ku8<7J5>>ns!aDW1Sl=PEpaCol`;FMvFp1spw#W9{1ID_5Pl zdSLC+WowqKUcPGJ%++hRSOlHw z@Cd?K(ik4D)Q477N5>Y8jO6v9#g%5Wia--yQQc*m^7`27d^8^msT1F52+M~n8xW_F zOF{t@p6&Jfn~mX>`Ea9|4vC1h4|8?Cq1vb?W2Qvfhc(dXAa+J58s3UZt3J4KNux26 zH!EY+P5J40v!2&Nag2$!3A&3rIeJ7n}^XhyD0s4H8wvFb=oC3@7T z1?sL0KBf@mK_M*~+Eg9QLpfS_ru0z-V-;Ff0pZ(02)wQ?Y1GH^&A|*T0;7T*gd4Bc zMUd_4wUq3y)`C@!__m49s%E}E9~;~l$!_GOk!Tar>Pme>KAPywWTb4dXqc~S<;_cE znb&q!cdf6ACe1%x}?@{Ta9!I$u{%KnA60TNVll;^5Nlps2Wi#JOlbmQGCsYGcIVy zfP7#;lTMB-ioTVrXqz60I)>ymudSe6LmJCA=Yy>=I&hOx7jz+sP)KVo!PwDk)T>)a z-ir{ZPoXKg9+E)%+mX@<;K>qV+QE@qwy8QeCdPG9*Mul>3rWj%>xe5yNB*_>*qX7* zSW9M*2)uA|WLH)$%0)y_w$&lbV+sb}E1;{olc2KW=V-6FWOOVarahI2sqL(FmgdN@ zd`YFgXmBte9bJ`+sanX@)zIJ1%Rf}ZkO4rZRUNv`BnArC4^@1rTL3tXc$(jnNTztvaIw@mS*ENN!(P(07 z+*EDA{O$QJ8mXS0H%H|pf2o!uuAT<3L3OS`5`>ncRYTrh7|TZn!`ilorx#Wtu3AG( z4NSjHe_0Z`i{6q85&hGWIkhorUhsre;o!;1$qPCXBdA<_Noyb~#EG+|v#QakiOB?V z=v2JD*7*1()B--frk-sS8unbnxa%w+fomr(R8+2DY%F(A>rA|I2i0s3OApasj@8JJFIaD*?5T2hr05wp)+^ccLX<{sL?v z52A;0IiTdZQVlT{FK$(9Ln|5^LdppNFPx~p?LqWbZMh z<67PfXa{|W=GoQeSSv;y^VIsH08O7oe_0YmpCv@}PfLa;lulfPr0rs3CrUd|ud zP1Xq$73-%hquIBpMJ)Z>mIAY0+>-N15XwM(A5Zg)$}k@vjlB5nNmvwTMaRLF<0CbZ55KLcT&i8bmVz*NnIA7uoEYOTM~x~bIebq^#v$P4>8e58*R(x6J=s)vAICc+!$+MxoioJN$>M6g zBDRZ+lI8Ac+4V`3AreMhwFVh7O=0+M0cl+fTXG_<9ksj~QccMAZ%0ZiZ^@OIcF=Og z7^@nf#(#?c;0mw}R#--k%gY z+Rc~dgOFHV3#TaH?T2wzeb`BpbkUMEJ0o)6D<6g%|ZpmNtefRgcN_4TP{FA94tVOF$9>Aa>tm*XXF&UIE=z zd2@ZEiE{>nd4FYatkDd{H!2khv2k%835EKEG)KjD0I5BzxJy8Cb!Bu}v)O2(1vm+| zd;*$lu$0c%TAPj@aSv&YVJT}}xGb?cL;V8U1NpEx1JG*bYp^8QaJd8iZYc8f#s(Y( z5o5K^BFov)u?IZUC~aq7hc7nq-dIsUT)H@ts|Ks^!p;cEN8%QgCgZOOU+NH_pr#Wk zsg#Hp_wphp9o_KJzhkq>wAb{36=Cca`a} z6s2=WXLYL{(lO43j#zRMYe4*S*LuZ_I~boN1|!@x9tn-_e)1Gr(vN zXQNYe+{ifdJ~gIfC%}AaP*(#2^qa_r?y0NY0xEES#*zT`YM)_fu(SFOple4pd$lD%;o*vub zzi3mXTH~`eyi5;XBw8@$?s8%f@D7GS%7Bv-b>b7$)bHSyl5`Bp7iCIF+>#|;yaJDn z)pC9)MPy{m&u2Sx5XoOs8L6zR)~ae#yNIUm8wtVW#z3XHLH5PQNd#hkDS=2h%LfOs zbYrDKs4Mq~mZ@(#-9ksWMl=R44AMyFh|bxK+8AD}Vkxw}Lt*firDGTjB_I$9CzwxA zPvjGP1No#~BRrvef=WD};Av+BZ7OFDdfq-SMEv$*9F07xZD@A~bP*p*vlzY6cCGhH zp&s+L!U*gUDfWO5)s$myh&bd!HH)z+q#g6ML*U!+UW)rGjo|e9R=?ZFxyr?dD-3sYPy-gUX18D5f#ao>oyF52X+KjDI&>2?r+7!+{J9$NV%hvYH`si}6l+adl`+OmG6OzJ{F~qe7g09TOQN?xf4Bv4~nRbxrsh zu8<6g6EdrfL_8#s<7%_Y52f%eLf~x&DOgs|n;UR)92-FLZG2VGBcg{{Q$D(}QR54F zqqb(@<>{`d4p-$XpjJ^Xo}`mIDr&^l)<8OCdJ1wKuYi}Q>%2 zD|}4Fl7ddHLeFB=BkQ*bNxG7BBUzrCBW-x4pwC;>3K?4Q+k|9#FCNEMvx$vVJpqb& z1$FJ1Q&r*<)a3K=MoX{~QAC&q5*CVvH(JBlR-TS@h#@gQOp$BD{Fvxj*qUFOuWN1a zvM44K3~kxw!CGq+2MI&L*q%Xsk{P^JntVbes8>|iBs-*CcozC{j6w}D5w=|dl1|1O zdE)MBA^J}E+Qmqc(avg9LHCFji2*Xt)+BxrZPVs#MW^i%Pf;^;za4bfOc-mND|9P! z?9dL%;i@>+-I-jwP{mHyVq8SyYiO|`iBL!3Evxf!A{iPNX%`kYrgQAOHdj`f7l|<| z7NR3a!WC!X#Ht?_h-E&S?HYs@_cmdOhQ4Uy0(Lmd`QYDdjqt@gnSjE@mRy}zhU}ykF21C! zn)kiVk5 z-K>xJ8Ew%$ndS^_>uqu@mFm&e@k4{*^3Q3$M35F2x zp_=j-z_QK!3XWHUBA!NH7B{jD%Im8&oM~L&;7vK`ZnRV>iVuv%{c0`C&U>MYu*Z_yA_-oesAm*{?Vi@=N*`};hz`Fn+M zH5#Us6?MRM@?~`{B7cprcJrk%sVENW6bi&StYZ%5P#UgfJc4>?q2|zz#a4K1g;lco z+A7%z8Dposm9dex!NwPC){J{QW0wK10Gr9Kr+6y9f56fEAxd9DtMWh5a$$$S)vIfPL;`>`lOgFCy&# zH<@0hV0fl1>sb}8@-aKTO)dj#m;Ib-(# zi+9P`UBH6z8T&BMw`<0(0VYm>4)7ds{%#q2061l0#%=+o?G71W!laB{4NTo5W7h&x zCTHv_VEmpL+X6fXTrdTEz^c76b{}x^RPX}*duQx^VAVbudl2Z~H)FR0Go~TVz~uch zb~*3@uyKEc0i1dO;tHHFJ!5wOvk%PJ%|Krt_<^YhfghMO15aS$!5O;}*zS;wZ3dnM z&OJ0^j{z$VLwtaf4$s)Fz>Jw0yAGIqM8>WFUIyw%BADQ`z2ZAGR->#`a_Tvjf<4 zb|C9x2eBFKV0H*QlpV$nXEWIm>_~PLo5hZ1$FO7BY&M6@Wyi63Y(87Sj%O#Z6WK}Z z&)CWA&)F&LHEbbU#1^w9Y$;pD`q^vQscbnrjh)U`u$Al#b|zcJUdL9mHEe*L#m;80 zXXmiLVCS;GWaqKJV&}89>;hI{>)0S0VmVvSHn5GXih6YstFd8LXAL&O{+cz}C>vuf zwuxQLHZvUBVSmH^mi-;Ol)Zsn#{QmN&i;X2!Tyn5$^MDGk^M9K7xu60P3+&;RqWr{ z)$Bjmo7sP|YuJCWx3K?a*Rr>=>)6}a_3VGx+u1wV4eXulM)ofDZuTB_6MHYanZ1v_ zpM8Mc!am4uWgkNK{1J8=`zX7eeT?0~KF&VDKFRK6pF(%~8Fm-@EW4Y1j@`pP&+cVk zVE3^vqWk+2`!f3qdw_kFJ;=Vs9%5f--(cTl53_HvN7%R7qwG8EyX<@HG4_4-I69ag zvLCT0*pJzh>?deTKVwg^pR=dgFW58em+V>gEA|}wHG7`@hP}XEWG}Jbvfr`avzOT) z*emRh>{W(Fm}S{^+4k8E*^b$`Y^QAJY?o|&wre&a+bx@z?Ve4__Q)n@duCIzy|Ss< z-q}9czS*>Fzij{PfNXkpVAhu%l+DNv&JM{A%?`^B&t_&vWJhL4WwWxQvtzPjv)S34 zY;Ja3HZPl>Ey#|~PRRbhru~0S`@gAa+v8nttL&5M$)>l0s#i^;0%EDlev|M5krx2D zEwN}YMr3R3fVB^G>+N?%!=`4k;_0GN`pQ~6=`L(zMPn-$`L@QuRR=-x#Xp#1Ya=Oa zQniuUI-63_K*fu=Tr=StJEPKAR8+jYk*8=-<%}~0f^k+TNSl&A9T`>%gd1$C$R}ek z$xTrW^t7+Yge$O*f1^FUfr zE(0S@rpXwWt}Og`;iH_k+B%OHnCPMdON1uwvX7A$w%W@jCG^^=Ui9-K)48o1Pt~ai zCi+QhH6|2s>JCL$))tKy0B*2`vOtF)@DoF=(^iC=+fsxhzZ`5gcH`WM^%osqp(K_i zxfwbC+HK)4*E2o6VB~m!$izf&eAO`G9qQFhOA_%B5>7b_6BbVH2{$)_!m(qUC_24D z%g%2kkaXK9P_hl6)e0HWnI?o1n^g!z*5atTyy6k{Hig(mB9;V^_^57qD7Iu>&FYd0X_=GIN&h-4!sZr2z$d!bomR=Po& zEhRlYl^wWgP{zNlwF;XDu`FwQ!VUnd;T_tN@br>N>Cy&-CTnlbK}({J3iAdja}7kc zjU@&mi7dA4FF0Nunp7ZK;ZE0XM{fPGeR-)%s|Y4oF_t?WuT%;LRmzZo%>wokMhQIK zj;#!6P0}|w2_!ylDmyX@pc#r=^0?HIn?_kU^E%*Mk@s~wxyA2m@zIRm)%ShSG~1JQ ze9c~m-_gx^M>w$2R`&Y44pHB-ja_sP$f+o|Jr5g{35lZDY`d)W7=^q zaThwM$>*w(K)IKf6MQCO-(z-ywP&JRc>}o!v$(3%$(>@OXYT!)>%+cRU@BwZj#W1Z zjAqfOP#T?I@eS)c8E=KnFFdLny8&fBJ(0$*h-XLlZ5??p3gzz2($Em6&hEh$O6q16 zvNBi@`sNfLW}MrtMHob=T_A1mqXt9eC1I|f`Vw1xfix0S;(jCY0o^BQs-`;@k%(Ve zgra?AF~p7XGvOt!utF>mMfgUe2)}p~;n^;V5OGQn=}NtOc6HIwWD!WV_R=B^V=L|- zTI#x<^WL{GUrI$V@Wwsm=o|O<2H%8e$$Nb?j=1a+57-R82@yF0#|6kC{fH2<2}KG| z_a&C!g@cdNCBhLEPu`a^DUptjq|9Q3P~3gg=*EMufOjg^~$OC<>u={-PT5 zJ(bP`+py^)*onaO{#+#1=HPTB2(zXPe{;RT8prQ!er9OJvo>wS-OCY2t_Sfv?f2zv8Q$gD=`$|)nqCCj&uxd}f;OoXn z`|(!p#JI#+dqZgIf|mm0!$&-2?b27{o}C|xCmU{UzoIC*Hko(kqff=>7v-KzyAL3e z<4UP8)CrlCh&iGZ5R#{eh)rtQfNm>N09==7!Z;l=PvoIc5~)5O2AukJgFyE{lO2EeeE34;2Nxt#87G{6k-Pv*h$m9a0t4 zW`r#!`j_I^kmae{rse1~4W!r{5;Aha(bE>Lp@TuTp8YKnycqB=Dw4R1-!yRHKn?)$ z;0iZ2B{7cZarumress(~?qj66mka^vcD_4iAMw#2nt)~HcNG5vr0Uhd5jNru86qLz0avTS8`dAFl zTKzXUfKa1lv{%4h_3;A*q`FN4D$4hvB5Lp7DvwZ4g=`e55Pl}_#W4<(YVoTe7eC!&F@Y}vDJR`|pgE9^37Jf$jN&sU z%2#w5L2uE(13sZ;M-gJ^u#4~zN%p77#Ltgj;Ao4G4;^X|9=f%;=?uLN2nDpkK53_S74<^=L@7h5=myj$FJzicjgmhL zRrwpbdLp8Zp6IW2z(*W9A-^paqwgIw(HfXMMn7*t50?;m@gCuRq9hVQ?Kn4*`0t|C3^0bc!jk@#ALeY}I{$hzQtZNfUzLBy?kL|TlAFmWqlZ6Dq6KHWe+ zYM_duXjHsJoP!!v-u>``=$V6`Sl~UY`w0bJ&l7>`Zy+Q-XJ9@vU~)j6qbO2O4pxM| zO<-QqOi+jw^4VNiZg;=})`VoA&c8Ka9j&l@ialCsBU^ljm8)QT$8}M=o369o$kC@4 ztnjh1;K<8alB}MNm$#-L4AMF(VUdYL5|#_r!z0G-%^I~62aMnJNQ4!dwT9{%qqmMh z7&TN>t97VC5sjOlQZRz#(1YU^(x+D39dSMZA!&lPbGX6C>sdwoKm!*oPAeeNNyF$| zT1+R!)Sxhl?jKx$q;a&9SJX$asof1oTh8W~g6ok8isy>5q^&=B!IM)QO0dF0Tch0M zAc7Ec%zGO*LgF4fkTM9@G01J)oudZaU9~sQ*ft{Wq*OOS__zlPctFHEuIO9vMQBBe zNEp@u0vf)5bing9aS2*e24m1nGThN&*oLyn5dn|I-g&pZc^T-$1S0`~U?kukhy+|C z7MJ6ZfR4#+BrQDmPHy%iNnBC1HWaBfYvn$n4SuOpw4{)px2dm`f+mBNpU%uscP{`k~mvO-W5R&&c>7}$80zx{d`FtD*d68kj^G1S^{ctX(@76XvnB`u( zW3bjVG?_IBqu6wZj&Zd^HqJ)f%RhuvbY&r5J48RrB0k(uiRT6}|SS%06Hn+_Q4pT?#x@r>E>0AsI|9+2BWWI1<+I)GP;5U_ zAjDbZfU;bP8fms9ui|jNkCmk@KH)=u4;32hRs70JjW#Jq#L8F0UqvGv zi&tRk$_v>y`$o2M?U3mvuI-~=YEDgW6D7w2986GpJxBmW&U#0|e4 zvPS)cWgXbp(m&L5A)B=LnCcKp>yWwHYYP5=*(_L+k^i>GV@lkoqJ;xCJbEgZ!CZp8Q6W|sX4mTd|h2*bn?hfelF#zd4i zY8R0}$)sfFaTzGVq>yr%U$}7h8T>!>kB81+)~NWm9n2Lm946Q*10;zGGu}#c2Q_By z{JEsYY$s~Wrw}z6WyV4^W~fMq_>0t-nlYQb?~?xVoXZLkHnh`B1t*4cU=Cuj zLz`HbQ!ZV_;P^<*;~m?MZDK%9C3P2vT1GjyqB!ym#C+%3YjhqvTdc-9b4VKR?_mjtBRIcHH(^IU0@fJ*6J_oi@QI z%e%FUN3fGBn-i%qYGNt|>YyJ@erfI#!lf`9soy=tO^y+E!=D?#*hqr;wIcIN~qrMI%8ll~ym z7xB?yN`yUpKD3)y%#ZZO2O$=Fxv|iz6qK-Z-J2gt_Y5#j^%(U*BuZH}r&IINy~B+3 z=q#l!pV~+bKNzXq$K@`TX1e1{l6(-nJ@)rF!-t}2qD%8f#_pL&^8p^eedM|u-uV&U>ENBNc*W|tS#Ru$E%FY-U(yex zY~eyJUwachQXl2BPNR2~!bGodkjH3tL?~tXR3>_b8I%)i-F$~j9HsAyF|`VEFl7v# z)bEa2FA6$y+#^wg4sn@f_M#uu7Xc`B*%mrK^6q;mm=C4Q?ll8STE1)(ec@q(RqjN& zw-_jC*_6(Ua85%?4=0>*Po^8DR6Uw$Gx;|#C~@h0P0iMiaF~-9Ji3ZSYW{Mh7YpzH zNLSp)M|yt#ws@4yow%MtqfoYZ8OV0ByhKeMxb=c z6TR3mA^vF6d&zvP!>o6Ni(T7p(mX%1f_gaW+iaWJy>&!sOP96!ERwfU$Y+yt9A3IZ zO-ak*ge&nK?G^R$0d6H9K@#>N(Rf6+v zChIa;H>q0?naA<1k_^KKuFJEU9ZRL2I-MtCd9dDu@onZ6!tVfy`>5KPAPM;_^ z#Wq$?-aqC=@|ADEo>Xx4>g^-(w)1BN8@+?rUFk(el|H2IWSdXf4)3l9NX>}<+~Kn> zP)SVucU?iRUp|&nB7<;nj zo7%-YO>m+4lVdKJ3Rq`qY1WplJ*j!u=?=G8a!!$#F4xr9vZ7!Q9vAK=)u}#krREJE zIi|Fw)0^Tw!|(<#`*g*d>U+=BykgoHy5~=+%hs>?i7{@K&D_IUonn2RW)0dX=qjy| z*)D(MUaeWf77!?H`C668<29OBujurw=SbD%0nHjdoZnq)6YorCXe>o+mm6WygbQL|oe^QV$sx@M(DrgJnex$~bQFP-O!_m02N+!4DRG_ys5Zug;tmHERnV;4w(xU*$D<0R7$g@_KzdUW@JM)XTVHPSE-HE zS3XP02{Udoxki8}ZTXrT8AJJg?s}WowlUNlYog!TV6##aMqP2HTFpkAv$w4#QLAI47_ia*6&!!orEt{jM9oSLA+Ot*c z$Sj?owjI-aJzc*}3O*%gxf5TsP&J>Txb6yxV?OyP2@{CQno6 zrZ1svm((l6)j@8B%IUz%wvN=vn=OP9rRQ7mqs3UhH%QD0{>El@jRbbb8ySuHZ0~Pv zUio^YZ4h?Hnpg?@JDb%paY$*)mTv02#HBWC&)y$Wb@&Z}lPt>w8)@lOCdRVMT*h!i zEnj;hYe4+8^xq3k*S2I=<(Qb&T~670mAVAqKM1~JnJF+CDV>|Cy`C!^K0Yb&Ok>i_ zEzO>Kr~gNXU7GczbY*L6>K*J#hfQo75~HEr@TS`3KRLX7^H7>3?S?T`s&8}{`8$1& zOEt2q--j0Z&knn`r`lC&BlAAKI{PnzFKMlnQunk5|EuJd248y9-zIv8H&M1;^+u6* zB7U3f-zcAD;)Z5u**2OQH?MLTiS?Xq#rf|JCkAU{g|{2d)amD|1t+oKTbN7ouF{#9>t5sW5#!2~ zv~(&{dvpKgv1;3$g`s6iUb^Kba`-JSx3Ub~4PWa0>Azh*d4r&^J57m8=koE9lydVs z*SgG}P0ilqcB=ip)r$kS+}Ru6)SAw99`C4k&8=*$N~QB{E+0&IQiI>LkjwI>R(7s; zd5twqx0@=g@Cl zloc;N(U2gi=NALGoSGVfDVNKNF4=C2ahaI|0X@c_>$!R_j zAV?QP4!G44)FFc zpA+YGJ`!LKb%$lgn^bCV3vgN=9O){bQ$58;U1qG!^r%)PR^D&dtk&mJXd=tzZff50 zG0GOLL}eLMquU*nF;scV_Y8?X@Z*LJg|{?cUuo5}g?*Z(54V6R(hvmbw08~>B{GAqR;)5&DYcSiPT#ErxmYW>ESy#n<~{>WnC@D z;hDjUx}S?i`UHLx^fNwt&)RD06wO^eum0jrcO{%?v7hxhd(um!Rs`<$Ii0UF(d3qG zwW+nl&pFJYZ$EX#9hragyA}6@xc!ehQTp<2InhskKE!Xn8AWN!=6GuFgFoyY3J1Z}Zyqr4Z{j zu3cXaac=Y4^_387Pit3dWO~41N8Rel$dp=f`>N!`(BJ;@YS|tu)s7!5c#B0EjQ`Qg zhUI&a#N6X+1$*18v!zp=;C;yOhQGo`JFe2ak==N1U;FEZw^&X{Wp-Ka)EnnF2zNb2 zTE2u6D}diDSR)@O>?*GlXG$I}nDxZEr>Rcm^|uUn;OmMs!R7NhQL7&@ta2kYmE5v5 zEA>retk_KLX&$HSEPl_h znyu8{a!1wz_!nayE4bx`YHyh%?~L4r+4l`I*--5%duk2samgO|W+|;)Wm|J<4ebY% zu_tfGDYhRPHn9L&Yc#vvK}fOvh_dx0ZK=1xCnTHfj4`vLpNs9L^iuuNj~z}_)E?c* zPMs-v(qRo{ZP_}STG{xC}5N4_ykX-kjpsePKChj{5HjJsk^)$FHjX7hPtN?5jZQ+>uSY`)&rsKnU% z4BJl+TT^}93zE%Dt@&7HebAg%(bBah zwf6L)&)l=pOuaY0>S@0?hO`WF0_ zJgo#}`>E8s-S7PvxL>NL)TPTHRr6o=xk)|f_I)6g>OX{;P3d>bof`XIQQY!%7OlM( zEBTRmE4NGgN0%>JwaeD4R4QM!SW&OHaVq~L`N*q!&)x|VtDG$3Uv?d}S$mhik)1HU z`e6-|Yp+Jgltr5D1uM0ux@ZI+haXy*!*0T$e9aDjLbz;=PK{bSBx50N zq%O`kRQWgSyX!F{?_vBqx;wUy!?OBNw=qb4t#e#51~6oiSBCTE%1ZO1yz3R6)c(~@ zIv&E(tT=^=!}zjB2m?i3pPfmqI`1smm*(qQ8x{>-RBv2d%ZE1P!X9h4xu2Q`??QRU z#wvpwwgZVBLO(8gD}cF&=M)kvWP1_zwkFL>=Xh#_nQEAIUiUD?kvCF)Ds*qc%L_Y| z)uqch!Ml&)ZEt|OG;d;;Z(qYp<_@VumoDYh3d}TzJw;qP)rooJeh#m+r_yNDQV}Sh z@YG%5U;P+^5ymQ4QCQkHpMABq6bM<)wpwaLRUV8iFfiD4i{3` zqdnNvmnaXmSeuR3$dXEJuvIg!LuXfNt*$MSq-L^*crifnc3GyW*~pK8!9R)B4V7U*7Mfd_vGAse@%YBC#P@K z`_@x?)v4t@hn!Sv4{kVl$}9Tpf5*=`Vcwj%ube%1j_rP&`g!~b$Iq32zitX+lh!hJ z)k?;$zL>GM-OkuMrYT&GCveTq3YX#uymMCt_`MUjb}xl%@dVyAUEz8>f$L@|+>9si z?xPiM#S^%Gw!$rV0`ECa;Wj*h8%|cZ6Hnm1Co0^HCvYRe7q|gW;C+1xSK{qx6PvApHm%vqc0w12La0#9qFG}1z#F)V2cmf~6 zZ-FQA1a4WU@ED%JN6%4s7*F8V0fl?<1U`0>!iVq#K92MWT!tsdW{FR{fiZ#K;R)Qf zLE#}hfln?}K%RdRxcyLttMLRrh4ctqj3>u)5_cdR0mS_d;L~R-+>a-4=RAcQ@dQ3I zUE!^Gav)zg?%Jq;xZMSOwxRGep1|G62La^4-N5ITD|{YL;PWRc+=3_Yg_#Q1;0b&Y zyaHF^3496Z5x5jj;GS(2M)Bl8*>HS$l>+Mimw|g{C?Je`fv@bWunA8N)E|!fkS_wr zlly?Lo~dvjp1}RcZ-G1T1im(30dfBt@W4R|$fE~Tej9ig;R?JTPvDVh3aHzUa6B*Z zXooR@r|<;6bG8EN!gql0o~nQ}e;4@Pi3+H5-vb^yTH)<@0^gsi@D@BdUXplxBVz(D z;t4#_RsiP{z>{YyJcuXo1Mmvmizo2Z$qFCG6Zqlb3RmF?JiV*JoABg7`Z#`sd=q#a zPvFOADLjBD@RP*~s4G7KetL|;^>_k5+gIUoJb`DXC|ram#|siaZ!#u;a{f8+i0!p5>H_Jl?r#_3G8sJ!ZmmT zJ5EtR7&~&jBr)Y;mT~MfS>aMVft|NixCl><7bSK%mt_Lr+y&Sb=@UTOb_I5uuW%cl zz|>xj*UA z71dp`F>iKO#iOtQ|6XGBhJML8vRn8=d-SLnV65s0He-3_hk1ok-dPJ)XBOl$CQ$eIy_2)J8_fLXiv$M#EiQA8ytfqGt?+ zM!ln9yrx#~oZb{+1fyj6$3tn?JNc^`t#!4ANnt2TuL;QoLS58oxvfSzg=E`#t>ZMY zMba%Qy=-J8AF4;x3eSN4ViaGy<%|p3F(4lp(4>)u&n%7p)t|5&j z7vzKE9XfE6QWta~iBL$ZFT~ix|0`5`h;^j_iusdAW!P%CjhIx zg7P#>lQk7kx%iUf>!gfTT~`;bMx%|Xaa*+k^S9@_V6=W_-X4>a{KZ<1xOy7A1~s_? zNf26&Rtg@Y$2 zCokwojG%JyC9Q#|5GT%-&Oocx5R(by&?$I(ZSe6+s0Dm{O+DKvH0-&Cao1Tu0@qGn zsHj}Q*jVnK*6DcV4yxH6mL9rUZ!W6SAN2FJO%oGI6-rJUTbeg9%;=n=zjSEAO=riZ z-1=$B@G3$l6iU{0wsdHelst8^WcB<^#K4}JA@j)s2b#v*ORByC3{->mbg^)mj@Zn938s8~O38O^>)En?~4wiKB4;+C98f=~wP z`*@lwYa@JoH1guNCuJqq$P!miV|8o1J(w@(blUZGn5pwQn9=k6Z0*%~XF+>-d<3sq zXVx0yMf5xmPfxsz;H7eXeb^8}i#mHc;+$nbjL= z?RZdK^#zp^j zq%LgMM(Tsh>&^2F`(ol|OCvr@vH8&9FF&n1CNX@^NX8-H<>{(H+Sjx_JU!VIcpt_) z9K%PQJ)M+O;A zF-ExJ#yOKoT81c8_G!cT{YzfbKxvUf*hC zpL{T1S{v-N+QIlnr9vUL&d(#EP@jdkSa@)c!hOM6}8lyI%@f?!3O2{@CNMm@H~up z1=!H*g>t2ECOS*6<7g}oAUfQUQ@xvHnODzXa&dvrRZ+jGadTfXPf{nF$ zgD=VAWpVI4(Sk8Il@o)2cQ6c62ArIz6Q7`_e#f+wq+?LNC{se>mMrn&6}Z!BT9lSWiQfPaJ!r(1S$1oU5Kp+xMFrT2F$S3#) z@=3c!ctZIEm3ThE)6NLmR8AlCynSAX`0d3w8hKRP(C!52B0iR8F?yrzT8s2TJ?3qN z5!mBU>;WIDDaYIpama^i7GqOLJLYSLz_*dT6fdng&X$hqmxGEyKTBIz`Jj5(lj`GT zE2QZ;QyuxX6;ccNP)+XzD-_5CLRNVERVx^V?_ueUjR&<{XR3o&u8@xBOm)obR!B4A zY3a+CuAp|vi|Ufst&o1y+k}BSfW{78yw&Nk%X4$2&DbgholzBz=xP?V8ZExkI7WZc za3Dj&F+YurtY(PZLcEinUmxm-2~NP(*RYdgREV>$VGcNL{`HY8B<;NjkZsqDEY84Wv`1ry$qy3V3@xL0voM zRF(JyHTitJH7;0*C?ZS)2@6HT8*AWfCQnB?#E_UDrpUDse$4RQB5b<^B%O>m^2FWM zLiC;RwTqD?qn*{Jg6y0(Lmd_~73jALWa8G698)Ex9VM4cSR8TzpAcH%*{i zR5|&wIw5>vNydST7p|_x7>=b7QSsB35m#WGC{VY6bU~*hAruN{qQejPBY#DCyICLc zGuoniGR+y<*4tz!`A|*LLPP3N zZ$FH% z+F;&Zasj_W;?2t~Vw7!be3wN9-_jnL4pAq{c3zpbtcQgY%E)6+4rJm>j;H(>QvuUq1LZ1U>W3M-1b7h`x*4`|fT54TCJwOf7TDMU&iN>8 z`Tzs4zjH5e(#K%K3h0AfoXddCz#BdRdoRF-+h89ESolfA9XJ#=VXg+IehT>jJO^yJ z1K|T_e;T%pfO&Vq_7*VxGmO0zcoEom7s3M?pM_l=;GDY=2VnWtFXfaEW96f zs(|@lL-@c!4cnR3}IPwQ* zKLK6f>?dK@2{`o!C>!A9r(n+uIQ)kwUtrg#8G94(JTUqr7fh*N78v&To)* z;H>As11$V4c!2)jq5cBL{2sR4fN9U8Jb)ekfVu%}1~$Hcaskf$BgzFh`$d#Hu>4P` zzrZOkq0R&Se?}byW^P8@fxf>mb_D>lF2-EXK$fw4fc{AtyA_zVO~!5pW^9|W>wu}- zW$ewsW}rDaV^0Drx6jy}z_B}I>>6Ooj*tOf0xq7CvB^7S>{4Lcoilb3@FH;TF3<%A zb_E|Wf47X?1{^aLyg>i%8M_^rzemRI0{ZvN*bTrQdu8kjVA9?hYXeUMXH7#Gz`T7B zS77G88M_ha+b?5R15@`0KQQ?K@B?f*p1@{cd>|oZ%4q=C~!`R{M z2sVQq$!4;n*wO4|>=^cPb}V}ZJB}UCX0cbY*=!E$XLH#+b^<$*{XcdR`zv-b`)f9z zEno}TBDR<|dU-?DSq z-?7)Swd{4Q#@4YxHpFtao(;1Ntj^A5=dlJGL0xIFQTBS)W@D_w#@R-8KD&V7U=I6x z_6GJ3>>~C?b}{=$_9pgE?9J?-*(K~>*rn`W*=6kC*jw1Yv$wMUV3)K1WLL2NVpp>N zW>>NQVOO)avA45#uxr>m*|qFl>^k;tc0GF!yMeuz-N@d@-p@Y3Zekx~A7US7H?xnh zTi8e0t?Xm$?`a( z_EmO2`x<)y9n*vC8|<6xTkIkBZT2vGggwf>!@kSD2a{FbXOFWd*puuB>?!s`_B8tu z`!V|o`ziYwdxrg-{eu0HJF?6~as?EmW4|JALn zRkyamJKGl7$0lwNnsml0+}tdZ4?buyim)v557S&CXE(F z#mgHTiUw8ANKhacXQhF(De2RdfuBIQ!KR9QGKPB`9pUA^$Ew%}zb&z~Q*el-R4=W< z)h89s-5JMm81SCYcZ$e00ksveC6;ZRu)8-?qZQ?HF5+aGjB%-CX~qj5<+R0CV!Xga z7adsYGjW%FjJ#Mhv0PF@uaWA-J1;Wb+g$Nfor++hpR}4{LJ_CuP;_N&qIdz|re!D# zboc>3G1NM3MYy>wML6{6+#vw~YcN%lTTZkP)3}LMXB6 zgg|61j;hP68c}akh!aS}k{}Wv)h%z0g5*wW0}JU!pA+SkOT1w-5`;mDMMEJMQ=>9l zYScl;Le75e1|;0vx+xryY{bOv8sp{{G>gniH%PNJq^GB{12+xI_)lmx!RA3M%i2J& z1Hh_wmo^PNy<}1<+9J?osCm@Jfa{#HAwXm&Fb_Z^k)M`P0LQD#d;UZ#+{4<9#;q*2 zFE3(wRo?_F25zU9l}h2DO5LuvS-@W6ro5-y_0@N+Njd^2fyBp6Wmk9AHA8WA8#gg> z(*g@;UI(0e+rDl$SId1ZK6LR5#MM`wElI*f{ zf4A7t6rBgZMC`bixE>tTA^m=)L;M%AHuqgga6lt%ZrUc>rM##>?YJAmpEYe1Qg^wORY@$Bjfsw3}3q1;tc z8XDr%-PO=SNnJKVRt5_~Unb(6xO3sN2!jYs6v%{EG=rh?k}$V1eTgl;QyB>=ape#B zfbNqtRnr}dNW^dPLD9a&7ve_wneY;~Iw6*bB7CD!gkL<0@SG@$5OGQn=}NuERl56U zvIrzwdXJBWu?1K1EOlMadGA|Tq@*Gkc;lXO^o{#_gKt8#gd&Bf`##98hr!3`65)u7C$Es1lt@QMQf4tiC@zT^1>+auHJ(m@UYQ_l zx7!cNgg87oktlK6-$YX1fU@slMZAk^O`d>vQ)vtCO4)%TvR zi3SqLHdAeR@sNv2MvV%D%7*%rYcbu3nmaH_VG-%Ty&yk^kSuwvCm4plcHO2ABBV~Rzri=OpK5?*y-IPal8}vLu^TuiN33V91IkIK1izXd)VsySgKfM zHPuJTb9b4~0S`@Byt8Pf3GQUr+@&vC7#DsYT%8l7P5p zgpvtNC<>v8{6#h9+Y{XhPQa#5B_{&YTV0V@o1?*zAk6YC{LP7flXvFieo?!O6BOJz zgF4GCEw2Vo?YkkXv-Do@rOwzpfr=t<7t-hfbY7wO!TFG$*s8Ku(y1!?=)O zit=o`!>T>`a;h66?Z;cW6XO!c&kdoe3tkG04!HhnN|b4G7l{y!6tuI}`cq9f!DF>{L&H z7z0A`n`>QtP>p^AEzF^UiD#%Bn9|=gGbeVU9O5lE^&K+N-eob6Q#yh0==q$0xAmE? zkbmgIU6!1_yh5siTI#pOME_Epz_C1a+q4{=rhycj^EO6KIC|Q`HFUQRi~ z#MKm($vASd1#Q^07UE$gQ49^HCW>E0JrqBSlPdm~a;OXW z7!LU!jkAMFedIKNxH!b#DHI}V?~W>uP)~(y6sZt?Chx_G2$O2@s~{IY-4iB(kGv=+ z-Fcunkd6tNOs0(D%M{92bQwW!(ZB;fp=C!AV(8q6@DNG%r^&?6k6z%Eh>#DR4-p=^ zwYljGy$%Qkw822?@ScsUe&KTv@;7w#L`0q9(2DR8=Qqf2i^b@BXEwA3CXdlCR?u@AL|(i{ zxL=)!M9@2xVZxcy5V7el#F1OpSqtOnoU+jF#X&XU z;`SYO1uP29b_Xfm7wCr%OleC&SCO8b0Iz1O3!=y()^L9q|%zPCih1*Sx;yjDv>sd1vYx%I8%(5xBN=PfjZomYKQz zpTPO-N-JMTf;9lMy0K0@jXlF;TeEeTZ*or-aXMr2CVS0n(xTVsfny0j+eKP z9t_etpJ0)RlL?j!7NH|X*UiGP69){M^c;c}nzeB08b!Cx9vC%L6{~)#tQXCT9}O^q zXS4<+c{BSSv$gjPX1T%b7)~_}O=buU0KN24$+f9T_yTzx7pssM*QsU?PC9$csW)#y<8GOfGuI!lLX=JP;(;N z45zg#4R3Eh+PIm_gNT(}Q$ctplHF%$QZG9PsZV&v*-^JPbR%JjXa#neX;3`xQZ^+? zWcakh#1nc}6ojX}o6ND!!2z2=>Yjo zkC=<0NTm#Sk)4L+qyy&jMlEO=-lO*Pw1~5lE1V++@!%(_f+;mb-x{-|cW_0Bdj)oE zWsu>aldH&=$Pf=P))n`T`^dLBZ%zJ^<)vzVK&Tj-v@_o!?Zz@1;z0uQ_f-v%$O7{Qu4sf~@d{2|c_I7e$H@Mv9WvcH zwSDwUjS5`nwcbo|ocZENksvM1TZxH{Uda%3P#f$rg;6BYIFSSzH?&y26N>0s&SnqP z@jy8#D_L*cl&@{HYD0N@?aDKjFJHUnlrvVIws!G?H46}R<#X&%qXDB?$Im%o-kiCw zoIQ6=pY>1L#_H>vbKLRD{dH4@F4=~?=d@GS!ekC?%gBGn4sioyhp$#YVf_X6hV&2h zoX@5%JhncB(mH&W_Sy@7z(f7B>$@AKZeYh| zNB)!TVkPQGw&&FC;R?e~;@_^YadXM&9xLjD?bcXpeW!2cs-yarFI#z<^0yv4;}?$# z2aG0(zoz2!`%usp_|e@5Y%gB6Z5jYbW59cf^Ks~OM3zY zseDVcv^gq^L`&r@jsBa@=IfR=Tf6fX{fPddo?LD#2Znl*QM)J|0~1X zF9R>{d;;8Zi$|Vzq12_*JvYM5D<%;yplu(7y28;JEzb@_( zj)izRmffOQj+5rWtnA-{*-E5Myc}`ZGaL{5Tw|qNNR1AB{vW6d&ap{KZ)HUntspfz z>>XhbAIhu}i#d_e;Sj`Nni~tfN}oPEQL^t>vi zE}z;+4L=OL`?=i3(o8kZM49gIa)u`UJx%2Slo7oMIm}xvm65sK;ozI@afY+DbWKW( z9{ly(^cW(Y;giTT(WUt#WA{vOALQ}d=bWqX&WZ4HPX6l5(=ArV&3a=cw#YjSXX>MD z;X*B6dlNkp?*~oltRvu+Qkdu!ct2<}njH~JSw5ACUV-;Uh7)Vue1}V%U$4ZNS_R>) z)G&sQ$5&(4i-OJ^_h_&l;WEqYML(!d-&5+cEp$%g-Iue@pv>;seo9)tY!iLqk%Cq3 zM7dW7C~4W0&W>>M5r+3iIO8r(>&ZQtDom+*#8=b^lYi}j5|_@`)NGyCCCZ$=%l|Fw@ATbcDiPi=pB`#h1{fW_@uOJa- z*@px!w9p8YPI;midqs#pn)F^W^EYLha=jy5?AlgI^PI>ED(}VkDk@=iuli8h(q%m- z(wp&dim!}NUb;g~Nz3N4*}8H@sXp-W^;*Ii9FR-%CDO<@W+>nII@x>kBNtWDm>MPc z+B`V*y{d{#sS@NXu7t_D!Br)7{gHVb@0a+x9^nk%52NJe>ujQ}%(L0u>u1$?6ZMF% z@sm{dtVfz1>5KWB-id-!Y-9E0{bP0{U-=sz-vn2e*n0K$k$BrVsbHfQ|EiTGK`lUaW1U z*6$YBe7&l#k(S8cZWh{%nEKP(UX}7S(kA)5V3E!0-YBEArTf_aNNeG<|HTDwXmuQQ zj8a>_x~WNt_w^+Xb8vc4I<2WbcBx>rZ|YV_U8;|LRl(J(I-7VaKc!&vuNl)Mmd?*q z+gw(#in`OI5{#siufd*LF!p54H?@n$PrHfBz2uk+rUKTPTAHGY;}PdB{5%RZHOQ+@A%<`vVv(B*ncUABJB zNsMu?wwZfat5dA2G;7c>K&7-MM%LAuHEec((w47Pi9BATdG(4;&w7qjT|Pszh7adg zOKswvi65nnnPi}(WovPNWHjJwN3SuQV*Ms&@PX2oFX2SZI?LwoZT}qE&*AHpf1`QH zt@{*t={!%ocbu)cBgQajW|!uT%oh3j@ZW0Q!bCtS$)(FVHE%peaLZ$!y?EbA%^Ux& z;Of=9F|xkH_sd>eu;D!i85&*yEuWi_^#Fb*WUb_E*7-FTT3yPsMA{}_<9MB5k)2?0 zL8UZXYX7)K*uuqTCAEt;r_}sx9bx0Ir_G^yJ0DagEu2g(%bQB&pv|l7%g`b$&6XHZ zhHSQIRT20;N>}{(_lS(^6t|_P9^T#El^-{L=XGI$Wtf)-5a4b!%RVmL@k8?iZ-TE`l3kX|p z@-%gp`9jKeVY4<;ALN#yD$X%er(re`MwFg!+m04v`Q9KgPXE2l?3&K2#+#Udy}{;{ zuPoX!Up3Z5-}etTt7G_&(w6S~5_9^CY}TI5=~H$1je?Ubl>{4U=~O1hvWs2Da6>I$ zdn4;I`~>qq3QpJJVx@9S%ZMQsNnmq?ucq zJ@uab&knmZc1h{V*4ET}(IpO>*vccOG^_BY+U37EynI7X8tSaVm@3sv9Y+41-Q!YC z%=`Y;Vb@khE2TEl!}IfvmkB;%hqN#=Y1uHX#2@LA_z9JNbNL+$sa1GWJ0Ndyc?)xV z?m@uvtu|4<|L*e&^RAWBnV2`c)#D?^V<~CrRHk;t|HEU|wh0SEvXs1RPn9}*bh*o| zEZSfz2n-jr+VHuBXj zvt@Zun|iL4+lk)tZ5F#@NSczCZpW#!p>L+?^Ic(1x05K%nHu@7 zvp8`)$o^oum8HF1zkfGn?#&s$NK4_TU$0kemI3zO)ucp!^qvs2HvLYUP3dw?^+z{& ztmd2a70WcWw*Owm8Zca6Ev1Q8bED)mCh*m~zeryJ~P?!M3Aj8>$cvc6xk7H>Hi zXN)=NYWbU3E&V`%m3%ay5^t(!x+%ch!+cJxU4AgY9P0e7hSbH#no8}50-V+d6)NR( zs;Bs{%Z$a69@VPIngL(MyjinapM#)@EStNj6{wF;wrC|P%b40Lx`i@^Do^>IA<+kZ z)UcuO7U%10;|*;=oMvfx&d93M4Dj7bIlT{YP};KjntGS|7-jW7^HGU4u`2y>#p--S zgwmDI+eDxH37fB{?-QwY)7uoUUg_bh;Tvo9Mr~aq$3cX_^QzB9BYgs2NBpGE-m|uv znzh{S^Xe~(R4d^`i~W?(*^^!(wIXnb&*^-9h9$zDf;I91wn}*&IhXk| zus%>Q>xp$wQ=Q7|uN&^b*W+k{%jb2XRzGN1q**DZ-b9XHrW|tW=B7VS*7$+{n7UwPE^z$-8oI2DS6ys4P|ZFI+|M9ctY}d z!woVyxQJewJF>FC&;30a=ZyHU1YLWl-TAJ9><(kUfpV-X3 z$=yiX;U{T+s`&g*@6sfeEzi`5^E1jA>K4oAXXFil@9#Y$+2nVEd(so8<|aR4R-}8#we3ZNwEBTRmD?d5#2bV8ewaeD4R4QMv zSW&OHaw`8Q`N*q!&)x|VtDG-}SbLYhk)1HU`}e01b8kk*L|uBxW9`)lnXS8QIE9MC`07jX{NwJG&(5S) zohM0l@m=)=gXc9{=Qr}9;au38teX3&dGI!rx6`Q&ZrGY^+Y+|Lc`w*fRtL8uZ0NiB z*CfX}c^e%}v_maD8%d2=lL`Ma)T(xSe6+F!pUBwJhjzR@;aip0hC-Fc>Pose^O~|M zyo19h`^~ELC$*-rqr<3+)NPt&Q<-WZQ!LK0@e%iePWd{MSXbDIu)T(VfH9bwiR>*cE>=T-7 zE}i445oW4k)_L8-6i41j`Ki#|2`?|~R92TR=LGK_hIe8EyQO&(yL@{ZUNUz`CAxGe zr&eJ0a@bSErBj`lNAB(LO4}5zRwEUG@+nW``?O#TdfKid?=gHYXCH@M&s$T;FP;C1 z+P-hWuRoC$ZQiAM6K7-gGrViswZXi-Mp_ll3O-!^yeU6UvK2gn(JH6 zB&2MLXGh+X4?syDAX!!8&fN)J`4mRD_&fP@hYKm}(H?B-OM3@eto+qyQLQmJ-Y~C2 zXVx0ydDS^rYUX^97XuV;#WGFJMh*_~MKil>ZAzUG?W4@1(+nEDDqpJ-_u~%nIKzga zN~bfyd8o(9Cw;X}tzxOBW;};^oczlTy~*Fy7<_n;w?}!Lcy~C$;|*QQE?ugL{%Hnf z)I(lxzT2LtOGk!Sxy{UqbtzS*GeeyGCA$}OIFZ()6t5W6d6UW4jSu%=giF2GA1yh} zfZm%HliL4%nZwwlwvwp1$2g2w6k5{c*W@a%b*6fxmj_r!>y6g1Kb)5D@guYIBQcLS L*5}3O7Uun5)Ai^$ diff --git a/level_zero/src/ze.rs b/level_zero/src/ze.rs index c56321a..ce675eb 100644 --- a/level_zero/src/ze.rs +++ b/level_zero/src/ze.rs @@ -270,7 +270,7 @@ impl Module { }; match ocl_core::get_program_info(&ocl_program, ocl_core::ProgramInfo::Binaries) { Ok(ocl_core::ProgramInfoResult::Binaries(binaries)) => { - let (module, build_log) = Self::build_native(ctx, d, &binaries[0]); + let (module, build_log) = Self::build_native_logged(ctx, d, &binaries[0]); (module, Some(build_log)) } _ => return (Err(sys::ze_result_t::ZE_RESULT_ERROR_UNKNOWN), None), @@ -346,12 +346,21 @@ impl Module { d: &Device, bin: &[u8], opts: Option<&CStr>, - ) -> (Result, BuildLog) { + ) -> Result { Module::new(ctx, true, d, bin, opts) } - pub fn build_native(ctx: &mut Context, d: &Device, bin: &[u8]) -> (Result, BuildLog) { - Module::new(ctx, false, d, bin, None) + pub fn build_spirv_logged( + ctx: &mut Context, + d: &Device, + bin: &[u8], + opts: Option<&CStr>, + ) -> (Result, BuildLog) { + Module::new_logged(ctx, true, d, bin, opts) + } + + pub fn build_native_logged(ctx: &mut Context, d: &Device, bin: &[u8]) -> (Result, BuildLog) { + Module::new_logged(ctx, false, d, bin, None) } fn new( @@ -360,6 +369,35 @@ impl Module { d: &Device, bin: &[u8], opts: Option<&CStr>, + ) -> Result { + let desc = sys::ze_module_desc_t { + stype: sys::ze_structure_type_t::ZE_STRUCTURE_TYPE_MODULE_DESC, + pNext: ptr::null(), + format: if spirv { + sys::ze_module_format_t::ZE_MODULE_FORMAT_IL_SPIRV + } else { + sys::ze_module_format_t::ZE_MODULE_FORMAT_NATIVE + }, + inputSize: bin.len(), + pInputModule: bin.as_ptr(), + pBuildFlags: opts.map(|s| s.as_ptr() as *const _).unwrap_or(ptr::null()), + pConstants: ptr::null(), + }; + let mut result: sys::ze_module_handle_t = ptr::null_mut(); + let err = unsafe { sys::zeModuleCreate(ctx.0, d.0, &desc, &mut result, ptr::null_mut()) }; + if err != crate::sys::ze_result_t::ZE_RESULT_SUCCESS { + Result::Err(err) + } else { + Ok(Module(result)) + } + } + + fn new_logged( + ctx: &mut Context, + spirv: bool, + d: &Device, + bin: &[u8], + opts: Option<&CStr>, ) -> (Result, BuildLog) { let desc = sys::ze_module_desc_t { stype: sys::ze_structure_type_t::ZE_STRUCTURE_TYPE_MODULE_DESC, diff --git a/ptx/src/test/spirv_run/mod.rs b/ptx/src/test/spirv_run/mod.rs index 3976c76..7c790eb 100644 --- a/ptx/src/test/spirv_run/mod.rs +++ b/ptx/src/test/spirv_run/mod.rs @@ -116,7 +116,7 @@ test_ptx!(cos, [std::f32::consts::PI], [-1f32]); test_ptx!(lg2, [512f32], [9f32]); test_ptx!(ex2, [10f32], [1024f32]); test_ptx!(cvt_rni, [9.5f32, 10.5f32], [10f32, 10f32]); -test_ptx!(cvt_rzi, [-13.8f32, 12.9f32], [-13f32, 13f32]); +test_ptx!(cvt_rzi, [-13.8f32, 12.9f32], [-13f32, 12f32]); test_ptx!(cvt_s32_f32, [-13.8f32, 12.9f32], [-13i32, 13i32]); test_ptx!(clz, [0b00000101_00101101_00010011_10101011u32], [5u32]); test_ptx!(popc, [0b10111100_10010010_01001001_10001010u32], [14u32]); @@ -225,7 +225,7 @@ fn run_spirv< Some(module.build_options.as_c_str()), ), None => { - let (module, log) = ze::Module::build_spirv( + let (module, log) = ze::Module::build_spirv_logged( &mut ctx, &dev, byte_il, diff --git a/zluda/src/cuda.rs b/zluda/src/cuda.rs index 469f8f3..1eb08d5 100644 --- a/zluda/src/cuda.rs +++ b/zluda/src/cuda.rs @@ -2454,6 +2454,7 @@ pub extern "C" fn cuModuleLoadData( r#impl::module::load_data(module.decuda(), image).encuda() } +// TODO: parse jit options #[cfg_attr(not(test), no_mangle)] pub extern "C" fn cuModuleLoadDataEx( module: *mut CUmodule, @@ -2462,7 +2463,7 @@ pub extern "C" fn cuModuleLoadDataEx( options: *mut CUjit_option, optionValues: *mut *mut ::std::os::raw::c_void, ) -> CUresult { - r#impl::unimplemented() + r#impl::module::load_data(module.decuda(), image).encuda() } #[cfg_attr(not(test), no_mangle)] @@ -2736,6 +2737,16 @@ pub extern "C" fn cuMemcpyHtoD_v2( r#impl::memory::copy_v2(dstDevice.decuda(), srcHost, ByteCount).encuda() } +// TODO: implement default stream semantics +#[cfg_attr(not(test), no_mangle)] +pub extern "C" fn cuMemcpyHtoD_v2_ptds( + dstDevice: CUdeviceptr, + srcHost: *const ::std::os::raw::c_void, + ByteCount: usize, +) -> CUresult { + r#impl::memory::copy_v2(dstDevice.decuda(), srcHost, ByteCount).encuda() +} + #[cfg_attr(not(test), no_mangle)] pub extern "C" fn cuMemcpyDtoH_v2( dstHost: *mut ::std::os::raw::c_void, @@ -2745,6 +2756,16 @@ pub extern "C" fn cuMemcpyDtoH_v2( r#impl::memory::copy_v2(dstHost, srcDevice.decuda(), ByteCount).encuda() } +// TODO: implement default stream semantics +#[cfg_attr(not(test), no_mangle)] +pub extern "C" fn cuMemcpyDtoH_v2_ptds( + dstHost: *mut ::std::os::raw::c_void, + srcDevice: CUdeviceptr, + ByteCount: usize, +) -> CUresult { + r#impl::memory::copy_v2(dstHost, srcDevice.decuda(), ByteCount).encuda() +} + #[cfg_attr(not(test), no_mangle)] pub extern "C" fn cuMemcpyDtoD_v2( dstDevice: CUdeviceptr, @@ -2926,6 +2947,16 @@ pub extern "C" fn cuMemsetD8_v2( r#impl::memory::set_d8_v2(dstDevice.decuda(), uc, N).encuda() } +// TODO: implement default stream semantics +#[cfg_attr(not(test), no_mangle)] +pub extern "C" fn cuMemsetD8_v2_ptds( + dstDevice: CUdeviceptr, + uc: ::std::os::raw::c_uchar, + N: usize, +) -> CUresult { + r#impl::memory::set_d8_v2(dstDevice.decuda(), uc, N).encuda() +} + #[cfg_attr(not(test), no_mangle)] pub extern "C" fn cuMemsetD16_v2( dstDevice: CUdeviceptr, @@ -2944,6 +2975,16 @@ pub extern "C" fn cuMemsetD32_v2( r#impl::memory::set_d32_v2(dstDevice.decuda(), ui, N).encuda() } +// TODO: implement default stream semantics +#[cfg_attr(not(test), no_mangle)] +pub extern "C" fn cuMemsetD32_v2_ptds( + dstDevice: CUdeviceptr, + ui: ::std::os::raw::c_uint, + N: usize, +) -> CUresult { + r#impl::memory::set_d32_v2(dstDevice.decuda(), ui, N).encuda() +} + #[cfg_attr(not(test), no_mangle)] pub extern "C" fn cuMemsetD2D8_v2( dstDevice: CUdeviceptr, @@ -3322,6 +3363,12 @@ pub extern "C" fn cuStreamGetCtx(hStream: CUstream, pctx: *mut CUcontext) -> CUr r#impl::stream::get_ctx(hStream.decuda(), pctx.decuda()).encuda() } +// TODO: implement default stream semantics +#[cfg_attr(not(test), no_mangle)] +pub extern "C" fn cuStreamGetCtx_ptsz(hStream: CUstream, pctx: *mut CUcontext) -> CUresult { + r#impl::stream::get_ctx(hStream.decuda(), pctx.decuda()).encuda() +} + #[cfg_attr(not(test), no_mangle)] pub extern "C" fn cuStreamWaitEvent( hStream: CUstream, @@ -3630,6 +3677,37 @@ pub extern "C" fn cuLaunchKernel( .encuda() } +// TODO: implement default stream semantics +#[cfg_attr(not(test), no_mangle)] +pub extern "C" fn cuLaunchKernel_ptsz( + f: CUfunction, + gridDimX: ::std::os::raw::c_uint, + gridDimY: ::std::os::raw::c_uint, + gridDimZ: ::std::os::raw::c_uint, + blockDimX: ::std::os::raw::c_uint, + blockDimY: ::std::os::raw::c_uint, + blockDimZ: ::std::os::raw::c_uint, + sharedMemBytes: ::std::os::raw::c_uint, + hStream: CUstream, + kernelParams: *mut *mut ::std::os::raw::c_void, + extra: *mut *mut ::std::os::raw::c_void, +) -> CUresult { + r#impl::function::launch_kernel( + f.decuda(), + gridDimX, + gridDimY, + gridDimZ, + blockDimX, + blockDimY, + blockDimZ, + sharedMemBytes, + hStream.decuda(), + kernelParams, + extra, + ) + .encuda() +} + #[cfg_attr(not(test), no_mangle)] pub extern "C" fn cuLaunchCooperativeKernel( f: CUfunction, diff --git a/zluda/src/impl/module.rs b/zluda/src/impl/module.rs index bdfcd86..98580f8 100644 --- a/zluda/src/impl/module.rs +++ b/zluda/src/impl/module.rs @@ -87,7 +87,7 @@ impl SpirvModule { }; let l0_module = match self.should_link_ptx_impl { None => { - l0::Module::build_spirv(ctx, dev, byte_il, Some(self.build_options.as_c_str())).0 + l0::Module::build_spirv(ctx, dev, byte_il, Some(self.build_options.as_c_str())) } Some(ptx_impl) => { l0::Module::build_link_spirv(