From 4530ed3603cbe7f7848cd8f8fc7f17279d903556 Mon Sep 17 00:00:00 2001 From: madrocketsci Date: Mon, 13 Apr 2026 12:16:10 -0400 Subject: [PATCH] integer vector types --- build_linux64/libamsculib3.linux64.a | Bin 836798 -> 898946 bytes build_linux64/objstore/cuvec2i.o | Bin 0 -> 17392 bytes build_linux64/objstore/cuvec3i.o | Bin 0 -> 19960 bytes build_linux64/objstore/cuvec4i.o | Bin 0 -> 21720 bytes include/amsculib3/math/cuvec2i.hpp | 26 +++ include/amsculib3/math/cuvec3i.hpp | 27 +++ include/amsculib3/math/cuvec4i.hpp | 28 +++ src/amsculib3/math/cuvec2i.cu | 167 +++++++++++++++ src/amsculib3/math/cuvec3i.cu | 183 ++++++++++++++++ src/amsculib3/math/cuvec4i.cu | 200 ++++++++++++++++++ test_scripts/cuvec_codegen.py | 36 ++-- test_scripts/cuvec_codegen/cuvec2_codegen1.cu | 16 +- .../cuvec_codegen/cuvec2_codegen1.hpp | 20 +- .../cuvec_codegen/cuvec2i_codegen1.cu | 16 +- .../cuvec_codegen/cuvec2i_codegen1.hpp | 20 +- test_scripts/cuvec_codegen/cuvec3_codegen1.cu | 16 +- .../cuvec_codegen/cuvec3_codegen1.hpp | 20 +- .../cuvec_codegen/cuvec3i_codegen1.cu | 16 +- .../cuvec_codegen/cuvec3i_codegen1.hpp | 20 +- test_scripts/cuvec_codegen/cuvec4_codegen1.cu | 16 +- .../cuvec_codegen/cuvec4_codegen1.hpp | 20 +- .../cuvec_codegen/cuvec4i_codegen1.cu | 16 +- .../cuvec_codegen/cuvec4i_codegen1.hpp | 20 +- 23 files changed, 757 insertions(+), 126 deletions(-) create mode 100644 build_linux64/objstore/cuvec2i.o create mode 100644 build_linux64/objstore/cuvec3i.o create mode 100644 build_linux64/objstore/cuvec4i.o create mode 100644 src/amsculib3/math/cuvec2i.cu create mode 100644 src/amsculib3/math/cuvec3i.cu create mode 100644 src/amsculib3/math/cuvec4i.cu diff --git a/build_linux64/libamsculib3.linux64.a b/build_linux64/libamsculib3.linux64.a index 1b5956fb5789a2fa05119d2f5033b4cc1c526d69..534f417443cd74888ce66c6d2ab424749c2318a5 100644 GIT binary patch delta 39709 zcmeF44|r77o%es|{+~N{=H8h}G6RIn%^x5HNHUX<1XQ>YkZ7$&@DDXu21Eplh?OGU zT4z*R!s513Y$xto9W7h3@479ht!#H+cT~2dyKXn_x)R&2J76nx*VaT4s0H$V&z+eN z{@Akn)V}*X?>o|G(*2YGge-b*8aFl56 z8~^p?nJb9?drR`SZpbb^N*}s>@*uGfTT(wDc0o(>LJOW-PMGywx}c?J9SI+{BtIGqpJ7<>rMWb zj?5VF)9-gl{`LxeC9~-xU6N<-V8uG}{o$5n+v5+$pB256uHU|_!@GT;US(_2v9za5 zh_6XM`KTwfy|Y1rX?el)!-7d%GcTTgGO3pFd?aF{3AScCYafTr3#)W2rAD@Ac`e)R zwh27r&dW>F_GdMP$445fYt}4S^szOx?8d8Fm*0NZ&7WDh{Pwlnvgl)1Ene1o^Jg~R zax)4(**0NP{G5_2H=bMaK-&ZnCAYUXgwXQ)KCguf2NAeeouIC{{Yt(G&z-$y<@Nbp zFI%;I!-_TQ?*8=h^~-K>+<+5%VM*d8{^v_s{1K= zOA(q1-G#N=T~6z0cfsU#7fdR;%iZg)Sn2kBIyN?4=(MJMr|~+`X=&%PTE5rHkzeRE zoFDXA)ug6#y5gKpLv2sCm7m+Y=kys09%w5U&}FgK(Jsp$2)b;%)>P~>WVfBS&sNQB zp4rm8{B)l+t>JwZZ)%#=(ss*C-e=ptT=$0CW6f@lg>holI2^LEN55*EFu;Oi;P9VqGHjoc(CYLyw5LLZmhaF{mq{$3L07Yq&iiO zwKmO)Z*SPyAndw0ZNypK^H=?wl3~U+WvjYP*t$@9c~wtZ|954+w#q#9;wtksFvmZV7`C>Mr=VYI` zHI3Q2c5LPDZVDTB-PFGHa^|ka)-{#aRm^r{A5zz<^ffPN zWlisy_kNsKneSwAD8(c0cwn#Ab<@l5nH3)(TgBq<*m}IsMe57x$5%<=PZbKQFBcof zi*E~?-zr1dt>JAqppVM7-XNWsuReNg4*Gtf#McYaf5%_${N;7uWn@wRc=s4dAjTX- zYeV^TO=x@1w+agHeBFj7s^4(S zrVaJ!Yd)`)CEX!)!E7=olVMy=cjPbbE6MUtCG%=hv}&4zLGo%cr;=h^O%1oNeVj_% zFu*30`8Wmg38{kWS6+WP`K!sCM|#8USI#598yd;WF~TiwT(*F`@+2}BlH9QV%7tW2 zUw>u!LMs1QHMZQ?!JEKLRDk*}cWX`e3@hQAW~|wCl$kYnhTX-?`J|xsiT1p>O?8Xo zKja%}*hq=SE~g+5{8!J9FW5MrbfJ>WxjbVo`559B)1n)_%m<^#yqA1(B|UU=dVRZ4 z(%PWZzvo-FY~99X=%Ed3ZeG^Zw(N}2e>05)&mH>H4gaLBOg1*HU-J%`x9Q`}Y9fD) zj3$vWQ#ar8`2CZ~zm|Ac-a)c9k^Gmb$W{Dz@a`0i&*G~B*s6$K?E{=r)lUH=LND-f99@JePSI%@@a8rE!j zL6D7QWPXw?^Ck+LV|nS)WWOmGNP=4%%!C9j=#&OXqX?nfqYQ} zqPfj5%x{vfVP5RSJLNfHd2-T(AK2WYm?v;jU`b6058FGZCS0EO643Q`|c@@NLZu%ocDtY8hM5=h6 zqas$#O-~?F!%c#Q*d%URg$OQ63cP~I6rLCNAXdvwy@*WZkrRj{>v(3Dj_5RQ`x`{2 zb5o@ku^HTSKO*%!@;V|7Jg?C}tdX0xAQIz|A0slJM-C%00g)SaN<_h0p(=}n?ZdGt zwL-}Ggb=tVLKCAA3jAk;)Y_yFU_J zlWHi8NVpk^W}ZMK+=N6EPar}&Gaz*2XFuv~@zr3qu`WIw(99>*7-T9EEF<#zIwMC{RPm z@D!xcgW<`8SP-9#Z6;w|Txw2GxCZGOG^ZNt;!^c03s=>UY*1yZv|+*>tP!14XVD11 zFq{#5Gzy)U#Xp3t7U`*2>aa|!rO=ncCmwK_n{-c1emwL*kHXSI{nzWCd37^{x?3 zuehcsGomfj=*ef))Cqg9NF+sQ?`Dxu0|l3gnrc(yDf@^`=!NRdicKNKq0ESx8=+K` zOkB!Y`%RO!jEtDhL5%2>*wvX+>4cpm$s$9;e&1D76*L&t7(=NMvpm2e)O77T`T%Vo zR#PV_8zt``Rrzt{zbl#=8YjKpHOvUdqSA=Ii$cb6sv6NV^w5x+d7oaXz|h3n2fN5$ zOR>#9>|}lFJ=OjX+Pg#>8W(OR8X6Llbt+EVcbA&8v9=E`J)k=OAsv}3MR8~;@s=)8 zTtY)isxvTqKSr>FONXX7r|3vI8RDV)h4ylBHtx_N$e9vMDIzy^y`^SPW%ouSUVRCR>BjT5l+5~#=ZWOnu3`0ir!qtlvbr5~ z2dtWcfyuvSP=46ugMlHkKPZ6cD%j~CwxVp*mj(dw}c~JDV$?f3XCyh_D-9gdN*tzq$5v|%1x?%_aLj%_4RD`u!%E( zMvMNt_pq4ezq^e+Ec&-lFKg^Nnql#XN5lr2n#MIDs_B6w>p7w-r)Wtz;p%S`?Jk=Z zlT_Q~NeL{&UVb$i5w%vb#q3h0Lx`#3K8J>dVf7m$xK>HS`6aQo3xoGHc9*E#-AE2v zaKfhDil>!Gp)S?anxSDy@!U-jv9ary=+qRt&Gg?o$YP3;`~_MriUTZmM70xAm8o}P zFUlc)v=Ul+K+R-?N>NkyfO&X3;VmXRP%t7BAwK1|t zLCR1`s#{w5OEqznUP(#SgIx)!B-JQwkExjvwqYq5ZBoc;ld24(R(k57nR%1;&W-pC zTZ;Mg{RsWm{=?LGR26b!yJTAXhiJEM9ixbft9*8A2U$l+nTabE|6PQ9q)foV*QAk( zFeR^KkKl$9NeLXoA(*6;4wGUZqvJB}FT(e%uF<7-*N~c}nBm=rt@TZk<;^M6wWEc~iq$DH+$x)rKsKOQ$Q$}YB_tuFV z;n*xi=wpk2j63qzq}-LLp#BOnNU1iogdp^@N4hUjD{*&3X;7B3a>(ps&QRDnMIujm zH!*2Q-W?^L@TP7F2Wm{`ukwLWP>TA48O) z?(dgl9(8{f6|32U4#gDDmL6F`WvX(RS?_1qBHV2n?m`UL)GsrO0=A>9coBQlREKw! zos#?I(3T7d8Ce>Tn@sc8ZYEel>?E%EB)%yb}nf-u{vX+2L|N1BkF;KUzf5{Nb3|ro<6Zg zLBht@fgq%)>Gdwkh+>b_xp`RKcY+T)$zCZ$y%N3yDB-c5%SfSqTO->P+Nfkl2PIJM zl+$C2BCc*7hl8))!@|QSDL2A!qhUlY-_cKI7EdsHv!s6JSv7HrdZHPUHc7Zn`h;hL zONYrhN&``e_K?)Z?_i8yQidwT44F<+4Pxt$MXf9uJn)+-Yh>gTrelnZB%=!*<2fNy zhUu9QQQ8Ab34?k(vs<;5kT^sk?Jx#0aZ4@9wu0|QBc-CQE?!Lu4=nqslNk-xQZ1vU zUdGLh7Hf_;6HtQlvwN0V%6lBA`Kn=5QwmuN%K+sCRi7;s`I=|GA#NJ@%{*!-^PA|)s^ z(m|3VL`NK4dg$-zbn}GzByN0?cG;IS@L90N2qP0op*z}W+|k1m{0_~fVf`*UdJ5`K6$I~_s9^=LL@>nGge$JDx@>msuhFXD_{uxhK z^Wb)#eT2tqcz|l8da2x@=p>#z!E;M^Y%+p|1-PC>UNA|~Dct%cp5M%4wLI9*lk<6O zDuRXzi5I_?r|WpIf@iPfv1vT$IQfJ z&+_zhJg1Qd@ABkc9*ZHU3@NY5ev;7MFghiSzLfgq%1Y&@vOml3iiSRzytKCJwp)qT z3;axEx~A6gNigQEhZTBRDpxFL{KQSQVpf{1+x*;YMJ;ui3Ysb=vr>7`q3B?W|Dg0De_6*#dH<5xeCUQ19C{xSJ@IH+GJ(A<6ChAaac>oR31Z;CjZaSsdm9@mb zm7*WRelE$(z$RBtCBl^zeVAJk^}G_dCzG#pgWZ71NOW43QA3rO!aeZh)LKSa$(Lf$ zPhmF|%~{@=XgT^xs<&ribGI)~B3YN4xnvp|;d8`>OZkbaPqZN4?aVJAKbqsMjb#!1 zfFETm-+3IihOnGkLeluw4li$E6TdkrdcC`cHpH6wc3^J0> z>qz=KuB|wpTVrvmsEDs*6VqE;)td408|Wt15Wne7Cd#!GjMq{_(@n>i5KRxXsuwBt znoUoWnGUw8&2lu>$0&Vuo4Pw0ZHp)q@au%nj3|P65rv|$h{D`R{8yejqfVI;zo~@E zn&v3%YJrtfRWqKh=T7=`d=23Vb1oiw7E|rC=IbcfQbx0yrZwL{Rjt#@d6Ahy3blV^HvB(q$;xMAhShp=}Zp5FUI*gLl* zC%dbBm}pzMW+S#OmuF(nSozJx{+all?au}NTa(AXLoj0$s)*ewOl-YVh&J6RDCWaB z#`PGQ|BQUvLo6EmfMLcfjEi025gw4ts|5<(hs@7Rhiqf3B=*5~yhE7d z+NO>P*DBvpapmF1pJS@IK%Rz!z*W{r()+6UxKL|e>eh3qTTh{y$Ax(?^D}#?JQWRR zDEL+LSHX{|liA#Lw=VcIp~P25{$9b1Q;GQngD#4;e8H&pKSJhBLYetB{7PCb%(-^m zEw|R+HBpF8x>Oa+rwulBTJy!Y=%)$hCc^qZSf~DZT~Ci(uZ3^Cpj&$x{+erUU3U3Z z*DhFeRr{4^{6_kr|B$8E-=Nihl=VAr$rb6kKk`@~&v&Q;q>+Q$tCG;zAGx}SO3@uj!FLP0ngav$`x9$6mzF@cCkD?clFqMuDR5&p%kWu z)G6@tMCaYiA8Gen`LZ=Ab5EgcBi~?w-})jyiU-7ne*32Mi>uYrEAZUwek_E03Wdko z@I#IJ8Eea(Wh+66BSN}!o!XI(w6nVOsWqCCZoU~0;=wvyEeUu5arCxtR$8%y`Rpp+$hP=f!+Wx4C(_C#MP*fVF@|FmF_6pp^Q z#S;a4?^*U81-m5{OXk;qk+)Byp0hH(S=d06&fMTxwmtWovgF?*JBY5!w;PWLg;M|a zA@=tR_Tq8Kh4MbMlh<>$eNw@`%WVe!QY@6eG;dGR*%`|U8)OR`@Jon)K3lL4oMnHc zV9yoo?wy75y9;(3YyK(Pe^9U|TzeA#7v%HlKMNU|!UjdVh7RRNmOIN{Td>=O4tH-H z)YVb2CkpmrJ*x}$%-HfiY|vfE$cgFopVUj1;?MEIxhmcsFBVEPo@M{Hg1z%B`>}$3 z{aN;a#Lu}fMM@9c!pejj^T|JxzPgH)jsH6ot{gr8{Km+C-jvdf$QNWeJ#;g!^eJ*W zfwX!#PJJa==@&I=Vz_z5&GBW8h0lzZB4PCHCs~Z8>yF6F(l=bjfWV9s^t83fWu*R} z;jz>LA(I|DB2TKm01PXg+^0;uz%Lr6v42hf)i-2c`i*j-F+KY>W;}mt4ReG`-YuL! zeht`I`2TJE6E9rHXWKtL)}Bs|m%_=ho4q?Wdu(;M*7)a~>ORE&%!k;A&$Xu)Jnadv zndxO$uw|?z{jDltd-{|AAa}5VbnzU9kZD0De7CkdFf8Yr*om+WRM_7A=m88!It6>9Z<-aq&Zj!)YgSVj7UJ6J@>MAMy5Y4zzhUXv}su1Tk+vB}#% zcpTpav(js?VsnI!ny*f0lhZG(N8wGg(rU`lh;kdXJpW@pPj*r6Bqg2AMlq#rE%h1HE)ndhKF1S#_t1fCxG2CBkAEsKrv>z0yLqosQ>4K- z>ZOSbw5{_r;88A(Wk#L!eSf8mPyhSZRjr~(Ot}T8om23|r?tq+^MYgU1|Y)5R?Pvb z_G>TSdgQu}ttS_EZ2gzRyf5-R*F{?wAO`U4(PN*Jwl0v|?Ge|uE31GsgZ0l2L^18}?P!T`6w9t*g=HWqMu=C=WE zlZ+A%@Vzwv+`#IY3yjGC-0q-K4K&~FdBAO19&r1@nSk5Zeiz_2KM%Mu6+cM&IOC@O zHpVSI^8@AM|I+}s7r>{vfZGf(V0pmp5ky?T?LQH50k=y>0k=nCasjvZ5aEEE&3uTR zpKw!ag@B@9q71|hv|Dvn&QcQp3nO0XpwrI5+(M}ivN`02EyczU!`!mJ9dnyl*T&98 z+8tc_V1Ai_!PzL*f7d3>tK?u=wBbC|*S3I1%@@u=q`w zE`W%m9a7a{<3XMOGk-!_{0_D{1By$3Yfvvseu9BOOOp7Z9NhUaIFN3qF}Qyi#o%`T zhX1AC6O~4^Efy|!$)?OG3^#1kJp9`)w(*QMs^v6XavlmRo!rs^K2Ei1dl4Z}6QB)1 zNeuZ;19Io0aQ26xaKqAA6wcYpQ8-jtVBvb`soyaBC?)aF|A?N#{=vdkZuVL1#lfy; z1+WV9zXqBa2%C}Of2?sN&L|LZfKX1-hOd!1+otTnr9~jlFcLJ7>U?)f<`C^&2IK}D zT(%t0jSIl(=Kg0m_;yr%=@bW~ydcW}z$LqM^Ldwm^XfL+{T3joEbtTE_k1Iz{J!Uj z18^1%bc2-}QnPh700FN>%K^aoX18_&S}byK?}xJ|By1{zk0oj0Qm?*=ETDaN76>|r zRqFdYT|8$<0zLgr04i*4`Q4S>$U$r@zUDV>B8 zZ;n}Nj$JX46o7H1h+YN|rUT^4lpL#4k=}jMY(FI(aVRkRR-0ft%wnQ`_hI(bu!%z+ zhysyIfwD2oR?u5Mv(;g-gK8Fdk*R{0G|UcG_~O9#N6k;V23X)&Mq0!vLun9TTpz1?nsThOI|+{M zEfPa|7U<9(jJdl&Itl2>SDql>?41^c!dVUw_OPce0Y4=G@YJ*cR^`#MEbh_MS@vXs z?DQLHbfamcHTJIPH}XW;Q6N?f7o{IUhbqy7-722UX|PbWlH}lG)G#tIkcEBBI+_!I z^ay%YyA?nZ&_XdIR;gl-=y0;o7M{|?{X`rrM8;F#3S9tk*vxSLsP-z$s z2kr9Vl)Pwjl*0OPTo z-~YzY$mW+z7=ePy5-9`*cKb;+^Cs(_OX40@#c8@KKZ4y8G3)?^OKN`o&@L7>SIcym z1Y3S78ZmK-dpV8>YA@#T9|Jw{^NiY#JL6b_S(s$a;}mMACp47Nw)_3Zz;@VO<`VXv zXy694)$Gz|R61n?T=i(1w4f&otaWyJ7g+}~v^0QmZEO|TZte=7<2 zs;KmDIj&TBFgDG^ze5Sp&JO0(K)Q$RCM69f*Y5HtX4!fk>d_b>v@-bq4> zNc$aO5k@oxcOK@Zi*QF0{zwQN%?L3PZ{>)is&VXKap2y@t@6CvX!C1NHqW z0f*m+u|_;Dgv8Y}VYfq;$C1dP@g4_|Jc5`@-_G^u3qg4amM<$HH|%qWH_MMTE>qxj zK7EVB{x-DG0gvmjnh;VvGTU(c;Bx(P=Z#EbOi@(6QDNWA>=vO92Jmlw-gY2yp=uOIf)W#(CC$i0Z#a7yMsY3~yeRl!KB?(xm?nRjQ*nBTlRM^E|8B1wxc zk>T*<6Ko(FF;sksXlj2_wiFE$ahsGC&@lk0O)zl~m%3+5)*EWcsgh-|lo!*$OO<`p zoxqbN&bT+`lI6gRHPs<0V*vOH;9sk8TgYH`U>?9*Gvf=1nt&#zOy;PR^@+VcAhNci zWPG+gLaNlJOMOYSLK{+@_c1-pQAoAL0!G`zIP*FmM@>s)At#pf4XK?jRGjsT18R+k z;}-_hs)J!CQ0GD2Eryu}^NWM< zLORLN&V3o5p227$A`-=-R8Yo;Y1Le!o!CoA#aF`~AKHz5f&ga{d8A$V&q%7FpJte{ z2S?hkK7j8Wm-;;}1NzQ@2vYYRe#;;(llDJzCjJ#?D; z?NLgpTM4L+Pf7@Fy0S2l2SrH!-a~DmP&%Xp3n|Q zZciOiD{Z%?5i1fS!JV$I6zK}+hEF*9VRr2Z&ohROX7L|EOj)M-j~8Sbrpz4+F&G%RWQZl}bXY~L_ z_ahqCTYF^DE*-CBybE8)e9$(ny%;wdRY@z2`TfT=7WMb*u^o20Di}}QzS*=P}GDJBglucOF zLrq@eSi%Y&%P9dnq0&$ZH3pTG2NOwVy^YIwZv|#Jn=C1(uEGTMpgrhVA@87C<^qdRf_t+p2?jokF@ynu2QxnomqDer zWr#3Kl?E8tJ?jN>xVFGWdma`%#RfZ)ddbZe} zkVg}wS`Vs~8a+laGC`^JSVyz&-QDOxzG8aDHr<|LB>J1cq|CUk0v0HApLXhyCiC=;tQKt)1Y zEn^9dGUF@7Esey#)NnJiN1ow=GkjrrBCJw_XeX)>dzfGP|r5{$FxtinC@TF&Oue1 zg1Shn%tm+A;Fqd^w1>WJfrGyRRIRo2sLh8>ny}EO%OGll!UuPznWevS@&YDllV0j0 zP1r>&gQ;v(8}Ld?O#inInB)o`oWeI$^c5Z(!_xnhr)|tX3NZcKJYC6ypY!4;c&rLR z9;UB1%AK(=JyK&}`aIh{iQBN96up?oCi9K+F#VN0J%tBf;@L}itdbdzj$L0AA{5CsE)LV($ zW$_4CM--8I?;&M2Wh&M;oypXfB*dBH6!7pt)Plp2r%obgrkG7i<(5Nvs(zL|jk3)g znlCW*;Nhq9j^KUCp?YZ+i1|I(Y7ACiV*k3SmJ{*iXd}n!fyC#bdIRZ)d3q)yb-aq# zk&-64?eLKRpq}@C#l>CPLA0Znfm!Fid*HO`X^J^ z#XM|+MpHD06#C*!q#ietOVnE$xgW2VC+Z(-%M3^!$=D~GGphl~ZE1LcW0R8M5t)W&1Ue+N1pubNT1JHj(5W?{~3O3bIL-S34Apoa; ztVUbFar)PR(;u$Unx*>KO+QN4OwtY{!|_K1@qS@OEcSp9ZC<1*=02eASF3rTK6Ygx z5|0^<(WKVLZ__B1iJ0 zu{Q-Mv7Ux2)E3mXe-ZYn?6I(K0&3kt`nlI*9k(ytcy1Kt@NW)w3=jwi|@ z?L!L8{$Q!g?0+WD?B6AHc)m>Loq+Xw;lBc_FL}&QgV|r7XZG(BDv_!1KHA{o`x{(* ze}lmB{SCs$+>qn@8-yA~@Sv=V@2|stljxXF0pnl6*h)Xg_-|wW>p6pO_?M1h@PBNy z!eNDMeyvFH=ZrkX|E96t^PXh>!klS9JaEWtHPDEkn=?F*gXe#e`NN{i^J6v2^M?V9 zu7UzVHsRVl<^Lck|3@^z{C9tm^8Zu8 z{1Q>c%*O=bR#}-t6|HGO`VNI!=LoXP{eSfat?UyQGX6B3js0K9M(I0l&>C*Mkdptc z>YLJ|)c;3#49O*0>ql6<`3WuhQFf8b{C|X}fHVIeVfQ3&)TVrd)sNh$MQ{D+dgn~! z|HK*#G>t2ODX zR;$fCW+!W~ha#Qd#&>rHoxj2#neVsS!Teju+9B`)7*+Et*S6<8wj!)wIi_5$P_Aez zLEaI8($}n07yf@*ztB_U$_ZfT|vz0|M7_uVMc-8 zPvq(SC-e0F-y^${j6x-Z9>{-wQK0v|1$uu*f!@Cwm%fv*8Ydb^xwth z$vWblq~|{E37#6Y4IKUo^N#ep_KQTv9Uz z$?usu2FbVTPb2x-ygZWMRzHg5r%z2~TKe6qSexLW<0zD0 zkSvWt`Lpx+caDMb`xcCW^5Of|h4SOE*0>Ah_kMg7%1^QMoGpSKJQK=q-$q;0OPmE=f4|{JE~rB3RLH2D~m>o_^zbZ`p!v3+T~f7j1oU+X88ox$Z@mME@S>awOh(6vbcW zko~@G3k(?k8M%BtaOKt|{I*iBP76Y1EW>0x*0{6I!^0%sBV?ylVe*mEAvH>0DQH}Tqwo1c%z&*y*6 z*t%eJ^Yqm(sbxv+Fuxew^EJkeEY4ZMi);bFze7P0v*1S745%;njP)CfJOKK46aoEn z@;@BVUwMPJ?FMuyT+JFUuQH?ad0XI!_QMGMZ3_(-)yGj3N&O4mk!7*ut zuVAs#1@$wZEP(pHjH6Nm&*;8;(8wbG;Na9g4r^ActbsXWpjv_oLaS{G2@V0{o-&JfG|6mne(Y|>; zP=Bmz=pv|phF|E|8Lnt6yl1$g{e^D+seZ02+E1QA>QD8a?RS>h=enYO;Ve>ra15zG zbOxyp_GyeO+V5N#slVJ9o?Rl z*2j`@!&hSdi9qmH`-sS)Q8Hh{Ir##(tyihiyb5k>C;Gz1RPz6vpA*HLWr%QvQ4)=9 zqKJ1ond^M65;?8R;(*OLN z@QFOnU&&|PuEW|Y#9W89)!gJdtgYcD?yz_T26x72dF--Y_Wf=HZi_ckI;JZ~PS`PlK-L~u_Q3^2D7aA9yX_>{;N4d37DSjh_0O{%{MGsZ zApF|Cymwppu64cJ@?sB(u7BHUJq7=^(>GBd$sOE+>*o$`C)T6zce!(0?zi?&IPu@u z2Q&f4#~s%OK0k*1e-*YSl$wSEoY;&^xb>Ylv*6u!c9@;x-F6mu{#p2YG>k~N4T(0M zXv41#;Z~$tc?yxmEx0%U&QEUPh88T`8*X7KqW|0{Zow}ucZut|#RdA$F@5e4moKr? zOVPt^BTQI|9&Q`p;kKbI$vFEz>Khl%YGS!l@|X3Wm^!PE-xNNv<T0;DkoGxrR~`$}VbcMdn17+AU#wLo zCargKH1LXuFZQD2+!S4N?K>7z!R^}6o?!9XY?5EI*)W?!z;}z90cZtM&`QXxD zs+s?k`k$cFx`EZynd*dIphXyx%Zohz7~w#b@E$zoV(jN0b3q5^ozDWp=bm%DdIziY zJzq;h3^?GaEC95EZXnqeab4z;MhCztUC$~{(W}bk_7v;HR%hu3T*W{Po^%r=2K}B9LN+*_qc2;7j-r*wU1* zpcjbX-e;BwH3p;tWuQVqIqqWFqF}&)O28rOeTUYA9p;*WHoDY@RR=~;de3j;tOJ(@ z)|gifoE8lj{?-h4$NQpHlng*5kaI->{YiKs&9Q5M+=$|3a=bjMWPxmgpn~h*jO(IU zh0F*Ng9tT=VkpP%#3=#3=gToj1_p89BcQCnvO7mpQl+TuV%2^AJs=PdvmFl>{i~WV zn4rt7EPLu;I4^zSxLM59lgF?!tYZnr-OXmHoqX1@A$%WO-5B9J+4iH*VPpaOtG#Hj zd3PIpcA9s0LmM$%a6D|7ce}8^O?P&K@m6OaN|803VLyUPTM=HDZLSdTM}#IrygSG4 zPeD5Zyc?STDn(oZw#J5|Q5HN(Vr4|kJ0dpFISIwr3c@K%)=Lh>Nd=_dsX&bpL0PMh zb-s-|ER_)}wSAkIqbNGt!5)M<+6$xbQkbt z9`2#tMGv(QTHtg*VX*}LYNRbDIP9glk)q1N@9$?dq9-BV4;Qtl1^(wmOVviT6;-S@ zA>uZScr2mQRaW$~*?{a5qR=VFp{Rgl1TLQ86C1yaasu{l5-L5d9deC|N&|LIfoucx zzKNrP7NnL-Vrtxr?E5$zC@B&mfNM~jdH=LQ2{)6ecO(T+vW`68hJ)Ia7#>WBAt51N zYnRb0SO+tk)T;UKbgMn1?wNthcY^#=WAKMQpqhZ>qZJudzk}=wIxJ>XBdTOnPs3Y6 zM)h8ONa$adQI*id3`5BUb_-?H=v-L34}@`9;W|!cR07>%sW)>E;?$3;AwgoF%C!-F zYG@*8a_A$tFrrVLs_)Cn)JAu4kHoGdV%+hp3_-*nmB!1%Ixh;Q4=Zm({s%2#Z)s4V*}E zAw){$5{Q#!vt0?q?xF<3)jz~s{R1fUR923|XLCvw4;OS0Kp^F0Nb6B?Ur2yZLH4_m zFBTOdxLqO&avm(oL-w#bzb_%?i-88MiAK<0XkwMeZ;6#ZSF?H(V2X%{M!8irzOC9B zhC#(mXIRPr-IGzWOLJ^K7k|Kmd4A{Cfec*?+}hoWrNq2GgYWwQ7nNbeb#r2%SW{|# z|JD%_E{3oQEznXb*E*yKIWcHDG%p5wPYMGX6A$|y;lea1JEDZE*-&Z~sYCu>uCL?0 z@*0mGbspDvIHqTlDmkW^ zQf0_I5IUBaQsLHKjvd86ID&+PO1jjo#@?q~$}8a_T4ITf+Kb&qz#Is3iE_K3S;k01tuS+3y2@^won5CW|i#8S(p9$lY=NAfLM zitfvTBljb4PRmd+ee~m}juiV^Xdy^FHN~YMg-8W9Ku)dmE?&{#;HZL9TTmJ@tzX(o z^I$Tn9XI8ER~aH?3>i9tHo2r*lQ+i-zS@`Q(h`52VPrI^QN^b>s{;$Lpt@mFr>!xIyyLY3~R|_juq6& zGsYu3w)S-BmYU_dkgSAcNbJzn1QT-5g;=ty3&}wj!et>C_Em9{KY+^|qL744Ifz0c zG8cvPXbu;Jqy(e)v?e4sstMu#)}K01GXcC!o2}}8)puqwj?ZJ|`lCPZ^)6LbX_@Eq3SH}^SO zGro#3vS2P%UZR3T$eJ#&n`X&G?ST{K}UzXW`8US#PzbI8?3R z>!S2f&YYjc)NU8UyrfL+V<-eCpn0sIaKS`e>S3YhFk{1% z1ojC~p1@_g#}%}U+v$X=Jz|E5t07o0GalxMl=246^#P;5gb{#iLmWb{c$P3}eIRti z35CT?0gG~8vEPHEzM0EfMD2Jm2n`zEAR;;0ZG$kM-1Y16q4;#Q-SEIR6qi+=aaj)ISOT?s(5eKpnCq$Ca||QBS+B&n6TRV}R}uo{xgivaqZi^zQP6 zHh1UcEO-ot5+Z`P0LY0*U}uMbQ6WBubY7%`Chrh27d?^kLfDZASBZG2Ks^{{I#06P zh~Q{~lcb6rS&!*33tq#SYVA&FWv|#w%urzSYF|6`Y<@3SimnT`F7dI)%S(beZ(_$Y zP)^%fpJ1ssq`3M8x1Ejs*l6jd3GFFY`hu56;Blb|@c*U>R;relKx5#!-0(pJb8}Ju zk{lz2DKi|sB;e>G9)@j+a)NiTL>ef0aF`LE1XSWFuQSY?xehbjEGeZ+J>~>ow;8&| zRwCkHiE68x7fs2GRA!%`Ju;q@_OFI^W|v&`u9}r`2S^=i3*PUNm91V3L7i%{P8ScC zaQU=T4#iBb9x@J>1dSe@->xZrA=l6J>Q!E!o5b`DUD-1}q{8M$B>tZIy%b?$3YnLo zXT}%AGnfnIu0jE-scOx zo-MatwZ{pTXO`a|)W)&QxaIh1016rlcj6LN>>j@noo~NoAIU`~nEfT2i+bnY@gcK& zJnQq?fn7OIBlOJK5&`#+-tFpU1~ic-#v@!igBevXYldtlqp=({C4CUklmw6ECW?p4 zLql0SYaxrFky19u65(bDWi3rwg2z%^YDaS=kj@N1%3d29_dK*GK1*P>z<4TG67sDl zvNcyNGk$-jENDU=+>gsZm_#K! z%@5*HN(}jy$P$%GPDsj%&;w=5L|;bqQt26Dpd^*n>G~#FTF(o!PNr@uhaQMZ3wj_F zFVP*bvLFYd(rVSD(!3f7g`nGtpl}%$R}JKqAc?E>^8Ge?jaqN?Xa-{VS})6p6`Qk$ zVtXl7Rujdq?oz58MWkd3uYvJG1w^If$(BU~r{XU&&lYq*fr!^k)RcyD(91+XY(QWf zf*vRvnjYPf4NWqg((kg|ks3U@GsTOST76!3Sk!X3P^b_`XBtx3dd4~VkrG>QPc*t?zjgF#6ZA-H*J3$GPJpX6Hf4X4hEdLuxif_ZQ_udno3-sdiH1r&>;? zFvY0ByDo540o5+IKh*;7JGB^?r?lhRks@rmRo+1?Ll%AzjBc@^p%_tZpJ+NVWD0C# z4$IWeLy!3#JWjQ-WT~00wW=~yqV5^w{Z5ZqwMvZU2K?who&0t6 z0+N+&r7FCYcFV}C!L1(`YGS2y_#!da>?)B0*)P{MfwnvAy&i$*{&!4VS1CGB6D4sv z@^ALo%VSkM;F_EKt2(NYEWFs`XFR8d2aoaWBRn>V2R}#h5~%|1ay2l8H+z)utSLOW zofmgEtmVn)c(R$trt%pa-_;v@H(a`WSk85>H8use-Vn-tNGF6U$%hhRQWm zq4P66&wU+8w2ha*-lTfF6C2Fp$!;7s`h&*L#)OSVKMb>@B}m~5&i&MS^3tYz+6ylW z8T+rbD{VT@3you=FAF&YS<~iOo@z+K#t9SC&&=}7YLW2}{%>5?G+i7cYr-^mjI8Ne zq#$q#Hbdkz^QP6J7){sD_FTpgjuQ>~|LDE2(=NyBOaa> z9&Lp1h+DW2su=zW7eZ|mL-E;$xluer2(?jkg;2?jq6LU6I%}iYqPn|kC7MJi-5g!J zQG_yTqX=bGA%HS!qlh<@ATZxnzEluKy`#9osCV}x!=`C-9M(8 z_q+P2`(0%e!lYa1Pq;wpxZ(<=j^_nZpA}KX#b^>%K{Hmuby7VAoz#~LI;o|YJT@&A zqge7XsV>o#NqHV+CMaz2X*MnvHcFfJ2+ zp$bvWgKb_DQjki$=1Qnub0t)-xe}_^1TLX^4HBwpI?B2ds)77Y|4Fz+f$r%aghlF$ zjJ;uTA=USp|J!*X)!!6_R9tK{=+7&uP85|?Z}{^{sssKlo^`T$bI?^%Ee~=f6>S#H zWx+OcY0&Y6h4E{~3;v$GpbD!|K~+y)P}SoKs?5>V%jR(DLi3%{E6le`mq=lO-IFZA zueDrMWvgRERXwh#%2v6Ysu(`1yy(e`s-D%&iRx)l)o$G!uezeDEn=wc8Ql!4F?A`w zQJP6c=4P?ZEOGaN=dRPjt<8{Fm8kA(a!S;^#%i_6O8OS0P4JH_=6C7g7Wz-V_Mc^FHtM zy0ha(hv3}x!sh#a-{*bb^Pck_cg}GizdsKnbfp+=rcM2pGy+(r8WDiP1~0g4FPUob zPs3zDEGq{=`xOf(!R=QrTGl^%JgnAQ*uC$bdeG9n&&^jX+dCRJU$LNaHY2}ik%^k!rDm#86hmG zO9*S0mKQuZ!L$pGUMbw#7a(N= zNe_x@R3X=47wKgVF*EoZ0;V|WA=2Sm7!#fqkZUJ>#h~9I zJqUNoJYhVHuRNlPF{;2*2K1+-FOt5Gw}SMFbV-j${R`>q6faLC2bB4t@zb%A`T@|> zpEk?vjRaxc5D?Y^LRjqMi^9S(dDsgef28~m4zVYe59K+&=lTQY5Hu9WXg8*_P` z1J`>($+SP}gK^TgTRqG_1g+dKAl zzx}w5#Tsp)6N0Z2KfH3XSUP<^Py9w6Jps#*Cu9HfLO-^xVk1uj{zz%7Bev%WEDwJ! z77JWdf-IJDdD-(NAggA$hsrme12>gAHo|RD)*uIVSRB`~r-XqDaax?%LZ1{Qwk9~r z^^4o43;>@6)+`8Y&46tR0$XhNveRgw5_7WUO=mWqP6T@F=3$M59?U1Dk>nWRZRtCwGL4MS*?>{J*%}6*0EZX zkqoPK7lz;lDzg#6YMtOdDK8^BSgRnagSC!Cjj&cQs)M!8Ma2%vg0)gHAxc=Qqe+w{ RtX14B%7Z#sYch_o)*meu1xx?{ delta 2814 zcmZpAZ?><~c!FiUiG_iQ0th5Xmfe{xj ztz%#$3fB{>5K|qKUKRsUxK)mU2^X%%B|#_gdYmq!j|1vaf`%nv6$5>7J-sbNv#*{^ zAG6+@#()pAF}-3S2G#@3uPL^}^nAI}?i2=m*18F(X?s_78 z&%J081Fdj9tqdm8Tpdz;Ho0-S@#bJ2)i5S~{mG3v>XRqrus5&BXH)c^nh diff --git a/build_linux64/objstore/cuvec2i.o b/build_linux64/objstore/cuvec2i.o new file mode 100644 index 0000000000000000000000000000000000000000..89cffc592994519e8c1df6c188d65a8dc47847d5 GIT binary patch literal 17392 zcmc(G3wRXQmFBrs)m7cCR#!_6r~ye;dLo3>FSQY+iA& zwC7Y;NljZ~@MLy<_1$yN{qH&F-h1w=>UQZ+Q_DtCO{a*?QG(TER`2Fl6CA;stA6Z?t z`Pl!_#BY?;(`nHG}o{*56{)^`Tod`G@u;y@~VnMfc&c*VFqF zI(?)6n_rd4zo%s-m>i$68_SQ2$E4%ju}#A1&M_LV=0`UPhu5)Yx#`u@^X2PzZ9aDC zuI6L^xD$s{^WcZ|@>hi759P1ifrERYQ=a}or~DH>w6}iG>9Lf_AS+t^|F-@|pHjd6 zxKJ#Aw+W8%^~X093&;w`8U<{M)`XN6X#M+S{iOfRt6w(sk^S`gH`WiIx_(1n>&XJu zC&%Yj+W+b6H~O!6;->ZE&|jvZZ$A8Mb8`4;;5+5t;_}A9PYz!5;WOsYUTyy1ypS;3 zn~z=7r-p=cac+OOpYj@={NpnqKY=GApLg*OS5Zo+ZD zc6$47?)obR4mi7vrak}%j0G_*b5ahVSWhL~-2zESqgiT72HDo8Wd6BB$&Im!6nYu_ z4U8gFFi|H#KLktx2tp=Y=7Zp;?b1ECe{g7EDAERQGpsn+V}>aR5AN^n7zKClh`Xz| zv(EyT^V<7*2ZtiPLsj4&ijEGs!}o2ijr0t*4|ha*_ciU~?tUkUOy5Y0uf4uNeTuLRMy4Lm*FPs1XOX0cL`pN6btM#u+u$K9L3 zUIcO@SUo*0jbPWP?a*elxzihNf(%m$$eY0A>2KKtjtZi=TZ&L|q?ua4WYkAC}G zK~}mpLpBk+(*HPeNaV_{(JM?Y2GM?s$s0h#+Va#!zY66${jZZ1Gg=W9A88`7Z`0P(Vs(I?cu7raJaqBUmfW1b%uT3@DZ5F4z~}7JKL-4Y9sz=*ta_1^@q{n zf!grqt-G3bY^`sp2B+{-DDc%-;D3iZqMf~M-F@v33^qLsktMYcL*LN;foP=TD6}tL zeH7NW_YWU~{oAT~BfX;#@dP>#LdUi$9NH;pmyuuZ@9#PQ5pPw0*Ka|--?~gL07vZ`8H%Q71ANunz0y1 zwOt$4y&%ZN8Fr;*FW7k;<=Zj}#X_)0z>b5uJ0sg(2y#0(Xo&B$vWkw84!~BKr6>&w z$em!uq3zV_?1U^6PM-gkQS9mI>V#|oGyf~2Sm}Qavh$J1f0iN01voBkVDWT~{taic zg+cy2IOM&MBj?((#3Hb_C+)7b6?qzcXoM_=YzoTT)9Akk1cIzYGNAzKi$=P;P0-K{Mco-7U$JG#45UBu3&<1@^ByvVR2GzwD$lpvdg%_1ZLDs?3!4{IfM6dSzmHNp5lV6e{q3S z#$_=Py>K7)Gp_)_|Eb9pBzOtYQtV+~(HLYM#wf$9SGEw2)R`(3cqJx+s~^#g=FSmwgM z<$1BJ4^y8Db_yOd)HDl)LXiznaD|X%gH4E47ei?^6hjU&IRVTBD1%H+6=td^gAANB zIc=EtQ{Io!Hy1%Rv22|$Gqqooa!?)DS|T?%;22H_^}y+*D*WFd_T0oNw)JJ z(22!ZP=c`(pG9|ZJ6}hr46`dSmSbF12F?mhD>33B=k#FSizyC<(~l}0!p@(dN&um1 zj5TKTBRp~K(9+5cyMA9_Sz&6MJoHm}bFf%c6NmP5Uj~RBqEi|^$haaEPQxh$%)(U} zV&_8^OcUINkRUo>lAn~+d^p8}Nko)?Gc;qi=HEFMpD}MT!?WL$V%OC(JOD1-Ge>wp z1CYIkH%qwEr`1a;pbv_Uh$=Y6AjD_ni5Z9~Amc1fsPD;ebY?~lzKaJBw4T~NA;EPO z1P9p6SnUT=EFUf?X2wuxM$XN`U8QQ*Z>{Ixi76@eF(eeoxB!KAobw;>vKiYZycwEe zfIX35#(D~zwrfy0V~xXOlT!RLyq1sK8mqq$0(%+wj$~pb{jm=v^;d9aw|R1D4nCAj zPV(jW7BJM}_7a(p*?D@H*2JrzX zBU3ubw(qf7bMR41l`>AjL{I3FR63o&N;5PRRJaqjn2;4WTs8E55V9||;44_sh{R%<_daY{z_W1D&VfmfK%Y(G%P9_Q^x z;RN%BuEd$&W#Li5EPVs#1228!92>hNiPNw<7jWKtc{QZMP7umMmKew4?DQLoi#OMT z$|v@S&74o-Uk$<(Hzobq3{F!pv&eOqd~AW3vy7HxhNzU zrgw4YILf{)1TBJBP%?rRR+>L#l<*xPiy$~H zkAPri&KY2D_nwawtT0%A2zQV^Y)ATGy9DVRY*8Z&jUCvhe5AZA^~ z4#*%bo&r(53fD}yumBHMs~H;9(4>?ApDp86)V7Waj*NupK;)sq>+l9H3hOK4QQmqU zPT-Pn#kV5Cr#NbZLXhV(FjHl-Ee{G@D<=e{U=A-Kl+DL%+he$>u1^5=%pr2)Z0B3J z6n_WeU~4{{%?BHZMK*IQ$DL(QK61NMfa^km3nn39a?0Z@IGGci2A*;mqfD4I9r5G% zaX5Y#Oof^9SyQ1M3B(z?OL1k&zkn0C;xjYHS++Dm3+LhMfcDJNdEA&p>HJxf&mx^q zpj|2PZV-H;<>;76KwFY{kvT5M*$!N0W?X)_N5feYb3m3F6x)7)6`G|WIZBP0&YGM@ z|IeA$&q(jI+O0xDaGJMsPRlr7EMi8* zBY@*#Q1xaA;ygblY(Fw3y?ULtn-hW)P6)W~L#xH{a$InpRn4HvTpcVr*$k~HHXR&O zc^>Dr6+44x?M3eKk6~hl;R?fTIQQgPkQ4Z7btVDQ@JmwbG>j>65Jm-@Bjeml*?XoS zcm>WW0=xo39W7nlwuCt3^Kp=aQZ_1YQ5*@d5wXiL$IQ&Va?mz2b4b|~w7ty5m*BG) z=Mw9)_5g0={`iZMDmwW|aN>0UT02?>CRN1K&=!->R*1|jBH~b=`#5e!bg?;Zj-AF; zjSV&j`AK^Sx2PQ4!b3GPfLpmgxaW#AF%8%uL1^xWXL3z5yHV7!2e`#@3+zK@-Ra`YA`?PTZ3OGTv)93mwT2Tu7eLduiG z`CGjB%V9)5apGZDX$Xw78Rc%u7JqTDfKd+{*%d`p!MJxM<@fj3-8|t(9h8@#6v5~>2VhyQ7?*cTF4b0!DV@) zWqd+dV^wKjj-i$Ecs{H|C^tI@%7qx1Dhtt)<`K^d8V843iK}uih$tgt$w&}o%1DMB zsQf)3w@*|+jJ1>j9Yo3&tmgK3C91Tn1i)FPJWfcfn~Z3Fe3SC3 z2Om9@RS5>3LOzDG|Jjvg3=%?Sj42ObE%|{2c@o==UJ~4DHEL^){1~(46V8XHfg9XiR=DJ6s1A|?Cy=tA?YjqDr2KGgV)WM!`wcqNFbaV_v2M5E0U0;Z* zH32uwY}W4ARClDiYhSM#P~F2-e)VRzMEYIZty~m6Yh2XPZPAR7dR?U6uKKHM);A3a_2_YRvhR&30pgBH#U9?ugBQ^)p*CWo+;k->$0=fPaH;^*T0Eg;!9wv$23Tu-dzEL&Mr1n=Y}g!LGLMC{{k)+pgj@mKUx6$ln~8*#4Zmhtctvytnl;)gYk4Y1uJsvgFkXJ{^oM^B zc2hYB&*Q%1$EypUAM*449W1Y^gDHUyw$#=J(8X+PamKF&r{Y_Sleh}GY|NV8#x=$~ z{=M%p-c$zJ{xa}X?Y+u4#ebD8C$fY{{w&!_+=^M! zx+cbA(V>2K^g;LFQmF3tRXk<-z!WU$_o;0I5w*A_9Byc+2?y3%-B8pUb>n-t8?tu0+`)bBmNM4lY;OlEr;q z*^*6|llz=f>*SUZaS(CKmw|nyzgxQ|c8PL1UQ=FYvL6q2^0e#{L8)_z%kere6EmEw z`Pd~gjq(4hrQB97uJG?&1WT&ciEJCk7DHhGFQe9%4EVbMuVmKaMZr!eTNStqvTK&W z>Z(-%oJ+M8xg=8k9pL2STrS>49_JQoG%v%Cdzl}y5Who1N>gZS-QaNBagJF`O4V`B zBWkAnWUA&k=U2Lqb4!Ti&yuagt`?Umz0*rva@M^E~oiv zm6vjxEMEgThjUrJ0rFx@qGs8}De@9dwCn~f|Dvbk?xkF|Mwe;ya~dtuXt6jcp=_Oa z7syWT4k?4N99sZ?$8a+HG?OdAW*0cQ8gdk0H7t`^MvWW=lib8Dm)C(=@u31TGG1!m zXE`FsPi0oi-^^TVsbI^xJdE5{31&)WEu7>U?Ra~Al3Qv=z1ieX!?m!AlScv5U$O1hFI%n2;}QNNj;|Yj8Cmem zn>*mNpQKVxIBd&j(2it-w00!V)WN3Dq#dhg{ej2$TuYJz^#9*4^xV+&8H`7uQ>*wh zo0s%c+h@~nt>VwD-}s#Onaq29h(v#D--8_2KHF)2G?ScQe6aKA?Uef^k(-2+pSXSg zj<(^>&hX%HPxyh(a!4(zP8Ria_4Y(Y7b+c$4&jf%>7{;OQWt+VAL<)u>+0?6M-kr* z_UD)Jm;ZbBlk(ZkGQCeX`AvGkoi{vZZXQ4U$5>4D-5+LgeqWNm|6kapF{1EdV;jC7 zeKOqI*#WgJGTbW>yF8nx`Mw&C8PVb$k%SGm2F zzS?0e?^~%OqyT{Tc;LWj!nkScU1TCb7>|*{qGST>3p;>~I@ofXWFEshNGDe>BV7td zFzT<#fUg_)L7m-knA>J|?XasG?XG&eqrvWKu&eDi@(l*wL8ZD3@(anQ;;Hh;*W1Yt z)V5bJQm`cx4SP-JjrNkQP6KNRD_e(lf zk`%uwB{vB$H8Y0II0bv&Mi|7RU{>2uD|ct*!hq~RxZ++l3cy2799_*feL7dmb> z9Iwm2gt###MxB)YJ>uCoV#$8Y?Vta=0Df7=jeZ&PGZWiQcGB?@9Y2M76pVehVF5g> z;|X2fX!AiGKbM9-spAtmuI1& zb^mMo8`jgT#xwX9T0(`xcOY3h@H`jl&4~Fu>nxxN#cOo{$XhX}I>Ft^o#a!iiKc zJ3Hp&S78`95h|_nX=S9|)H~EQbZ|#> zU)SJJbf6I*``aSz4{X=uX1|$j@zu<}nj7D4roL6Q_~yGpNO|pO@y-8kpDLoiZqI$G z+oHb+-yoEHH@!jV6Oi^IcY~JphDB}(gz-bxF7gd^-spCHcemyjUv)U#)7LTF9SwJN zpx^+wTN=GYd_??2@RQ<}#sHCOA~i&6iPYh9d8*O7BLk^EGQKx%#4mE^#qsxn`EdGz zcV2t}d9?ShTbI%3n-u9c*Wr2fdb+f0?F9;xci!_CYTshx zchjj2?IDHKPq+2v8m-mJyKWRo)z@s^L_hic_rm(0_;ftgxZ$e)mboq)4KvDm^pC;k z)9WEU&HL$E9d{>7HnnVPsBZ~x+qiLO)2{HY`i7P!bWb?B@DJdZ+qd|NcH=Z{>FRwT z+R@O}8yPrwiw5@h4G!&yb~frC+1C#YL=HBM;`Y%S=_YfLeNmWCW^?bnG8>{JUG34E zSxJ#tz{VD=JlcIT8z~a$Hng4CI3MVET#S)E+oWKeV}#RFmT?Z&B2NL20m#Q_$ZsGJ z1IaJJXyEq|h{3R9*bOJ)PuS_!agxu$XxMou1*i5@==r01n%3wo^MIRRZ_$-nufRP zc=Ek$uD@{|H|jC!ozU^?X?9fId1ZRJA>XRwyVCHuj>ppQ2^}9#!;SZbzfQxA_l9$6 zxbfcbvozdzZ@89%8)K+B$US(lNBhSz;~&iC{?%+C+8xpUkqq1&^nmSd8yrmD!eUAC z`~~+wACO15M6snuu)xbZU|t=!$)HxzZ_6=@|lr2PS`h4TA)hu!#J zUj5O5A#8|dIJ!UFi65s#HE;aEL`4fURnl)Df{0OggcfQ_xM3_1r#+oSf*4anoHh>( z#{8l)gz7QwVT^OURLAvcg}G!nn^v!Z>(koEYUe81{P*h9dpga?jKN03#y(P$y3P8*MHC_&g?yf zuAmx8FqXI4w*J#yBN&bK_^^pGQW(7&HJq9h+zT>6Ryb{QnZM`P!e*>tB&( zS53>GT|ob$u76vayivaa)A>*JdksKOi}i}r>nB+XM*3z7uNCNIOS+-wr8sHf@eXK0XYrtA0X`sTEOP1nDUy7P_SxITWy?=-Z2X)YN0#=MxXZWM!2 z{ddY&@ZO99ofGu_m99^FLRu3M!ST=s)vXqi0HH3mBoGg^x?4}cMlFOu1Y$5?8(T1Hfh3SE2}KeD z#+kSiNPs2DXk%whW}Uc=*@2v7C*XM3%uF`ZiJ8Fc@i{|gJVDO9+IC{dOtJwR1LFYN z?^aidrY$^V9>@2by7#}|{r~^|_rK42bU)R)dbz+DS%J|Kx}CF7qEpLn(=mgL(JU&a z_q37^UhU}#X@PwP%Jb{nzBG+zS|h2pgTHy4+7AA>`Hv+4ZC|S6_~Yy%Pc}aGI7pMX z?chHe?00e+e2@Olz43K4Z=9M@MP+8yhQj7YoA0se=JtD3`R{E@u8C8%#q{C9ALq8E z{nb1D-*!gf{ynpTYo%Amvbtx?=&!Lm`47nll|#b871AMgaE1C|{zCbhZdp$uvdvGaj2d|r>MLG1a${tZ)&FH_Q z+!zhKeC}lVp+fCD1rR=1c<2%Ib-}?$q(hIW265+T*8-XGtbFa-5=@B|=-4HETC7kF z0;XMml}{;zwZD4kPZ)Wrtvr(#h_M*Rt?E8HLcA+M(oqGG9c4yGx;M=WftNZ zE-=98`?MQGX?&SN?|Z1$R!cEz1Z#|fLZyz2R8tL3ZA(iIr=~wA8bqsGG=-N%8 zfzeR!K+g_4y+3Q~j)CEk&Vi8{3XOE{9tm|k_CP~t|M1qGU7Z8lTGy?5V0}l!)}5o> zTWfk-7m+rbT&|V0qiYmJy{pOVnM>LllGQR=0$>ekb4hlsq3Zs=Z&1Dw``B#KzCpfB zMk=9D*TXBxQ$|_~IjZ|tw~)t(T~Doq&h_DrR&v@ZNLxX+>cQ13$h&~Y#jDC7INHW7 zU>fA1Zl$4gC24`)m6XRLPhjvXpb%+Qy}Pe6t(;`f38pP08MV#QTcSEWwm#aiKKeSh z)X7Z`MApzW9{JWRi?)ugAcs&&S`#m5A~yqF%+$bWK63-KYo8&vt(5-j*3Q1ZZ9|=d z+dcIB;Lwhd9V2@NyBEu$_uL&F1EU>q=}7O^j+(}fj-I;Oy87DAo{mVkV;@cAb!^?) z5v{3NSl`%P*AZD*TU*lspVu~YtbAZY>$(S;SJ#nW`dccE)Z6JVJG#1i2DbF=*!sk9 z>ldiApy3O&V`TeKcW2jr+B$9Fernn}xbq-wUsE&CIk1~Lt807qP}iCo^zLQas)25b z4)*S%&T!3O@2^Q4bo-A4((mL$1|HeeZ+fm2ieqFD>9{mS$GM#cEt-v`A9lO=&*v%wqVW%gsdNX-Mbab0DCYNgniTmov{3 zByB5sc`t8vumxSCT?DICrJMIIX+5N(cYBOFd&pXBPl`y2Q&2^5@Bsm+ zPGRz&0pSG^Y!X3La_CeM%~XL@OF)W%=%=Miv|ttl^(tEtEhquydk#c^Lrxi#&Y|B0 zQOb$8K?HeOOaWEKq31!Ab4XA@RdA>W1ZEcbUIQ_QmqqQMDmio*#9U6?05Ol3wL3t~ z=g{{-EZ|U;6I2z4o&gc!#6N?m=4DG=pu!yb8i)ud-UBg%moYb}nV=pyhVx%ur4USU z*p0R+m4Z)}sgyolL~^B2ma?J9AWLKLe3%|XYiE@fk?$`I)SqZ0eKuMNK`~6ey(le5 z={ra(ib!9H*3O!pBHwqA&KVHMOo;Lig&znt>tH!^P7ygn+#$Sj=KMul^cnistz_}$1>W-bbfejms={2HUe z+2goC&xTY5QuCsph<-1!<>pnUv%)gaWe^Fz&I*I&WOGqzgY?|Y9V`qFVbG`YY9A|> zj!BQ7z}(r6U31<69vw#%kxHa|yS9lQ)4X&a{27}-P|1G_*G-jTH<0xTk#QH@C#?T_ zxV#(%6}V>Obo)zT{o8;lk)4Zc9L)9upX6y7hPX;5Xc+S`d>q; z7NvE#>Qy+NPF{a%deyQGzbUoHU`6Mpr~X!38G95hNIga3?-3=QBJA<>uJfuAyhSGh zq)OK`N}O-EBTWh8?UL-JOX4L(FQL;S*<_&N8|@S7|1EiNEIFa#Fn;nk6uU}e0ck?r z@5RWMbV}|M32G>B(`B7n&IKf~kry;lR< zM*C$(w_>Lins7j+-*ugW6OJ)D@*AdKrI!PM6R6%(Nu~STUiO3W_U{VuTeR%!*g1;m zk}E+N;Y61%x#r=PgZ;^-%S!SljeMQ7MLMP4+7v+J5@L6|%N3_{iM8*}NnW7S9Yt=} zh0h78Vz=WwC446~IkfXHv3FhVSCt>$LIbdKpQqzbQM*HXg^Fz9cBeyUC*M)@TdC85 zB8U5xFeN+SAE4F}Lu4$m=4+ESa|FHGH#AS#M5 zQd+Ue>DWzPmr9?4g3HQ}ZpB!siMQ-5od%Xy3v zGPS#%$4=1NHCdBQ~8J{3n<4&^hPuo4J5Ko-1cRK%N zi@bKzn@ZwVf`uK9>E98;#hMsnKkjc=oYx@O2f?Jyq7H{wCpOM5%j4}Lp;F~A{F=By z=q6G8moVEaPKZRxicKfj)5lLLsarAh)Ho|kI7rBMrkPhsvro(-*^y?`KrYvBP+OY$ z=cHLsLH^`Tt@)O8B}TqQ+jIhSkwrwuu8Zu*WeqJphtVY`Sj43@Aa-jFF&24CNu?>M zDHt9ut&IgeE)pVFmDwajUE+mC8Wb-OR&jKPNJ9KAx)LDKOD#-q_Z(*l(S3%JEVU`g zgc#|CEJhD1Vj(+#1-+29k@qUb*y$#n=sLT^ME5Sz<wODrPGAB?k7QQYU!i?763GiIp!6kGb!_JBOjg0iFv z44yfiBA*|-2R4!NzimIqeAklFx)Ks6q;ckxkF(&nDb3a%cuCRUA<^$hF~52WEAI(* z^0ELozV~lj5ZqhM|i2&nnwd?7f;6 z|LPgIRNY2NF||n+lVVU2$HdrWX*kFZ=Z+Ej^)Q3BocHBUjp%Oe=JQ5&@j;u-TtjtCEzGrAs!~MH@@WwW)w= z(6P%D55$B_!M$&HH|u z6}+BQBQ957Bl+!N8J&^7l~jW+*P!5sU-Ddo>e1u>2MZ!wRxDxR$&)uqemCaqO(pp@ zJJ3YpDHe3Q(fuAoCudZDKC$7br7e}lX3FBm~-G3>mKJD>b!(m%**EX{EL>FeBuB0}-q|hxLL(8u2R`cvB z7Eu>HdxE@gCTZR4(2t1{^|`plF~x+Fvj)rPaINKM)9!#p^}YDGGe zR9K3+lgczTsnp+~ms7T&OWrLN0rH-Xi;lx$EzLWYRJv}_*-hku4;{xTNOTcvme71d zI9-hCcLr;gjHMzXJ23C{n{(3YgH5W-b!<0_D4uO`RqL%j)ld|X-3H2hKj|uCrj>(vEtcccO zDiV!}{;NqLLgKa<=9Z!y5F@IhV4h*+#xf}?`!Cs~Q??q-(~z*o%5TxRl+E`GEX_Ff z*z;U#1^=(oN<^O9*Iu7)y+J zUGR(BY39iodF?TYVl)G%oS+9$#BRS3qb9p!MT`WmJ0+m7&>bUglm{sWD2m;V{W11E z|B4t3C^1&`yI{X?|m4vh=g=wd&iuXrJ zxfq0vJbN)s7#2fk5$2xb8Z6YDB%uRlrzFhiYa3l~8Xt?f!`5s$Vj zO-J0?1v-qwg4>Oa<8tBb8|+-5$faQ8P!+5wvWg?PPVyEU)a=UnI1|S)$nE=6PJew` zfoEC5i6h@q`3xP1f7ujnH3hq3#OG4IPsX&Kr~8*M^b z%s+lfS$n@y`i_wDi071Y5v>=_DZ#f2W4?KIsIGTu&nDSAw85<%cDrI4i?dqvqSwO` zZWu#_V$zLU0(-@xV@Y@O53&9r&tT=5q}!3iA|iqa;F3!y>BbqViRc(A@hjL4PPt(< zs>~x~$S$356G`RtcS$CnLOahYd$G+FP^t+Fep%9c$_moqdZ59a5GVO7^L%*;@@{%C_4)5vu`RtxD~@8>Utgyuj2K*h|_s&yC9|VAyj;{ zl4->90;lnG!GAI#OKGwl#|e8*LiYa*LljUyRbFO$Sr3M{xk|-_vQ*keiQH|4eE$;PSOUphz zuF$1?3g^iQJI;*E`%4Vq(vxVTLAT>9qonttA7EEe(q8{7 z3G%&i&MT+ALXqIZjmN_{7rQT{y?LioOi$QSSUQ-H#F-*F9Ydg0-G7O29+l29hGX<= zIEvwrHtmhHSJE}<607!jl6fLZZD*M??G?|l=st&h05^$DN_(pvj-)7~+c58xRKYw^ zgaxjRtXc+vMeAZl{{|7J;b%S46b$*X@40DoB-fXqCct zmqpB!EA8H^DPvZL?I@RZyDt$Jr6Rk0wm=Bm{l7$u!uD!jHQAGSj??fG7xO5eM+M$XX?MZPF~>flgqdg?5c9)W*S4!G0jb?aDrG8 zr`CXa&M9rQ`K5&5do!Wbijr-kl83gIyX1rOcNF@_-!Xh_AOVGq;_yfZQR5qKwKV7>XX2{*Y zScd&7qkt1zh5yM$YSP@;?H%-!QtjEODc7s;Y|0;eZ7zhQb7Aq+y!nG5 z{ijOFe+5r}kJ3B|`M_4F@qPQL0`dRNz|t0&6(b6K2Uy^HoGRhbEuMXiQ+QQK)t_WB z5(9t7^E!{7s&a*c-6+~2BDS$iVDvnX0jD!8W7&!`*^XO%sUBIad9zjkD z+{dZeh^kv*k{DRc^K*FgWnR3AQ2(hA78-Tz#SR*bL1B$tLT&=7{8ari*~D?DEZ z!tmr9$VnB38Ez7)gXkKm{DF}_ulY_Ws_ zFq#4{Bv3;OaF|SiG;;6+&8UyzIAsRRAhr>-SIq#}(x8tA^&{W!J=bV3;oOG%V_a2L zTX*yi_V#sG?HK6Wg9-0obg4Kr+C4PfyJJ9a422z`k$H@`xTXb zzaFitUwni{3mT44XXZ`Si_}yXiM~j?rY(Gt8t}gAe^7Dbi?pLr5B1d7zDV&He|hy4 z+7<36ZEbW}WZ@DvT7%bK_p`CuaP7kI@?}dG|CQ|u>mBah($|g3 zcMfdT@#?Lb{0$G%gwPUc9j$bP20Mp3`}NS4`q~Pz`8vwzirNu&%ym5J2n}^_LjgQl zF?3&a$zmUM*4KjW%A=(YM_<~Q>^3kScbi}`r8l1DoO%WxpTyD1lIGvZOr z$n$bgqkjpj@#58Zr`eIuJZcq5e})C9j`TpJ4hu&VQAy38vRSo%$^vzN%BE}G7-O$8 z*Ak5WB~%cIEWv=T;S{nIs0EdWtnJeRwF%YH&7T!S8y+V_E2+9>^H&9-FglO&YUWXO z?dER^LZIfG!puPQo5FM+p}OqoL7ts=XY}8SD*uU59zEZICXaUZh38&pOz++O25H|H zJi;9EoD#HZ@)btfgx!Nf-CcUc;7URJzHAFDd`u8Z8k5YEByEF0frdu%h9aB)ih1!E zGSNLU80!91Xm~o+4MrABQ+{KMRSZV-Ekm7p`Rb02rAzBO>K8jg=I^7MummQ@|?z_koG8J zpBd~kUV#6;qRqqG*SMfXDL)X2Qb}R-8A1GM%h zBho&w?b^}5voF-QTnfgIFw0FF9_}HLO?KFY2Omazxi>&)`S`Kupm^_EeHwfBQ&l4Ve}Y)!ul2rSYFb&ph{35 zrD+QvLjOO6F|=^;sK*=Ej!`&aG}@YMw3Vk7>f5my-TY&q!A_6Uo(@p}YrJd#fb zR*z51}!NxK*bO?%rlXn#-AexcR5V8;cm!37(Bu2tE;!DjTnz&wQ_ zbg4X<)Kw@h^YoBbC{Fi4!({Sp?N;6j?MB{u=^KoF(WixE-uWR}ha@WGRuqa&1~fM- z{+#?>6|(`ardQdu`SN_U3qwClY^7bRl=oU2tBN-MIJdc?AGfpwzJ&!!9$S1G%qK5*FM?7VoXGnWZDA2x+?@$g2 z8#WAdZwn27Q3zB#qYBzTdKl*6eVCblAZTAEDye%;5VqOmB~;SztRVewY)?yswT|C& zI5Ovi|A=k;j;~Vgy7%AS=#4LevUhrz@-f(vG03nZbG8_6{aEZ+IOz`@XKq{aarF;B zKec`g{ZZRvRGew{$5222%<1E4w^8xO)^414KbCQi!!-QOzc+0=RL5K2_9b^FnU+cJ z=gBX4^jn0iFAgKot}Q!zdOC)8_IEtdGmo;1>M}+By#xK7yYE&y+&zL1^Kwh0k&G<9 zVjS5qw54}o#~@fDcj%v*2WQ=XTa{cdxLR2&3DDbW|}GI!;>o`SXd=;#U^$i zFIn#Keo!T-p_Rm<@v|962~ePY$g(?18Ehn~r_gC;%o_jkeHH}Mll zd9cQ(8+LO&xp;+%r%ZmUCtB}OyH@>fheo`b(` z;z<*aab*aXO+1~0d*EN*F5UEV%;3fkTueM};>KEtvIk5&nS&3Qcsd9FvWe?ve_8YR zTP7YiaciwRW8#S%{B09Yns}*c=S>qomxC7~Kh=1s0`BAAXqds8-`tL=_+J2CN=}O% zp2~zjH*nX7@Lw3Xad0=vEPf{j5G4R_Ky=W;K=^xzPGtv=8 zsNK@N#UOa5svfTwj9(Ik_|J=ax+9^jgM-x5If5hkvd)oCocC8pqa6tN2~F3|{{B53 z#?yE_D(+c3%1dh^Mj01x9T@2y*|V;DTkr5l_fQLt{aZS>KC#x|oBRpW>PY?Mt-AF) zn(VuZ)sd+lDP-Mttd2~5ub<6g-hSWy;nZsL#`q4V%)ROzOdo=b8@oHCj5{vwmf8;d zsHm6whBw}7c5`2!;g?8VM@RpTuAP0|9lc#(9HP+bmN1VZJc{xNzwlb!Qp=+{9@X=x zfk%xvmuD+|xN|7mN7m23TJYrkWDJ3qSNvzR8$MF8fW#ROP+9xjZjf zW!-yOrhi0{`Ov*%O1b`CaDH`?*qG4`)hNtX{LUd3DE{ z<;&N%Zs^$1ymWOd+|!Yn{=4v7x_i7)yK}Iv?j3ldyK8ChK#?;@@a>a7kKsRy;UR36bDW>Htn+AtiDNq;zPDNY%S@c}7i8g&W#L?prMJVxbM^FP z@fU)&^nWf3=Xw^H`Tb^^tLLj({Jf2po^NI0Tu-@~KVhc1dj2ene>!MO&);O>T#t2+ z{ZkXq)pIe6e+Fnv&#f$+>){saotN+|T?hW$Br1@>~67?cnq8qNg#7pWANf zc_0htw($CSSbrBiomu>|z-Q^%m4$OX7XInG=!tT}5qNz%E{lKYF8E7#!QZ_L9>sZz zhY#t$3;xnw@OSTmM@&cM>bLHM_~WKJTvofR6B9R=Ulz{iAP*elXKo9(#+Su!^=~O> zM&R|<1zR#m{o^ z6gL3jdM26FpEhxqnNHfn2?RX{*G+tN4jwo0-8p#D#54D2x7(XG@pC!+y6L1FIe6T} zD@{3`@GEKJ_|(7%7M?coL=LW-vhX%!6F`UGm>`E3I`L-^(WVE51n3^5dUZ|~^A$2;Ab z{`&XEIu!7slKUhg28@TxWZ#BgPY<4u{IHSYKx!t1e~$5{LLY!f$! zTSf?R1ovH21#?(W7V&+R>-@LM<6}9)E_1j|>TNXZ-&j5g8H8oOg~mW|o1zwArepX5 zkOyo0XL7gOZ?$6s3OJqXe{U!wQDF(1=^b+c*cx0@`Co0x2<$fHxc~XuZrN}7e;Y8q zw&(h<(X8L%2a6G`yp?VNI#vA%vwo`{UcVieC1>Hoz_g}g4d`-wP5eV+P78wbCo1RInIM;tZ6ig+5yfq`?%n``R(*Fk=O!8;| literal 0 HcmV?d00001 diff --git a/build_linux64/objstore/cuvec4i.o b/build_linux64/objstore/cuvec4i.o new file mode 100644 index 0000000000000000000000000000000000000000..56697de147e6b64932d4cd94b7419b2723d79d96 GIT binary patch literal 21720 zcmc(H3w#vEndevCJ<~lsJ=61$kN}}+JpdtKG&6bvw$%az!X`!r9Fqtl2_%7RA!H)j<;C&+OwH+CN2UHdHA7)XNA z{i~-Ngd%+2$6X4l>Z^Z!^;OkZudlj$^kC!qH3DPgI7YY8^^}7W9r?(0+O$ZM=1~E? zs~7(2#jdWX9yws4JiV#uv0R>LjHjCp{rX;NI`qAUZ%F`}9;@c~z3e=XwtW0vkT!4A zq5oyE-^gk4ee4q-j<2<0%gl@lDzmG$1U7uE;X_v4(DETg{zv;Vy@@mR#n$1W?@jMZ z%ZoSaze|5Rtq0{7^@y1rpOG8$54#V^hlN9HrNivdTBYGI9a^isIMSs@R+Ghlc-F(e z5(aXc9$TwD{3~%F0782BSFA~IJoO@I1-v(m_rI0DVe`5}7w&F4^n*>99@dDw+o11} zO5fG@d<2tatV_T8ydD%=wl!>hkHTMnF8#vaf%rZrUg*;M zzN;WeBiMdg4|;E?{|egs9@_iQR(lUO(pN66IQ0DCMyX5hef}$#J|!Lc&f!Ls?ET4C zQ2fP9pHdE;I^3uM_$TD@_{{o48Dyctp646%pj3K;{2TSra9AqUztKpC#D>GGOC^h3 zC2?BG(ipYXrCOHOpfXFAGv2?cuc`7Z$2ut!sI#vJgU+X{T z=L%>ft9=%kZJtGDTW67wciwK?hm^-9L!FIQ`#>s(8m;zmH6Ci5NuE74zixW?O(i>> z_(bBXz?w|fMq|&nO#Sz{g%riH<-gM}4cB;uQj!fOZkmrY&t$bd3TK?fY{Ts88 zc_{DjJs6u=hwhOM--Cexp&Y&kqXI&EamspvXB)9jVTF`sZV6nPd%93;4687UPgaS zLm-NbyJULLPgSleGN}%%Ce`qLKia>me_-!G`*w;dwD5u53QbUJ|E}JSA&T}6M!S2v z_IT*Mc{}#>_7Al84pdTfpmS&-+WPT3Yuk7C@7UMT-n+AL;~jTyYOUR|Z?JPmRd?fZ z(if1Xt)o31gUIS$PXYfT(l?M>DWTf{Y#@CRxwQ>cvAgGMlw*ZKwt)1nQ81g33Mtz0 z(RJi6A$=986}#83BEJ=C53YmEO|jNS^19}ezLs1See2g!U zt+Rx#BR$f+jzT=}NBX`93W2_~d*~w5OUWYzq)1{BsT=lo_O$ntu1_+T{y35TMatE` zKp}mE<{7IJ1|>EnT1E1Ih4dxl(eEaABz`yL^B`x#>O|wuxH?bzRsVzqwo4!TYCpvq3nU~9j%ph zt*uqHHPzeews*D0W32~hGSs?bU+c2=jzmX$BGwvTR#g{o#b{L3wywK#bK}N48`f7- zSbCL;;x!)nr`C?nuHNlEdv@H{-}nf%&#HZd_6+RW+u7c6kapxQJ4p3A`t}{7T^lNU z+k1zoy`rk?0qWRLiE*Ex9XjaxL|^xQYL8X+b-zt|p?8j6LgX)XLBpiIq3XY{`uz*Y z-$Q(i_mE4ONB&R9$aVYo@NpKkMPB_U)=2!s8tb2sml-r2(*N38|qE|nQN`8oW@jE2zFH@0zC*|q$e0ka;^8cLA3%9;Z6ZIj=iNuE} z-@AnLUSEV2lD?f3J|Q0Q6;`Z{-$T9#0#W~(Pt*UDd=;w__mIRPl?Og5qK1;eo*oyi z?xB(%F9{0iLBCgLq&%9y5*DctOZX`hq%T1ktD?7`L*F#c@N$fnRB(noFBl9?)q9(;RGTPx40$n*lxfI|bM`NT zNO0m!5S2V{l@C-EhkhSKH7DK!F^A`geo%8kZ8-{KHZ&{*nHvi6K@09-A)J0AKlte^ znaro)zuIII86SsI=FJE8-F%XV1rg8!=zs|V;7lrbma6bfI3e_tW<3X-M}f5k^V!ls}2h{}+$2)-Qm zfdX;U&j2k!aw$B{<8Vb8g=2{0@Co=zq^rRCvb5<9u-1TF;@UI;vJSb+;Fkl!65oXY z<^ppoMafC*7x|Po53>O1@yid+Ub=eouZuh;QrpKLd{tj(9%4rN!8Y+;qU3`lWSF9n z#^@dIMY;GY9gC16UD7Few#9=uEsV8D?f{L8VWfnU_Qu zQwO1Qy&%Yvbl2v$158#(GB8o$y#|}_R}416ZliQW(G_$GYEnLFbcVXN#uoqUL0?Pdis5jJpbMm5r+}j71cmdX|#G zV_Q}IY_agJ*1`{*(FdvJqMW`;My zilKafg2PF=pRUq{0#XBX%xF=y=djtNw|X-)Aser6oyhynRr5fEH1APLx~awIJ$iyh zyxy~v?~b*+mtw0Rm>g-9u zD7cVh%_yadCi~v*7TJ3V1$$alHDFNuJvrUR=jQjIZ9hCQ9UC_jZ^Dn#prus6dNxYZ2HqYRrXXV~Mf$>zUE7wV;j8CEOv*grd8xxPQwCFoUDVE-VssYrAmVsMX?FPH%Pl^My)H+@E0&C9e?XuI>&9 z#cpv6kwG-dLKjIa%7;yJmw}}T&9u3F(43g2zzZfNNcJ6L&2A_)AH0E4*8C=xRx%|P zDW`^*>2|B^1T!OUnT@jWL`sy>Oc-V>f7IfZMp@vO235TrK^vtLOgO{ZUTv9>Ggnh? z?F1j)$N%KG+;)*(iqJ7uI;^Id*eAxZvePWAU?da-VGU{<&oi$YSg{@eao_Tzu+j2H z=y*alt_T=BkNO10!mmBC1xnSP2r$89FVyGzwM&#Q$6EaU)2vkVw22Qras_)AxjkFN z0Jb($GFlXaidk8ki1nB62?))ivXfG$*>dqLwTVKrE1`&I6@kI-Gg^MpCPJab%6<}x zJhe@(QW-LXM#f$2_ME^RN1uvSwM>X#luNDo^JmL$PnvQ> z7Dz}W6}RDjJScyCk_K)eA2uQp>hodWYj|Bx{WL9xFzw=g^&DoEZpev?@(ZX_JU32H z=@^_N7*|~f^Mw3|SHSEacf~b7R#LzhpmwPY(~QzW*wf~H-0#1H>yFWaedk>q!{$ZI z0ZGHQV#w*O9%WbvdD1K}t4(ajqEVJVae|6=xKIb(Vw7vDb^AZ4`Pxk$_j)Nax(2oXxiMQES?wQc}neR z7Jd_&QbuGf2S72NB#;YC0kv<*Gaw`z@zWWsh_gpdhCkjeqWoIWfM}a(nvnE zg%%cj`~k7(=W<%DB+Z{MqD_{l*qgGEV%SWqbc!W49eXHt#0*<)VAD9mvULYuOVKSa zb7`oKCB<8?v0C(_uVICoPAdMEY$WAu^??f{+yWa14N#IY`j<))a^fPLQoPvUJxT0i zOg!)V(&%wHeTANl5GJzbE^|vwuK)f`x$P>Q#IB_sk>gZ;G{cyotUkS!{6lg`9i?Kz z#(DgGq2Y=EF5mJ~Se$OUjJ;3YInIpt7n1L;dPUQY4zalG-`U1$6#q^Vu*(M`Y{;%|iICB$mH;Zj1>pgeB7x8*6zjDF!ZtLQqkhCZEFQMkDg7Nvv;E_YBFB zsq0DERW7*^ngKB>hus-h{34z|2@(=AsnW0=Z9aDpTBBiMJuA* zhdCkEF>E&^CSZ`Iy^B=Yq<{yHlS~T1iY@jdNO(NwkR+3qdDOMY_YF(P5BmVZm-!I1 ze1d%H!8AoqrAf&3BQWXl_bD^UuHc9z=EviiB>QjSwMiC{ldLp+&?X8|06AsmJ94tn zM5jVhT}T?&18*7bDs0x5b=PTkj%1LK(p+0yLY*sTSaBVQ=~+OE)BU*DghW}kkJj!+=v8-XpIx8Mpp z16xTxgHup$atGWBu5CurO(}n|d={qB7*>Vn;J7bb)22wKpQU}&L`CjXO%!;o&5st- zm@i(|g3C2Js$deZWH>Np_~Qq0|M3S8j*=Afo7kYc@H zAdcUZTS^5nyA2FtGK_8l>r(=6rCG^I!SrwXQH!GdGR2pC{DN#4v%*S>l`H5SiiJI( zJqReHlgWxTkWv|&8_DREntBS`oXKkZ`YBVX!r6-r1%^|xPur!IOz0UstY!53wR411 z0>(`P?8A8`RpCsBG3CqX%VhVtw6{cYn^fY)5s`Z?qpwb5c3~qMWwYc_1}5()i@%S1 zkq6tGHp<)?T?nXrtWwN;4;PgTx45Md7Je4P0Sg59Tot4^?!nodVzZB$;SvuBaaXWd z(^DR1ij?$*^uC~WI)`Buq%?6HI`M2_QeQB9!3^h#L%{%Lyy9sO`n=->?5?76ITXUi zR4$9ASgDAohJu1Ag1>Y^H^p$RDMsMcdZMGID4=7PhJuo!x=b-dN0?-IJkvyN$i8Z(+uS17zG@O~L7^jA+t zq&>RH0+Ev@lLN=wlvytsOiWSmz%E|{jU0G4lMB6tYPa~<*9vk%8E@M&e+M8-rPz zspbStRm64IkSHT~$8w}IIgd^-g2gWrPkGG=X4acbv$CY5DvjuKeZzYAF2kKKj^)UP zd|>0k)MQcmNjl-eedz2iS6DmdDt<>!yD;~qCb>oxVKoJgVCQL;%T@7w4$Q&@W$RM+CNs2ov?y*?`m(+qkXnG=?Q`9UP!mA{iD*q~U+nwpveOZQw8 zKZyP?<4UOdxg&EzdR-{4>ZeADwFLt&rxj@|C%Ei{tJLo||{De!0J(lVUvqT&nP0zOXREW6t|y_uqv;1DW$LsRr<6`Z8@P! znR(**g7A167b@g1G*L>2SX*8djGllZZNbeGv+QCfr`V^SVYzZyc=mIKC+rId%n%q) zWpcv4J|aUglCCtHb2c?Aq>sC_)0hJCM%?^-zB9~yn&qxF!lmX|N-TY)rLCaS=L?Y# zOqn7q=d-b}G{&lQ{c@YOqmVA(G4mKJ(bbGw#69jU-FQ6%8y-(p=K3@2sd>72l?@a? zIT_es{?SF+cuw$m$}4BW z;^iFaJbM&&)awSUDrB8uWQgk-Hoh^Ka;NdAG#9sZG%fD;owE#z?<57>RR!>-#PCtW zU>W907HBY;+B{W{;=&|`eOp`-sygJ7 z`ZzCZX7W%$In0CHu(pL&xIu_Hswo!fJKL~x>y%p}BM8ZgFjoeXLb>_inFSvgJqHc1 zWcs8sFRwP{)ylk)i>U%q%DgNk79L1zZo^B-S)~y?$Cr`?V3CqAiM|yNJX-Wg##F&lRzQHr18+0?NodGG~>HMZOh3{-w8ZWxG77 zH@_|!?5;oo%#I{I_K6u<+($U%dIib>hVRN0ljit&~L0^6o21gys)zf8x;uPg5 zBf3;4&yrF_9q;2Iq~xme7$i~D=jj{>h z09LY|cE8dxK@At_f+$^%Y@(5rS2ZZ=K9wrKMmxE1iAWg=Po{|(^H*Uzo=hP%Op(bX zP5wE58L@;u6{AnUa=V2|r7)PJv`Jd)goCk04_*XXSxmc@dkv%_hfH$KDL!KYDF~Sq8AIF3Gm9=;KmaXU z$Om>=Tbzb0NC(d#cy7%kxHkECME*8OF4`@IX}{kMF+r5#W{R|E1&vvS(96gj01_$w za$+_tB943!RWH4jOR>h}6(kqkMiT;TJ0DhKF-BP`T!X7nEQyBRhg8&P!afsEqt0h} zvTYp-^C5}6&y%U!O9x^%~$-F@9XolE!h_B>$V?E|Cng6Lr9-u~`Ay+&O$rbY+a_wMW* zF#315RwvYGdq>CK&i?+^{_an98Z}i>np|gnK4C=Ld%Ab_8dXMgUuDAhp!fXtrxYLN zRsYwW?=blL{z+--3xDHQBT-$m;wc)ORr?gRXW#ojP4(6B#M87tciGcai?{!OKm~PA z)1Epb+Er8aG_{%hQ-Ckh{#Zxtnl*JXM*12Id)FGay|k+98M?EpcFn4<;sXbEZ#BLW zsK>OZX9uujZ9qm{J^NGw9}L{aCRtT1wk#1}v$DFXsw%NMzU($OSc#7aZe=4?v8rXU zHLF*y_`d4`>+bK~-qVT7_x0{D@F7D5g=_DkNnusIaj;B{_OihHM+NRCo-VP`Fn3k+_oY}?KM@PJ3_QlReL(_q2lPiJlYUnVzg`{t$o`H zb@!e@m+R~E_X3KoQ2PgT3bZtkv295nsX0`!xM~&&kvasW@kAayxVLj)g{pq0goPrB z5*FmaIvzY4S-nOG;thTYYqG+}bSusCK|S_Gv>l(hWIq&%GCf9K>Bra<>!>6WucOjP zf(Mld3ae%d^QvbHk($}UY<(B5U!{W9fVI>>vmir-l3b_I7p{ z^ZULb=$jO+A|Cr^K`31IPt5;y(*I7tO0J_oG`{UB3)DTtlAQy6(VYX){@Il1i!Ujp zca-=F)tuiKH@5F>H%iyHwys=R(^|PgjZ#VH{ZX12><}HiT?CP!XQTn6W z6)H+Z_w{x|JrUZsT*YdP_Ag|D8yc)3@8Zhk58++vLllIc4WCnUj7oYtqj}8>2s*E}P2kupwFckDi1}7N5AxN=IpF7XR96M| z*}fj@JEw1W^m2Sk6B6|V|C}U2g?WiJqF5y^jaOA;Frywfo`eSVG7r;NP&g7V&Jy&9 zqL=HNwO5pAV%v)fy``|v2-2S_Y!21W-<$X^g!lCI6soSLMTy1tdcrk#)1t*ySTnUt z3OGmgDd9g++jmM$;xR+tge1gj3iB>A9 zToNI=m2}v(k;*}FEyw&V`+SP^^)ynreyg8PK3u73moiXYru$TCjmr3#~wZ)BSsUcJJ$n z_Uw7v-Qf8+>F>C2@oXjiJ@-l#-~L>7n-v{{522Mjyxz)(&`ORhm=Zo#-_bpI+1-HV zq1-LJSr3SuLW>ZE^mtA+K1v%A`S(vA|K~+L&PQe|E|Zm8#R$9`pOy)FiuKSw)dr5cBJnd|psK zM!CzL!>oD^Im=cNmfBwk(t`Ms-xZW?v@o9d3n9Xhio}Y)2?QeXSB2_|-RoWz@Yx;0 z0xSH870$Q91@2FK^;d=EymPM#ymN?5_ck7VF0%gI>=0k*X9X;DRsVs|E`PGp->D`@j%cIg=9s?M}H`UtZ>*_Ql+aH|_a{!Z8^KhM}t0#*3n@Vm_a=YqZz?+N~gpjT5~BwkH(^f7-e zJ^*v+SA@CxZKOovKk+jVy8fE~pr_ZR?+8_EXjyxxR{J>70a0%a)#+P9rYBFB(>+J< zkBAVd@aUSz5plJ@i}Vq3wjYA7MK9_Txoh>eb8pw*%-tg83GCB3c&o!FkRcl=jk6@xPPhO0kV&TDxJg(|}v5imD

f$>HC;n}P8&&^(|e|Gjp*(=oJL(R{LIlgl8 zf7aSf^h^MUcP3Ec|1!R^6ldY{OSdkF_igU&+!^i1NiqLd9zm}PGHeXDP`ElM>VH5~ zSREFH&2IN?n1g5E=(+)9eUUoV z0)qYaTWCkNK~_7mHdlY+Z>1f}uIU42-u0Gb?c4wVdf|3u<8NU+s=BO-v)%qJ)XzVr z{q6MIs`$6oZ`=cZEAt+g4Cvc&{QLi=!__?5^q4OjdGrPk{ta%*m+$dJ$M$_) zU9J85cDLTwRZdg0sJR<(tCNF*MV^03*<~27CIz1(Q`9X!}3GlqqwX0SwH_GnbzOQ#+pAk<)E2FWc z@!EY>I=;w8cmX0j&!Pal3x4gLcXJVGj*0Nx_2yB(SxbL=IXcG}CajE&8G@MIJo>eqgOwpwU}G_!v&J8IR9NlLe+1T~ z@v!s>yZyX9oVV%Uz&l{`HqRjUv(x0JY~HpRc!fK!GxIjQWb>|{f%lWwS*`Y7wt3gh zz`K7MZz-R5a36-Z`9t!SLf)rs-U==pbLkPT--pGueq*cE0zTS&yq>Z7@;DzH$8f{4 z+*Ir(8?%?F#h00e@%eti#+%Wfw+{XF5;f~K^?k>S| zIR;O_Ir8z?)WY+#@HKwEMmWFIPKUq3#vQp%Y}&Xj=9+w-7D;yCYwjy?i8_}KS4U{i-}e0W+8-y1o9z>2SDe%Z!T)9~N7@qByE zIVGO7@gvjlXKXy+Y|wUrKeBN=mAJwEsf{a+;_cjT0(Zt_YJtB69>R!a`$2fenRWm3 zCir_c?)1xUjE%S1cn#4S&PN$8Fr%^F*Y1`we@b9DE+|ndU)y6-%W$^<2M?lsMR@4oI0vt^aQ@~d>z=gZ zyDj{B51w`k4D9Y3>KbT`!By>u$1Q>mjY?cV)BtM#@5Psh8O2YGHKewv9qUc9k)pnKqfjh#EY`v*Gru7cUW zy?w`h%@*IapVF?6*Ic`@JHOhTdaGC;pXt%UlxxTO_{?|ysVw$2{Q3`g*V`B58ns&cYtv(Mc1#2 z@gUBF1P|~_uQ8Xna0pbmEVRG}Yj-#b-D&d-on;kURm;&`t#6V89TJ0pG* ze#po=@ySl_x~&6-V(n_D+GtX$s+^|WR;f&KXX_J_Q# z-8f~}clX}c*|D;_w|(yeA5y@sJ^ce4JG)ldzr%0XySM#;#vyDaz3n|*T>H*Wnu%v! z?~FXFI|sXWbbgSODI7PEaR(~z?D-%WQ#hu}z|Rk|jvt@=c?|z)BzMkv%yDjQIreBR zXzp+vP$9g-znX(^oPQR)gX7`84SDWUCx)M z@bf-8a{hP<&gGQa>0@?0UC!T3;paAxBj*QGa4yHWWBL&O*H^Dc;&f;#m{HJb$zjG6ue^@k=oG;u2 zcWy%5eyfIe`sG+myt&*vIA4R@ag5)79Nd{-4!<+LMVuLqS6>6~@NVyUHi9SEkFcl~-n$bNHS6!{%wYbAL!q z!;|*B8J&hZ_lLinhCBC%Gt+SA{_wMDxO0EFJOy{g&{Al$|AF1s-<@&({>=5iM8p3E zi}iPA@W))Zp&i}c-=BRNi=6D)OVPc1cwQ8xCyX-j7B;F|UUoYk8;ckb@&9muxhn)T|mgQECM%*v1U z?iuKe;!){RG-StpsD-C@_U?<~@4WPN?j1lwEWw?-TD$OZQ>Uek|8QVH0-q|=pG26z zxVs(qpQ-MKu|Njj)A^Jz;itOcHQ7{~J@5D#!s~IK$2jNs0vn&Ywt30lomMZ$xU;8q zHi3|zt6b*4w|pFBxlDVyfy%lOyZ$lfB5UJL`yD~a<^pusB<#lKte+qu4@p&7=XKkaeKJ957TY^L_7?D`iv71^;dEqxm4naZEG%il7M->Kh$ z&)DVdWEVj!}#W?MhbMQAn@O3$T|B2b<4TsT=r?;O$ zAeTA4e6w91KSjF1O)q~1xigL5h&_JJ`x?G}^WmNHoaSz(x;egO1rPDQ#NJ2wIl=E= o)64TcVOkS`zy@rvBHBOg9N_g$b$m|ERDO7UmdiV>z_cX&ACuH}BLDyZ literal 0 HcmV?d00001 diff --git a/include/amsculib3/math/cuvec2i.hpp b/include/amsculib3/math/cuvec2i.hpp index 172d7e5..a1a7d80 100644 --- a/include/amsculib3/math/cuvec2i.hpp +++ b/include/amsculib3/math/cuvec2i.hpp @@ -4,6 +4,32 @@ namespace amscuda { + class cuvec2i + { + public: + int x; + int y; + + __host__ __device__ cuvec2i(); + __host__ __device__ ~cuvec2i(); + __host__ __device__ cuvec2i(const int &_x, const int &_y); + __host__ __device__ int& operator[](const int &I); + __host__ __device__ const int& operator[](const int &I) const; + __host__ __device__ cuvec2i operator+(const cuvec2i& rhs) const; + __host__ __device__ cuvec2i operator-(const cuvec2i& rhs) const; + __host__ __device__ cuvec2i operator*(const cuvec2i& rhs) const; //elementwise product + __host__ __device__ cuvec2i operator/(const cuvec2i& rhs) const; //elementwise division + __host__ __device__ friend cuvec2i operator*(const cuvec2i& lhs, const int& rhs); + __host__ __device__ friend cuvec2i operator*(const int& lhs, const cuvec2i& rhs); + __host__ __device__ friend cuvec2i operator/(const cuvec2i& lhs, const int& rhs); + __host__ __device__ friend cuvec2i operator/(const int& lhs, const cuvec2i& rhs); + __host__ __device__ friend cuvec2i operator-(const cuvec2i& other); + __host__ __device__ cuvec2i& operator+=(const cuvec2i& rhs); + __host__ __device__ cuvec2i& operator-=(const cuvec2i& rhs); + __host__ __device__ cuvec2i& operator*=(const int& rhs); + __host__ __device__ cuvec2i& operator/=(const int& rhs); + + }; }; //end namespace amscuda diff --git a/include/amsculib3/math/cuvec3i.hpp b/include/amsculib3/math/cuvec3i.hpp index 90f6c1c..e98a341 100644 --- a/include/amsculib3/math/cuvec3i.hpp +++ b/include/amsculib3/math/cuvec3i.hpp @@ -4,6 +4,33 @@ namespace amscuda { + class cuvec3i + { + public: + int x; + int y; + int z; + + __host__ __device__ cuvec3i(); + __host__ __device__ ~cuvec3i(); + __host__ __device__ cuvec3i(const int &_x, const int &_y, const int &_z); + __host__ __device__ int& operator[](const int &I); + __host__ __device__ const int& operator[](const int &I) const; + __host__ __device__ cuvec3i operator+(const cuvec3i& rhs) const; + __host__ __device__ cuvec3i operator-(const cuvec3i& rhs) const; + __host__ __device__ cuvec3i operator*(const cuvec3i& rhs) const; //elementwise product + __host__ __device__ cuvec3i operator/(const cuvec3i& rhs) const; //elementwise division + __host__ __device__ friend cuvec3i operator*(const cuvec3i& lhs, const int& rhs); + __host__ __device__ friend cuvec3i operator*(const int& lhs, const cuvec3i& rhs); + __host__ __device__ friend cuvec3i operator/(const cuvec3i& lhs, const int& rhs); + __host__ __device__ friend cuvec3i operator/(const int& lhs, const cuvec3i& rhs); + __host__ __device__ friend cuvec3i operator-(const cuvec3i& other); + __host__ __device__ cuvec3i& operator+=(const cuvec3i& rhs); + __host__ __device__ cuvec3i& operator-=(const cuvec3i& rhs); + __host__ __device__ cuvec3i& operator*=(const int& rhs); + __host__ __device__ cuvec3i& operator/=(const int& rhs); + + }; }; //end namespace amscuda diff --git a/include/amsculib3/math/cuvec4i.hpp b/include/amsculib3/math/cuvec4i.hpp index ae44b21..658ebd8 100644 --- a/include/amsculib3/math/cuvec4i.hpp +++ b/include/amsculib3/math/cuvec4i.hpp @@ -4,6 +4,34 @@ namespace amscuda { + class cuvec4i + { + public: + int x; + int y; + int z; + int w; + + __host__ __device__ cuvec4i(); + __host__ __device__ ~cuvec4i(); + __host__ __device__ cuvec4i(const int &_x, const int &_y, const int &_z, const int &_w); + __host__ __device__ int& operator[](const int &I); + __host__ __device__ const int& operator[](const int &I) const; + __host__ __device__ cuvec4i operator+(const cuvec4i& rhs) const; + __host__ __device__ cuvec4i operator-(const cuvec4i& rhs) const; + __host__ __device__ cuvec4i operator*(const cuvec4i& rhs) const; //elementwise product + __host__ __device__ cuvec4i operator/(const cuvec4i& rhs) const; //elementwise division + __host__ __device__ friend cuvec4i operator*(const cuvec4i& lhs, const int& rhs); + __host__ __device__ friend cuvec4i operator*(const int& lhs, const cuvec4i& rhs); + __host__ __device__ friend cuvec4i operator/(const cuvec4i& lhs, const int& rhs); + __host__ __device__ friend cuvec4i operator/(const int& lhs, const cuvec4i& rhs); + __host__ __device__ friend cuvec4i operator-(const cuvec4i& other); + __host__ __device__ cuvec4i& operator+=(const cuvec4i& rhs); + __host__ __device__ cuvec4i& operator-=(const cuvec4i& rhs); + __host__ __device__ cuvec4i& operator*=(const int& rhs); + __host__ __device__ cuvec4i& operator/=(const int& rhs); + + }; }; //end namespace amscuda diff --git a/src/amsculib3/math/cuvec2i.cu b/src/amsculib3/math/cuvec2i.cu new file mode 100644 index 0000000..829f578 --- /dev/null +++ b/src/amsculib3/math/cuvec2i.cu @@ -0,0 +1,167 @@ +#include + +namespace amscuda +{ + +////////////////// +// Vector Class // +////////////////// + +__host__ __device__ cuvec2i::cuvec2i() +{ + x = 0; y = 0; + return; +} + +__host__ __device__ cuvec2i::~cuvec2i() +{ + x = 0; y = 0; + return; +} + +__host__ __device__ cuvec2i::cuvec2i(const int32_t &_x, const int32_t &_y) +{ + x = _x; y = _y; + return; +} + +__host__ __device__ int32_t& cuvec2i::operator[](const int &I) +{ + switch(I) + { + case 0: + return x; + case 1: + return y; + } + + return x; +} + +__host__ __device__ const int32_t& cuvec2i::operator[](const int &I) const +{ + switch(I) + { + case 0: + return x; + case 1: + return y; + } + + return x; +} + +__host__ __device__ cuvec2i cuvec2i::operator+(const cuvec2i& rhs) const +{ + cuvec2i ret; + ret.x = x + rhs.x; + ret.y = y + rhs.y; + return ret; +} + +__host__ __device__ cuvec2i cuvec2i::operator-(const cuvec2i& rhs) const +{ + cuvec2i ret; + ret.x = x - rhs.x; + ret.y = y - rhs.y; + return ret; +} + +__host__ __device__ cuvec2i cuvec2i::operator*(const cuvec2i& rhs) const +{ + //Elementwise product + cuvec2i ret; + ret.x = x * rhs.x; + ret.y = y * rhs.y; + return ret; +} + +__host__ __device__ cuvec2i cuvec2i::operator/(const cuvec2i& rhs) const +{ + //Elementwise division + cuvec2i ret; + ret.x = x / rhs.x; + ret.y = y / rhs.y; + return ret; +} + +__host__ __device__ cuvec2i operator*(const cuvec2i& lhs, const int32_t& rhs) +{ + cuvec2i ret; + ret.x = lhs.x*rhs; + ret.y = lhs.y*rhs; + return ret; +} + +__host__ __device__ cuvec2i operator*(const int32_t& lhs, const cuvec2i& rhs) +{ + cuvec2i ret; + ret.x = lhs*rhs.x; + ret.y = lhs*rhs.y; + return ret; +} + +__host__ __device__ cuvec2i operator/(const cuvec2i& lhs, const int32_t& rhs) +{ + cuvec2i ret; + ret.x = lhs.x/rhs; + ret.y = lhs.y/rhs; + return ret; +} + +__host__ __device__ cuvec2i operator/(const int32_t& lhs, const cuvec2i& rhs) +{ + cuvec2i ret; + ret.x = lhs/rhs.x; + ret.y = lhs/rhs.y; + return ret; +} + +__host__ __device__ cuvec2i operator-(const cuvec2i& other) +{ + cuvec2i ret; + ret.x = -other.x; + ret.y = -other.y; + return ret; +} + +__host__ __device__ cuvec2i& cuvec2i::operator+=(const cuvec2i& rhs) +{ + x += rhs.x; + y += rhs.y; + return *this; +} + +__host__ __device__ cuvec2i& cuvec2i::operator-=(const cuvec2i& rhs) +{ + x -= rhs.x; + y -= rhs.y; + return *this; +} + +__host__ __device__ cuvec2i& cuvec2i::operator*=(const int32_t& rhs) +{ + x *= rhs; + y *= rhs; + return *this; +} + +__host__ __device__ cuvec2i& cuvec2i::operator/=(const int32_t& rhs) +{ + x /= rhs; + y /= rhs; + return *this; +} + + + +////////////////////// +// Standalone Funcs // +////////////////////// + + +/////////// +// Tests // +/////////// + +}; \ No newline at end of file diff --git a/src/amsculib3/math/cuvec3i.cu b/src/amsculib3/math/cuvec3i.cu new file mode 100644 index 0000000..6e2d7e7 --- /dev/null +++ b/src/amsculib3/math/cuvec3i.cu @@ -0,0 +1,183 @@ +#include + +namespace amscuda +{ + +////////////////// +// Vector Class // +////////////////// + +__host__ __device__ cuvec3i::cuvec3i() +{ + x = 0; y = 0; z = 0; + return; +} + +__host__ __device__ cuvec3i::~cuvec3i() +{ + x = 0; y = 0; z = 0; + return; +} + +__host__ __device__ cuvec3i::cuvec3i(const int &_x, const int &_y, const int &_z) +{ + x = _x; y = _y; z = _z; + return; +} + +__host__ __device__ int& cuvec3i::operator[](const int &I) +{ + switch(I) + { + case 0: + return x; + case 1: + return y; + case 2: + return z; + } + + return x; +} + +__host__ __device__ const int& cuvec3i::operator[](const int &I) const +{ + switch(I) + { + case 0: + return x; + case 1: + return y; + case 2: + return z; + } + + return x; +} + +__host__ __device__ cuvec3i cuvec3i::operator+(const cuvec3i& rhs) const +{ + cuvec3i ret; + ret.x = x + rhs.x; + ret.y = y + rhs.y; + ret.z = z + rhs.z; + return ret; +} + +__host__ __device__ cuvec3i cuvec3i::operator-(const cuvec3i& rhs) const +{ + cuvec3i ret; + ret.x = x - rhs.x; + ret.y = y - rhs.y; + ret.z = z - rhs.z; + return ret; +} + +__host__ __device__ cuvec3i cuvec3i::operator*(const cuvec3i& rhs) const +{ + //Elementwise product + cuvec3i ret; + ret.x = x * rhs.x; + ret.y = y * rhs.y; + ret.z = z * rhs.z; + return ret; +} + +__host__ __device__ cuvec3i cuvec3i::operator/(const cuvec3i& rhs) const +{ + //Elementwise division + cuvec3i ret; + ret.x = x / rhs.x; + ret.y = y / rhs.y; + ret.z = z / rhs.z; + return ret; +} + +__host__ __device__ cuvec3i operator*(const cuvec3i& lhs, const int& rhs) +{ + cuvec3i ret; + ret.x = lhs.x*rhs; + ret.y = lhs.y*rhs; + ret.z = lhs.z*rhs; + return ret; +} + +__host__ __device__ cuvec3i operator*(const int& lhs, const cuvec3i& rhs) +{ + cuvec3i ret; + ret.x = lhs*rhs.x; + ret.y = lhs*rhs.y; + ret.z = lhs*rhs.z; + return ret; +} + +__host__ __device__ cuvec3i operator/(const cuvec3i& lhs, const int& rhs) +{ + cuvec3i ret; + ret.x = lhs.x/rhs; + ret.y = lhs.y/rhs; + ret.z = lhs.z/rhs; + return ret; +} + +__host__ __device__ cuvec3i operator/(const int& lhs, const cuvec3i& rhs) +{ + cuvec3i ret; + ret.x = lhs/rhs.x; + ret.y = lhs/rhs.y; + ret.z = lhs/rhs.z; + return ret; +} + +__host__ __device__ cuvec3i operator-(const cuvec3i& other) +{ + cuvec3i ret; + ret.x = -other.x; + ret.y = -other.y; + ret.z = -other.z; + return ret; +} + +__host__ __device__ cuvec3i& cuvec3i::operator+=(const cuvec3i& rhs) +{ + x += rhs.x; + y += rhs.y; + z += rhs.z; + return *this; +} + +__host__ __device__ cuvec3i& cuvec3i::operator-=(const cuvec3i& rhs) +{ + x -= rhs.x; + y -= rhs.y; + z -= rhs.z; + return *this; +} + +__host__ __device__ cuvec3i& cuvec3i::operator*=(const int& rhs) +{ + x *= rhs; + y *= rhs; + z *= rhs; + return *this; +} + +__host__ __device__ cuvec3i& cuvec3i::operator/=(const int& rhs) +{ + x /= rhs; + y /= rhs; + z /= rhs; + return *this; +} + + +////////////////////// +// Standalone Funcs // +////////////////////// + + +/////////// +// Tests // +/////////// + +}; \ No newline at end of file diff --git a/src/amsculib3/math/cuvec4i.cu b/src/amsculib3/math/cuvec4i.cu new file mode 100644 index 0000000..7de0e66 --- /dev/null +++ b/src/amsculib3/math/cuvec4i.cu @@ -0,0 +1,200 @@ +#include + +namespace amscuda +{ + +////////////////// +// Vector Class // +////////////////// + +__host__ __device__ cuvec4i::cuvec4i() +{ + x = 0; y = 0; z = 0; w = 0; + return; +} + +__host__ __device__ cuvec4i::~cuvec4i() +{ + x = 0; y = 0; z = 0; w = 0; + return; +} + +__host__ __device__ cuvec4i::cuvec4i(const int &_x, const int &_y, const int &_z, const int &_w) +{ + x = _x; y = _y; z = _z; w = _w; + return; +} + +__host__ __device__ int& cuvec4i::operator[](const int &I) +{ + switch(I) + { + case 0: + return x; + case 1: + return y; + case 2: + return z; + case 3: + return w; + } + + return x; +} + +__host__ __device__ const int& cuvec4i::operator[](const int &I) const +{ + switch(I) + { + case 0: + return x; + case 1: + return y; + case 2: + return z; + case 3: + return w; + } + + return x; +} + +__host__ __device__ cuvec4i cuvec4i::operator+(const cuvec4i& rhs) const +{ + cuvec4i ret; + ret.x = x + rhs.x; + ret.y = y + rhs.y; + ret.z = z + rhs.z; + ret.w = w + rhs.w; + return ret; +} + +__host__ __device__ cuvec4i cuvec4i::operator-(const cuvec4i& rhs) const +{ + cuvec4i ret; + ret.x = x - rhs.x; + ret.y = y - rhs.y; + ret.z = z - rhs.z; + ret.w = w - rhs.w; + return ret; +} + +__host__ __device__ cuvec4i cuvec4i::operator*(const cuvec4i& rhs) const +{ + //Elementwise product + cuvec4i ret; + ret.x = x * rhs.x; + ret.y = y * rhs.y; + ret.z = z * rhs.z; + ret.w = w * rhs.w; + return ret; +} + +__host__ __device__ cuvec4i cuvec4i::operator/(const cuvec4i& rhs) const +{ + //Elementwise division + cuvec4i ret; + ret.x = x / rhs.x; + ret.y = y / rhs.y; + ret.z = z / rhs.z; + ret.w = w / rhs.w; + return ret; +} + +__host__ __device__ cuvec4i operator*(const cuvec4i& lhs, const int& rhs) +{ + cuvec4i ret; + ret.x = lhs.x*rhs; + ret.y = lhs.y*rhs; + ret.z = lhs.z*rhs; + ret.w = lhs.w*rhs; + return ret; +} + +__host__ __device__ cuvec4i operator*(const int& lhs, const cuvec4i& rhs) +{ + cuvec4i ret; + ret.x = lhs*rhs.x; + ret.y = lhs*rhs.y; + ret.z = lhs*rhs.z; + ret.w = lhs*rhs.w; + return ret; +} + +__host__ __device__ cuvec4i operator/(const cuvec4i& lhs, const int& rhs) +{ + cuvec4i ret; + ret.x = lhs.x/rhs; + ret.y = lhs.y/rhs; + ret.z = lhs.z/rhs; + ret.w = lhs.w/rhs; + return ret; +} + +__host__ __device__ cuvec4i operator/(const int& lhs, const cuvec4i& rhs) +{ + cuvec4i ret; + ret.x = lhs/rhs.x; + ret.y = lhs/rhs.y; + ret.z = lhs/rhs.z; + ret.w = lhs/rhs.w; + return ret; +} + +__host__ __device__ cuvec4i operator-(const cuvec4i& other) +{ + cuvec4i ret; + ret.x = -other.x; + ret.y = -other.y; + ret.z = -other.z; + ret.w = -other.w; + return ret; +} + +__host__ __device__ cuvec4i& cuvec4i::operator+=(const cuvec4i& rhs) +{ + x += rhs.x; + y += rhs.y; + z += rhs.z; + w += rhs.w; + return *this; +} + +__host__ __device__ cuvec4i& cuvec4i::operator-=(const cuvec4i& rhs) +{ + x -= rhs.x; + y -= rhs.y; + z -= rhs.z; + w -= rhs.w; + return *this; +} + +__host__ __device__ cuvec4i& cuvec4i::operator*=(const int& rhs) +{ + x *= rhs; + y *= rhs; + z *= rhs; + w *= rhs; + return *this; +} + +__host__ __device__ cuvec4i& cuvec4i::operator/=(const int& rhs) +{ + x /= rhs; + y /= rhs; + z /= rhs; + w /= rhs; + return *this; +} + + +////////////////////// +// Standalone Funcs // +////////////////////// + + +/////////// +// Tests // +/////////// + +}; \ No newline at end of file diff --git a/test_scripts/cuvec_codegen.py b/test_scripts/cuvec_codegen.py index 134d892..23efc97 100644 --- a/test_scripts/cuvec_codegen.py +++ b/test_scripts/cuvec_codegen.py @@ -155,9 +155,9 @@ def v_operatorbk1(dim,dtype): lnss = "" vname = vtypename(dim,dtype) - lnsh += "\t{} float& operator[](const int &I);\n".format(cudadec) + lnsh += f"\t{cudadec} {dtype}& operator[](const int &I);\n" - lnss += "\t{} float& {}::operator[](const int &I)\n".format(cudadec,vname) + lnss += f"\t{cudadec} {dtype}& {vname}::operator[](const int &I)\n" lnss += "\t{\n" lnss += "\t\tswitch(I)\n\t\t{\n" for I in range(0,dim): @@ -174,9 +174,9 @@ def v_operatorbk2(dim,dtype): lnss = "" vname = vtypename(dim,dtype) - lnsh += "\t{} const float& operator[](const int &I) const;\n".format(cudadec) + lnsh += f"\t{cudadec} const {dtype}& operator[](const int &I) const;\n" - lnss += "\t{} const float& {}::operator[](const int &I) const\n".format(cudadec,vname) + lnss += f"\t{cudadec} const {dtype}& {vname}::operator[](const int &I) const\n" lnss += "\t{\n" lnss += "\t\tswitch(I)\n\t\t{\n" for I in range(0,dim): @@ -580,9 +580,9 @@ def moperatorbk1(dim,dtype): lnss = "" mname = mtypename(dim,dtype) - lnsh += f"\t{cudadec} float& operator[](const int &I);\n" + lnsh += f"\t{cudadec} {dtype}& operator[](const int &I);\n" - lnss += f"\t{cudadec} float& {mname}::operator[](const int &I)\n" + lnss += f"\t{cudadec} {dtype}& {mname}::operator[](const int &I)\n" lnss += "\t{\n" lnss += "\t\tswitch(I)\n\t\t{\n" for I in range(0,dim*dim): @@ -601,9 +601,9 @@ def moperatorbk2(dim,dtype): lnss = "" mname = mtypename(dim,dtype) - lnsh += f"\t{cudadec} const float& operator[](const int &I) const;\n" + lnsh += f"\t{cudadec} const {dtype}& operator[](const int &I) const;\n" - lnss += f"\t{cudadec} const float& {mname}::operator[](const int &I) const\n" + lnss += f"\t{cudadec} const {dtype}& {mname}::operator[](const int &I) const\n" lnss += "\t{\n" lnss += "\t\tswitch(I)\n\t\t{\n" for I in range(0,dim*dim): @@ -622,9 +622,9 @@ def moperatorbk3(dim,dtype): lnss = "" mname = mtypename(dim,dtype) - lnsh += f"\t{cudadec} float& operator()(const int &I, const int &J);\n" + lnsh += f"\t{cudadec} {dtype}& operator()(const int &I, const int &J);\n" - lnss += f"\t{cudadec} float& {mname}::operator()(const int &I, const int &J)\n" + lnss += f"\t{cudadec} {dtype}& {mname}::operator()(const int &I, const int &J)\n" lnss += "\t{\n" lnss += f"\t\treturn (*this)[I+{dim}*J];\n" lnss += "\t}\n\n" @@ -636,9 +636,9 @@ def moperatorbk4(dim,dtype): lnss = "" mname = mtypename(dim,dtype) - lnsh += f"\t{cudadec} const float& operator()(const int &I, const int &J) const;\n" + lnsh += f"\t{cudadec} const {dtype}& operator()(const int &I, const int &J) const;\n" - lnss += f"\t{cudadec} const float& {mname}::operator()(const int &I, const int &J) const\n" + lnss += f"\t{cudadec} const {dtype}& {mname}::operator()(const int &I, const int &J) const\n" lnss += "\t{\n" lnss += f"\t\treturn (*this)[I+{dim}*J];\n" lnss += "\t}\n\n" @@ -650,9 +650,9 @@ def moperatorbk5(dim,dtype): lnss = "" mname = mtypename(dim,dtype) - lnsh += f"\t{cudadec} float& at(const int &I, const int &J);\n" + lnsh += f"\t{cudadec} {dtype}& at(const int &I, const int &J);\n" - lnss += f"\t{cudadec} float& {mname}::at(const int &I, const int &J)\n" + lnss += f"\t{cudadec} {dtype}& {mname}::at(const int &I, const int &J)\n" lnss += "\t{\n" lnss += f"\t\treturn (*this)[I+{dim}*J];\n" lnss += "\t}\n\n" @@ -664,9 +664,9 @@ def moperatorbk6(dim,dtype): lnss = "" mname = mtypename(dim,dtype) - lnsh += f"\t{cudadec} const float& at(const int &I, const int &J) const;\n" + lnsh += f"\t{cudadec} const {dtype}& at(const int &I, const int &J) const;\n" - lnss += f"\t{cudadec} const float& {mname}::at(const int &I, const int &J) const\n" + lnss += f"\t{cudadec} const {dtype}& {mname}::at(const int &I, const int &J) const\n" lnss += "\t{\n" lnss += f"\t\treturn (*this)[I+{dim}*J];\n" lnss += "\t}\n\n" @@ -680,7 +680,7 @@ def mdataoperator1(dim,dtype): lnss = "" mname = mtypename(dim,dtype) - lnsh += f"{cudadec} {dtype}* data(); //pointer to float{dim*dim} representation of matrix\n" + lnsh += f"{cudadec} {dtype}* data(); //pointer to {dtype}{dim*dim} representation of matrix\n" lnss += f"{cudadec} {dtype}* {mname}::data()\n" lnss += "{\n" @@ -694,7 +694,7 @@ def mdataoperator2(dim,dtype): lnss = "" mname = mtypename(dim,dtype) - lnsh += f"{cudadec} const {dtype}* data() const; //pointer to float{dim*dim} representation of matrix\n" + lnsh += f"{cudadec} const {dtype}* data() const; //pointer to {dtype}{dim*dim} representation of matrix\n" lnss += f"{cudadec} const {dtype}* {mname}::data() const\n" lnss += "{\n" diff --git a/test_scripts/cuvec_codegen/cuvec2_codegen1.cu b/test_scripts/cuvec_codegen/cuvec2_codegen1.cu index cf7349f..e7f8dda 100644 --- a/test_scripts/cuvec_codegen/cuvec2_codegen1.cu +++ b/test_scripts/cuvec_codegen/cuvec2_codegen1.cu @@ -16,7 +16,7 @@ return; } - __host__ __device__ float& cuvec2::operator[](const int &I) + __host__ __device__ double& cuvec2::operator[](const int &I) { switch(I) { @@ -29,7 +29,7 @@ return x; } - __host__ __device__ const float& cuvec2::operator[](const int &I) const + __host__ __device__ const double& cuvec2::operator[](const int &I) const { switch(I) { @@ -198,7 +198,7 @@ return; } - __host__ __device__ float& cumat2::operator[](const int &I) + __host__ __device__ double& cumat2::operator[](const int &I) { switch(I) { @@ -215,7 +215,7 @@ return m00; } - __host__ __device__ const float& cumat2::operator[](const int &I) const + __host__ __device__ const double& cumat2::operator[](const int &I) const { switch(I) { @@ -232,22 +232,22 @@ return m00; } - __host__ __device__ float& cumat2::operator()(const int &I, const int &J) + __host__ __device__ double& cumat2::operator()(const int &I, const int &J) { return (*this)[I+2*J]; } - __host__ __device__ const float& cumat2::operator()(const int &I, const int &J) const + __host__ __device__ const double& cumat2::operator()(const int &I, const int &J) const { return (*this)[I+2*J]; } - __host__ __device__ float& cumat2::at(const int &I, const int &J) + __host__ __device__ double& cumat2::at(const int &I, const int &J) { return (*this)[I+2*J]; } - __host__ __device__ const float& cumat2::at(const int &I, const int &J) const + __host__ __device__ const double& cumat2::at(const int &I, const int &J) const { return (*this)[I+2*J]; } diff --git a/test_scripts/cuvec_codegen/cuvec2_codegen1.hpp b/test_scripts/cuvec_codegen/cuvec2_codegen1.hpp index 74f2e44..4cb56e4 100644 --- a/test_scripts/cuvec_codegen/cuvec2_codegen1.hpp +++ b/test_scripts/cuvec_codegen/cuvec2_codegen1.hpp @@ -1,8 +1,8 @@ __host__ __device__ cuvec2(); __host__ __device__ ~cuvec2(); __host__ __device__ cuvec2(const double &_x, const double &_y); - __host__ __device__ float& operator[](const int &I); - __host__ __device__ const float& operator[](const int &I) const; + __host__ __device__ double& operator[](const int &I); + __host__ __device__ const double& operator[](const int &I) const; __host__ __device__ cuvec2 operator+(const cuvec2& rhs) const; __host__ __device__ cuvec2 operator-(const cuvec2& rhs) const; __host__ __device__ cuvec2 operator*(const cuvec2& rhs) const; //elementwise product @@ -29,14 +29,14 @@ double m01,m11; const double& _m10, const double& _m11 ); __host__ __device__ cumat2(const double* data4); - __host__ __device__ float& operator[](const int &I); - __host__ __device__ const float& operator[](const int &I) const; - __host__ __device__ float& operator()(const int &I, const int &J); - __host__ __device__ const float& operator()(const int &I, const int &J) const; - __host__ __device__ float& at(const int &I, const int &J); - __host__ __device__ const float& at(const int &I, const int &J) const; -__host__ __device__ double* data(); //pointer to float4 representation of matrix -__host__ __device__ const double* data() const; //pointer to float4 representation of matrix + __host__ __device__ double& operator[](const int &I); + __host__ __device__ const double& operator[](const int &I) const; + __host__ __device__ double& operator()(const int &I, const int &J); + __host__ __device__ const double& operator()(const int &I, const int &J) const; + __host__ __device__ double& at(const int &I, const int &J); + __host__ __device__ const double& at(const int &I, const int &J) const; +__host__ __device__ double* data(); //pointer to double4 representation of matrix +__host__ __device__ const double* data() const; //pointer to double4 representation of matrix __host__ __device__ cumat2 operator+(const cumat2& rhs) const; __host__ __device__ cumat2 operator-(const cumat2& rhs) const; __host__ __device__ cumat2 operator*(const cumat2& rhs) const; diff --git a/test_scripts/cuvec_codegen/cuvec2i_codegen1.cu b/test_scripts/cuvec_codegen/cuvec2i_codegen1.cu index 57ea2c3..ad58ebe 100644 --- a/test_scripts/cuvec_codegen/cuvec2i_codegen1.cu +++ b/test_scripts/cuvec_codegen/cuvec2i_codegen1.cu @@ -16,7 +16,7 @@ return; } - __host__ __device__ float& cuvec2i::operator[](const int &I) + __host__ __device__ int& cuvec2i::operator[](const int &I) { switch(I) { @@ -29,7 +29,7 @@ return x; } - __host__ __device__ const float& cuvec2i::operator[](const int &I) const + __host__ __device__ const int& cuvec2i::operator[](const int &I) const { switch(I) { @@ -198,7 +198,7 @@ return; } - __host__ __device__ float& cumat2i::operator[](const int &I) + __host__ __device__ int& cumat2i::operator[](const int &I) { switch(I) { @@ -215,7 +215,7 @@ return m00; } - __host__ __device__ const float& cumat2i::operator[](const int &I) const + __host__ __device__ const int& cumat2i::operator[](const int &I) const { switch(I) { @@ -232,22 +232,22 @@ return m00; } - __host__ __device__ float& cumat2i::operator()(const int &I, const int &J) + __host__ __device__ int& cumat2i::operator()(const int &I, const int &J) { return (*this)[I+2*J]; } - __host__ __device__ const float& cumat2i::operator()(const int &I, const int &J) const + __host__ __device__ const int& cumat2i::operator()(const int &I, const int &J) const { return (*this)[I+2*J]; } - __host__ __device__ float& cumat2i::at(const int &I, const int &J) + __host__ __device__ int& cumat2i::at(const int &I, const int &J) { return (*this)[I+2*J]; } - __host__ __device__ const float& cumat2i::at(const int &I, const int &J) const + __host__ __device__ const int& cumat2i::at(const int &I, const int &J) const { return (*this)[I+2*J]; } diff --git a/test_scripts/cuvec_codegen/cuvec2i_codegen1.hpp b/test_scripts/cuvec_codegen/cuvec2i_codegen1.hpp index 1ba6b12..1b84e78 100644 --- a/test_scripts/cuvec_codegen/cuvec2i_codegen1.hpp +++ b/test_scripts/cuvec_codegen/cuvec2i_codegen1.hpp @@ -1,8 +1,8 @@ __host__ __device__ cuvec2i(); __host__ __device__ ~cuvec2i(); __host__ __device__ cuvec2i(const int &_x, const int &_y); - __host__ __device__ float& operator[](const int &I); - __host__ __device__ const float& operator[](const int &I) const; + __host__ __device__ int& operator[](const int &I); + __host__ __device__ const int& operator[](const int &I) const; __host__ __device__ cuvec2i operator+(const cuvec2i& rhs) const; __host__ __device__ cuvec2i operator-(const cuvec2i& rhs) const; __host__ __device__ cuvec2i operator*(const cuvec2i& rhs) const; //elementwise product @@ -29,14 +29,14 @@ int m01,m11; const int& _m10, const int& _m11 ); __host__ __device__ cumat2i(const int* data4); - __host__ __device__ float& operator[](const int &I); - __host__ __device__ const float& operator[](const int &I) const; - __host__ __device__ float& operator()(const int &I, const int &J); - __host__ __device__ const float& operator()(const int &I, const int &J) const; - __host__ __device__ float& at(const int &I, const int &J); - __host__ __device__ const float& at(const int &I, const int &J) const; -__host__ __device__ int* data(); //pointer to float4 representation of matrix -__host__ __device__ const int* data() const; //pointer to float4 representation of matrix + __host__ __device__ int& operator[](const int &I); + __host__ __device__ const int& operator[](const int &I) const; + __host__ __device__ int& operator()(const int &I, const int &J); + __host__ __device__ const int& operator()(const int &I, const int &J) const; + __host__ __device__ int& at(const int &I, const int &J); + __host__ __device__ const int& at(const int &I, const int &J) const; +__host__ __device__ int* data(); //pointer to int4 representation of matrix +__host__ __device__ const int* data() const; //pointer to int4 representation of matrix __host__ __device__ cumat2i operator+(const cumat2i& rhs) const; __host__ __device__ cumat2i operator-(const cumat2i& rhs) const; __host__ __device__ cumat2i operator*(const cumat2i& rhs) const; diff --git a/test_scripts/cuvec_codegen/cuvec3_codegen1.cu b/test_scripts/cuvec_codegen/cuvec3_codegen1.cu index 6b2faff..d9e5711 100644 --- a/test_scripts/cuvec_codegen/cuvec3_codegen1.cu +++ b/test_scripts/cuvec_codegen/cuvec3_codegen1.cu @@ -16,7 +16,7 @@ return; } - __host__ __device__ float& cuvec3::operator[](const int &I) + __host__ __device__ double& cuvec3::operator[](const int &I) { switch(I) { @@ -31,7 +31,7 @@ return x; } - __host__ __device__ const float& cuvec3::operator[](const int &I) const + __host__ __device__ const double& cuvec3::operator[](const int &I) const { switch(I) { @@ -240,7 +240,7 @@ return; } - __host__ __device__ float& cumat3::operator[](const int &I) + __host__ __device__ double& cumat3::operator[](const int &I) { switch(I) { @@ -267,7 +267,7 @@ return m00; } - __host__ __device__ const float& cumat3::operator[](const int &I) const + __host__ __device__ const double& cumat3::operator[](const int &I) const { switch(I) { @@ -294,22 +294,22 @@ return m00; } - __host__ __device__ float& cumat3::operator()(const int &I, const int &J) + __host__ __device__ double& cumat3::operator()(const int &I, const int &J) { return (*this)[I+3*J]; } - __host__ __device__ const float& cumat3::operator()(const int &I, const int &J) const + __host__ __device__ const double& cumat3::operator()(const int &I, const int &J) const { return (*this)[I+3*J]; } - __host__ __device__ float& cumat3::at(const int &I, const int &J) + __host__ __device__ double& cumat3::at(const int &I, const int &J) { return (*this)[I+3*J]; } - __host__ __device__ const float& cumat3::at(const int &I, const int &J) const + __host__ __device__ const double& cumat3::at(const int &I, const int &J) const { return (*this)[I+3*J]; } diff --git a/test_scripts/cuvec_codegen/cuvec3_codegen1.hpp b/test_scripts/cuvec_codegen/cuvec3_codegen1.hpp index d83bd0e..420855e 100644 --- a/test_scripts/cuvec_codegen/cuvec3_codegen1.hpp +++ b/test_scripts/cuvec_codegen/cuvec3_codegen1.hpp @@ -1,8 +1,8 @@ __host__ __device__ cuvec3(); __host__ __device__ ~cuvec3(); __host__ __device__ cuvec3(const double &_x, const double &_y, const double &_z); - __host__ __device__ float& operator[](const int &I); - __host__ __device__ const float& operator[](const int &I) const; + __host__ __device__ double& operator[](const int &I); + __host__ __device__ const double& operator[](const int &I) const; __host__ __device__ cuvec3 operator+(const cuvec3& rhs) const; __host__ __device__ cuvec3 operator-(const cuvec3& rhs) const; __host__ __device__ cuvec3 operator*(const cuvec3& rhs) const; //elementwise product @@ -31,14 +31,14 @@ double m02,m12,m22; const double& _m20, const double& _m21, const double& _m22 ); __host__ __device__ cumat3(const double* data9); - __host__ __device__ float& operator[](const int &I); - __host__ __device__ const float& operator[](const int &I) const; - __host__ __device__ float& operator()(const int &I, const int &J); - __host__ __device__ const float& operator()(const int &I, const int &J) const; - __host__ __device__ float& at(const int &I, const int &J); - __host__ __device__ const float& at(const int &I, const int &J) const; -__host__ __device__ double* data(); //pointer to float9 representation of matrix -__host__ __device__ const double* data() const; //pointer to float9 representation of matrix + __host__ __device__ double& operator[](const int &I); + __host__ __device__ const double& operator[](const int &I) const; + __host__ __device__ double& operator()(const int &I, const int &J); + __host__ __device__ const double& operator()(const int &I, const int &J) const; + __host__ __device__ double& at(const int &I, const int &J); + __host__ __device__ const double& at(const int &I, const int &J) const; +__host__ __device__ double* data(); //pointer to double9 representation of matrix +__host__ __device__ const double* data() const; //pointer to double9 representation of matrix __host__ __device__ cumat3 operator+(const cumat3& rhs) const; __host__ __device__ cumat3 operator-(const cumat3& rhs) const; __host__ __device__ cumat3 operator*(const cumat3& rhs) const; diff --git a/test_scripts/cuvec_codegen/cuvec3i_codegen1.cu b/test_scripts/cuvec_codegen/cuvec3i_codegen1.cu index 8417071..48b11c4 100644 --- a/test_scripts/cuvec_codegen/cuvec3i_codegen1.cu +++ b/test_scripts/cuvec_codegen/cuvec3i_codegen1.cu @@ -16,7 +16,7 @@ return; } - __host__ __device__ float& cuvec3i::operator[](const int &I) + __host__ __device__ int& cuvec3i::operator[](const int &I) { switch(I) { @@ -31,7 +31,7 @@ return x; } - __host__ __device__ const float& cuvec3i::operator[](const int &I) const + __host__ __device__ const int& cuvec3i::operator[](const int &I) const { switch(I) { @@ -240,7 +240,7 @@ return; } - __host__ __device__ float& cumat3i::operator[](const int &I) + __host__ __device__ int& cumat3i::operator[](const int &I) { switch(I) { @@ -267,7 +267,7 @@ return m00; } - __host__ __device__ const float& cumat3i::operator[](const int &I) const + __host__ __device__ const int& cumat3i::operator[](const int &I) const { switch(I) { @@ -294,22 +294,22 @@ return m00; } - __host__ __device__ float& cumat3i::operator()(const int &I, const int &J) + __host__ __device__ int& cumat3i::operator()(const int &I, const int &J) { return (*this)[I+3*J]; } - __host__ __device__ const float& cumat3i::operator()(const int &I, const int &J) const + __host__ __device__ const int& cumat3i::operator()(const int &I, const int &J) const { return (*this)[I+3*J]; } - __host__ __device__ float& cumat3i::at(const int &I, const int &J) + __host__ __device__ int& cumat3i::at(const int &I, const int &J) { return (*this)[I+3*J]; } - __host__ __device__ const float& cumat3i::at(const int &I, const int &J) const + __host__ __device__ const int& cumat3i::at(const int &I, const int &J) const { return (*this)[I+3*J]; } diff --git a/test_scripts/cuvec_codegen/cuvec3i_codegen1.hpp b/test_scripts/cuvec_codegen/cuvec3i_codegen1.hpp index 53a1936..5e6db5a 100644 --- a/test_scripts/cuvec_codegen/cuvec3i_codegen1.hpp +++ b/test_scripts/cuvec_codegen/cuvec3i_codegen1.hpp @@ -1,8 +1,8 @@ __host__ __device__ cuvec3i(); __host__ __device__ ~cuvec3i(); __host__ __device__ cuvec3i(const int &_x, const int &_y, const int &_z); - __host__ __device__ float& operator[](const int &I); - __host__ __device__ const float& operator[](const int &I) const; + __host__ __device__ int& operator[](const int &I); + __host__ __device__ const int& operator[](const int &I) const; __host__ __device__ cuvec3i operator+(const cuvec3i& rhs) const; __host__ __device__ cuvec3i operator-(const cuvec3i& rhs) const; __host__ __device__ cuvec3i operator*(const cuvec3i& rhs) const; //elementwise product @@ -31,14 +31,14 @@ int m02,m12,m22; const int& _m20, const int& _m21, const int& _m22 ); __host__ __device__ cumat3i(const int* data9); - __host__ __device__ float& operator[](const int &I); - __host__ __device__ const float& operator[](const int &I) const; - __host__ __device__ float& operator()(const int &I, const int &J); - __host__ __device__ const float& operator()(const int &I, const int &J) const; - __host__ __device__ float& at(const int &I, const int &J); - __host__ __device__ const float& at(const int &I, const int &J) const; -__host__ __device__ int* data(); //pointer to float9 representation of matrix -__host__ __device__ const int* data() const; //pointer to float9 representation of matrix + __host__ __device__ int& operator[](const int &I); + __host__ __device__ const int& operator[](const int &I) const; + __host__ __device__ int& operator()(const int &I, const int &J); + __host__ __device__ const int& operator()(const int &I, const int &J) const; + __host__ __device__ int& at(const int &I, const int &J); + __host__ __device__ const int& at(const int &I, const int &J) const; +__host__ __device__ int* data(); //pointer to int9 representation of matrix +__host__ __device__ const int* data() const; //pointer to int9 representation of matrix __host__ __device__ cumat3i operator+(const cumat3i& rhs) const; __host__ __device__ cumat3i operator-(const cumat3i& rhs) const; __host__ __device__ cumat3i operator*(const cumat3i& rhs) const; diff --git a/test_scripts/cuvec_codegen/cuvec4_codegen1.cu b/test_scripts/cuvec_codegen/cuvec4_codegen1.cu index 7884c7e..ed3f59a 100644 --- a/test_scripts/cuvec_codegen/cuvec4_codegen1.cu +++ b/test_scripts/cuvec_codegen/cuvec4_codegen1.cu @@ -16,7 +16,7 @@ return; } - __host__ __device__ float& cuvec4::operator[](const int &I) + __host__ __device__ double& cuvec4::operator[](const int &I) { switch(I) { @@ -33,7 +33,7 @@ return x; } - __host__ __device__ const float& cuvec4::operator[](const int &I) const + __host__ __device__ const double& cuvec4::operator[](const int &I) const { switch(I) { @@ -290,7 +290,7 @@ return; } - __host__ __device__ float& cumat4::operator[](const int &I) + __host__ __device__ double& cumat4::operator[](const int &I) { switch(I) { @@ -331,7 +331,7 @@ return m00; } - __host__ __device__ const float& cumat4::operator[](const int &I) const + __host__ __device__ const double& cumat4::operator[](const int &I) const { switch(I) { @@ -372,22 +372,22 @@ return m00; } - __host__ __device__ float& cumat4::operator()(const int &I, const int &J) + __host__ __device__ double& cumat4::operator()(const int &I, const int &J) { return (*this)[I+4*J]; } - __host__ __device__ const float& cumat4::operator()(const int &I, const int &J) const + __host__ __device__ const double& cumat4::operator()(const int &I, const int &J) const { return (*this)[I+4*J]; } - __host__ __device__ float& cumat4::at(const int &I, const int &J) + __host__ __device__ double& cumat4::at(const int &I, const int &J) { return (*this)[I+4*J]; } - __host__ __device__ const float& cumat4::at(const int &I, const int &J) const + __host__ __device__ const double& cumat4::at(const int &I, const int &J) const { return (*this)[I+4*J]; } diff --git a/test_scripts/cuvec_codegen/cuvec4_codegen1.hpp b/test_scripts/cuvec_codegen/cuvec4_codegen1.hpp index 1c61be8..b94c463 100644 --- a/test_scripts/cuvec_codegen/cuvec4_codegen1.hpp +++ b/test_scripts/cuvec_codegen/cuvec4_codegen1.hpp @@ -1,8 +1,8 @@ __host__ __device__ cuvec4(); __host__ __device__ ~cuvec4(); __host__ __device__ cuvec4(const double &_x, const double &_y, const double &_z, const double &_w); - __host__ __device__ float& operator[](const int &I); - __host__ __device__ const float& operator[](const int &I) const; + __host__ __device__ double& operator[](const int &I); + __host__ __device__ const double& operator[](const int &I) const; __host__ __device__ cuvec4 operator+(const cuvec4& rhs) const; __host__ __device__ cuvec4 operator-(const cuvec4& rhs) const; __host__ __device__ cuvec4 operator*(const cuvec4& rhs) const; //elementwise product @@ -33,14 +33,14 @@ double m03,m13,m23,m33; const double& _m30, const double& _m31, const double& _m32, const double& _m33 ); __host__ __device__ cumat4(const double* data16); - __host__ __device__ float& operator[](const int &I); - __host__ __device__ const float& operator[](const int &I) const; - __host__ __device__ float& operator()(const int &I, const int &J); - __host__ __device__ const float& operator()(const int &I, const int &J) const; - __host__ __device__ float& at(const int &I, const int &J); - __host__ __device__ const float& at(const int &I, const int &J) const; -__host__ __device__ double* data(); //pointer to float16 representation of matrix -__host__ __device__ const double* data() const; //pointer to float16 representation of matrix + __host__ __device__ double& operator[](const int &I); + __host__ __device__ const double& operator[](const int &I) const; + __host__ __device__ double& operator()(const int &I, const int &J); + __host__ __device__ const double& operator()(const int &I, const int &J) const; + __host__ __device__ double& at(const int &I, const int &J); + __host__ __device__ const double& at(const int &I, const int &J) const; +__host__ __device__ double* data(); //pointer to double16 representation of matrix +__host__ __device__ const double* data() const; //pointer to double16 representation of matrix __host__ __device__ cumat4 operator+(const cumat4& rhs) const; __host__ __device__ cumat4 operator-(const cumat4& rhs) const; __host__ __device__ cumat4 operator*(const cumat4& rhs) const; diff --git a/test_scripts/cuvec_codegen/cuvec4i_codegen1.cu b/test_scripts/cuvec_codegen/cuvec4i_codegen1.cu index 3d38904..736d65d 100644 --- a/test_scripts/cuvec_codegen/cuvec4i_codegen1.cu +++ b/test_scripts/cuvec_codegen/cuvec4i_codegen1.cu @@ -16,7 +16,7 @@ return; } - __host__ __device__ float& cuvec4i::operator[](const int &I) + __host__ __device__ int& cuvec4i::operator[](const int &I) { switch(I) { @@ -33,7 +33,7 @@ return x; } - __host__ __device__ const float& cuvec4i::operator[](const int &I) const + __host__ __device__ const int& cuvec4i::operator[](const int &I) const { switch(I) { @@ -290,7 +290,7 @@ return; } - __host__ __device__ float& cumat4i::operator[](const int &I) + __host__ __device__ int& cumat4i::operator[](const int &I) { switch(I) { @@ -331,7 +331,7 @@ return m00; } - __host__ __device__ const float& cumat4i::operator[](const int &I) const + __host__ __device__ const int& cumat4i::operator[](const int &I) const { switch(I) { @@ -372,22 +372,22 @@ return m00; } - __host__ __device__ float& cumat4i::operator()(const int &I, const int &J) + __host__ __device__ int& cumat4i::operator()(const int &I, const int &J) { return (*this)[I+4*J]; } - __host__ __device__ const float& cumat4i::operator()(const int &I, const int &J) const + __host__ __device__ const int& cumat4i::operator()(const int &I, const int &J) const { return (*this)[I+4*J]; } - __host__ __device__ float& cumat4i::at(const int &I, const int &J) + __host__ __device__ int& cumat4i::at(const int &I, const int &J) { return (*this)[I+4*J]; } - __host__ __device__ const float& cumat4i::at(const int &I, const int &J) const + __host__ __device__ const int& cumat4i::at(const int &I, const int &J) const { return (*this)[I+4*J]; } diff --git a/test_scripts/cuvec_codegen/cuvec4i_codegen1.hpp b/test_scripts/cuvec_codegen/cuvec4i_codegen1.hpp index c66849e..91f183b 100644 --- a/test_scripts/cuvec_codegen/cuvec4i_codegen1.hpp +++ b/test_scripts/cuvec_codegen/cuvec4i_codegen1.hpp @@ -1,8 +1,8 @@ __host__ __device__ cuvec4i(); __host__ __device__ ~cuvec4i(); __host__ __device__ cuvec4i(const int &_x, const int &_y, const int &_z, const int &_w); - __host__ __device__ float& operator[](const int &I); - __host__ __device__ const float& operator[](const int &I) const; + __host__ __device__ int& operator[](const int &I); + __host__ __device__ const int& operator[](const int &I) const; __host__ __device__ cuvec4i operator+(const cuvec4i& rhs) const; __host__ __device__ cuvec4i operator-(const cuvec4i& rhs) const; __host__ __device__ cuvec4i operator*(const cuvec4i& rhs) const; //elementwise product @@ -33,14 +33,14 @@ int m03,m13,m23,m33; const int& _m30, const int& _m31, const int& _m32, const int& _m33 ); __host__ __device__ cumat4i(const int* data16); - __host__ __device__ float& operator[](const int &I); - __host__ __device__ const float& operator[](const int &I) const; - __host__ __device__ float& operator()(const int &I, const int &J); - __host__ __device__ const float& operator()(const int &I, const int &J) const; - __host__ __device__ float& at(const int &I, const int &J); - __host__ __device__ const float& at(const int &I, const int &J) const; -__host__ __device__ int* data(); //pointer to float16 representation of matrix -__host__ __device__ const int* data() const; //pointer to float16 representation of matrix + __host__ __device__ int& operator[](const int &I); + __host__ __device__ const int& operator[](const int &I) const; + __host__ __device__ int& operator()(const int &I, const int &J); + __host__ __device__ const int& operator()(const int &I, const int &J) const; + __host__ __device__ int& at(const int &I, const int &J); + __host__ __device__ const int& at(const int &I, const int &J) const; +__host__ __device__ int* data(); //pointer to int16 representation of matrix +__host__ __device__ const int* data() const; //pointer to int16 representation of matrix __host__ __device__ cumat4i operator+(const cumat4i& rhs) const; __host__ __device__ cumat4i operator-(const cumat4i& rhs) const; __host__ __device__ cumat4i operator*(const cumat4i& rhs) const;