From 710e2f6108ed9fc063bf138902210b11e1199416 Mon Sep 17 00:00:00 2001 From: Ramakrishna Prabhu Date: Thu, 8 Jul 2021 22:45:32 -0500 Subject: [PATCH 01/26] code changes and test cases --- cpp/src/io/orc/reader_impl.cu | 7 ++++++- cpp/src/io/orc/stripe_data.cu | 30 ++++++++++++++---------------- python/cudf/cudf/tests/test_orc.py | 13 ++++++++++++- 3 files changed, 32 insertions(+), 18 deletions(-) diff --git a/cpp/src/io/orc/reader_impl.cu b/cpp/src/io/orc/reader_impl.cu index 9d7e82f0281..a5465090c2c 100644 --- a/cpp/src/io/orc/reader_impl.cu +++ b/cpp/src/io/orc/reader_impl.cu @@ -407,7 +407,8 @@ class aggregate_orc_metadata { CUDF_EXPECTS(row_count >= 0, "Invalid row count"); CUDF_EXPECTS(row_start <= get_num_rows(), "Invalid row start"); - size_type count = 0; + size_type count = 0; + size_type stripe_skip_rows = 0; // Iterate all source files, each source file has corelating metadata for (size_t src_file_idx = 0; src_file_idx < per_file_metadata.size() && count < row_start + row_count; @@ -422,11 +423,15 @@ class aggregate_orc_metadata { if (count > row_start || count == 0) { stripe_infos.push_back( std::make_pair(&per_file_metadata[src_file_idx].ff.stripes[stripe_idx], nullptr)); + } else { + stripe_skip_rows = count; } } selected_stripes_mapping.push_back({static_cast(src_file_idx), stripe_infos}); } + // Need to remove skipped rows from the stripes which are not selected. + row_start -= stripe_skip_rows; } // Read each stripe's stripefooter metadata diff --git a/cpp/src/io/orc/stripe_data.cu b/cpp/src/io/orc/stripe_data.cu index 13e606018ce..d6631ddf394 100644 --- a/cpp/src/io/orc/stripe_data.cu +++ b/cpp/src/io/orc/stripe_data.cu @@ -655,8 +655,12 @@ static const __device__ __constant__ uint8_t ClosestFixedBitsMap[65] = { * @return number of values decoded */ template -static __device__ uint32_t Integer_RLEv2( - orc_bytestream_s* bs, volatile orc_rlev2_state_s* rle, volatile T* vals, uint32_t maxvals, int t) +static __device__ uint32_t Integer_RLEv2(orc_bytestream_s* bs, + volatile orc_rlev2_state_s* rle, + volatile T* vals, + uint32_t maxvals, + int t, + bool has_buffer = false) { uint32_t numvals, numruns; int r, tr; @@ -704,6 +708,10 @@ static __device__ uint32_t Integer_RLEv2( } } if ((numvals != 0) and (numvals + n > maxvals)) break; + // case where there are buffered values and can't consume a whole chunk + // from decoded values, so work on buffered values and then start fresh in next iteration. + if ((numvals == 0) and (n > maxvals) and (has_buffer)) break; + pos += l; if (pos > maxpos) break; ((numvals == 0) and (n > maxvals)) ? numvals = maxvals : numvals += n; @@ -1502,9 +1510,11 @@ __global__ void __launch_bounds__(block_size) numvals = ofs + Integer_RLEv1(bs, &s->u.rlev1, &s->vals.u32[ofs], numvals - ofs, t); } else { if (s->chunk.type_kind == TIMESTAMP) - numvals = ofs + Integer_RLEv2(bs, &s->u.rlev2, &s->vals.u64[ofs], numvals - ofs, t); + numvals = + ofs + Integer_RLEv2(bs, &s->u.rlev2, &s->vals.u64[ofs], numvals - ofs, t, ofs > 0); else - numvals = ofs + Integer_RLEv2(bs, &s->u.rlev2, &s->vals.u32[ofs], numvals - ofs, t); + numvals = + ofs + Integer_RLEv2(bs, &s->u.rlev2, &s->vals.u32[ofs], numvals - ofs, t, ofs > 0); } __syncthreads(); if (numvals <= ofs && t >= ofs && t < s->top.data.max_vals) { s->vals.u32[t] = 0; } @@ -1571,18 +1581,6 @@ __global__ void __launch_bounds__(block_size) } else { numvals = Integer_RLEv2(&s->bs2, &s->u.rlev2, s->vals.u64, numvals, t); } - // If we're using an index, we may have to drop values from the initial run - uint32_t skip = 0; - if (num_rowgroups > 0 and false) { - uint32_t run_pos = s->top.data.index.run_pos[CI_DATA2]; - if (run_pos) { - skip = min(numvals, run_pos); - __syncthreads(); - if (t == 0) { s->top.data.index.run_pos[CI_DATA2] = 0; } - numvals -= skip; - } - } - __syncthreads(); } else if (s->chunk.type_kind == BYTE) { numvals = Byte_RLE(&s->bs, &s->u.rle8, s->vals.u8, numvals, t); diff --git a/python/cudf/cudf/tests/test_orc.py b/python/cudf/cudf/tests/test_orc.py index 894c84eeb3e..1a785d28b48 100644 --- a/python/cudf/cudf/tests/test_orc.py +++ b/python/cudf/cudf/tests/test_orc.py @@ -953,7 +953,7 @@ def generate_list_struct_buff(size=28000): ) @pytest.mark.parametrize("num_rows", [0, 15, 1005, 10561, 28000]) @pytest.mark.parametrize("use_index", [True, False]) -@pytest.mark.parametrize("skip_rows", [0, 101, 1007]) +@pytest.mark.parametrize("skip_rows", [0, 101, 1007, 27000]) def test_lists_struct_nests( columns, num_rows, use_index, skip_rows, ): @@ -1034,3 +1034,14 @@ def test_orc_reader_decimal_invalid_column(datadir, data): # Since the `decimal_cols_as_float` column name # is invalid, this should be a decimal assert_eq(pdf, gdf) + + +# This test case validates the issue raised in #8665, +# please check the issue for more details. +def test_orc_timestamp_read(datadir): + path = datadir / "TestOrcFile.timestamp.issue.orc" + + pdf = pd.read_orc(path) + gdf = cudf.read_orc(path) + + assert_eq(pdf, gdf) From c2efba96cee6475c67f0ea1a650de992b55df025 Mon Sep 17 00:00:00 2001 From: Ramakrishna Prabhu Date: Thu, 8 Jul 2021 22:45:52 -0500 Subject: [PATCH 02/26] test file --- .../data/orc/TestOrcFile.timestamp.issue.orc | Bin 0 -> 53588 bytes 1 file changed, 0 insertions(+), 0 deletions(-) create mode 100644 python/cudf/cudf/tests/data/orc/TestOrcFile.timestamp.issue.orc diff --git a/python/cudf/cudf/tests/data/orc/TestOrcFile.timestamp.issue.orc b/python/cudf/cudf/tests/data/orc/TestOrcFile.timestamp.issue.orc new file mode 100644 index 0000000000000000000000000000000000000000..554948031962d638225c10d94cf6d5f94b351bdc GIT binary patch literal 53588 zcmY&%pfceuiNsEhaMe-T965XLwEV`S`; zexd*N{a#wzH9*5Zi$jiI?LqZq7>y8aVF_rI=$Kj+zF~aE^fX*9hW=F&+sTiZEy~3j zpAr9+yH~EDE*%V?u$(jhfBNHR|5M9y4z$s%59`wtl8}(#Pv)Ph;-9JV>@i47Yy8jR z-`v-Ko>bWX|HNx282_YIWsq;mk8zE0tJSqDW;=TogCi~8L~eFMOz_%b`Q^<5dUO6E ziSh%iZ(vE%6eD#$BF9A}jy)A}K^(Q^KQaO$WcU^^CY=l}qvml^njs;V)Af^LsVQY8e1NHccnCQgk9HACv1yp z5?Jr4LkolQcIWi-cE4!n?Jh6Q)3oP!wY%>j5ACjE~XtKp}A7uo#%n`{a7X;uM;tMryx-!b~LD2V*}o- z?o2gDygZ|C%f*s4Bp&gP{Yv6T z2ioE_zuClVhEhiA3DF~DtoTU4VOA*FV6khOpAAWv71kwLp)WYA7KQfdi_%Uq6- z44h&RKQVCN4B5ZHW%|^bZ3oIWVHE_;i{xZM*ub~$6LmUY(1yv=hB;!j-BL(us;MdE z)HMHCrd)S>_u?n#E*`NzI6e{xJ!I#p)L|SVy<$|7(}uM!s}T|xt;n$#r&7EIHlhwm zR-&i@15m$>E3zmBckouY}{`=SRjk4g1^UL-gW-xI^*4iB9o` z!KTF==nTJjgSIB4=)T`aF&X4q*TNSujZ6K$q6>YXJ{2F5cE*+?cZtz7_ItrWa;97=N*H zp-pY(&FDUfVk^AZ7Bir3s^oAYFuJi>g<7$56w~E@8t(l;&P3PD9Qg_P$DkUO6%xV` zku}|`h!2&2VVpef!gP?c?Ns95PbWf_Hm9l5o99ISdrwB4@4`4NJVer*}Z z%3dd=-!TzcIq8A)8@Y<~BgP~BR_aIm@ijTW%M7|RT%3Jq zX(omdFR6ciw^ku8c6z99O=;=(RNnCszcJ*)XV>PFY1d{Clwi!HH2uA7)@_1vf=)+) zpi;1(y(3}np$@Oz-2yLsWVvpWN%!@D{5WTgHK;nIr|o*Qn*Y{%i}!f4UYV5SeO?Y7 zJAZ?eoc#~c8hCm`wH2yc9?C~NK8rY{phha+&U4J}FgZQ!Ha#7;5XN435MA%Lsw4q& zWQRgKjw3~>m%R!zYAZQz2oyM#{Ha{ZJ8q)0Uc;kRE2e&Wk6^FsUOi7F6-%!y8dJS1 zi|1@3t#euIs=e4^n#CPOc;|1QyTtkvbYeSJ4P$rLqPt?zVE~N*PjLGOXt^MCnsw!- zXzam~NFcaw=m;ciO952$vd+)*Pp<{@*NV*~24FJQ3_DK=`m>EoL}1bT?#D>#M8xPq z#B^s+!WSrVc}Kz28Iy=G37#`7C5JgC0@n+N^5^+>wOZMR5`xjKg15?)rNav$qSkAQ ztYu=#HXP`J3){i3#RBbT7)(N<6A{tl$`5sMC38EPtuZNSL?b z!%zI>(@V+-Su^NV4vKFEbKHRfMR~^yieiXWcnwmN1EQ&I3fa^}G`1rh74yI?80WFJ z+2Uh9KdeQzex}0*zgsh~nR@gxn;xzSwCapMrMxJC&n}HvWw`9S&V#kwv*lk(;boO} zC%N2Z6zbB#nFtcJ=(OjwyJwGSebOG%4wVIx-EI9pNbX>TTOK=LKO8y0LrS16@9abE zl|_dtmhpYf=ic|ahVP4Y4YQKFhP%2uhh6$Qhqdpi4{^^FW_N4w+e1>YQY2eH3u4f6 zA+EVM2f4ObZMs{JBxjQ!X_I z?(q|&y7^0N(Hi?02gC_c-}oF%5mZ(bmY2>F2YJG5h_cI5IKOs*$RE2_lG_78Ag6kz z(+1q-&Y$@KTHYx>+W5KWyCG(VfE}JMT=wj9C|Cw?0g;7lfzrDek`saX^HY~vt>alH zGk@`UUa?7b=5=h-*u%Q(I0yzY33sU!fRmkpzHQ8IDAxnypD-;A z(`u?d6TqN8lVJ*xa#Q#i)pE794r*o?Bm>VW5Qxpr!fRoec=ePxn_KU}zXxY4KLH)j zWys4osK|>oO3G*4D;lIcjT$5w>cQaD199JVNAZdx%Y;7m=~}Ol@XPFDD_`!BsD0jBxM#^jSBDy_z!S^NiEYR zLG>8jISKb7%Hm0}oq{wvK?KU>#W2L}#aI4qkyIG7D3PPr5wMjSPM#DF_pw(TVP*7b z4L&+5XHlZ~WCbgG>niy$^Metiy8)0vrkzIna?fIy`yxS|<*@zGwE)sDue`&uNJQ1F zOSuz;fStAbEW*{*2@ol~ph)XbDcjlMn1+|BY*bT5b?RIg;^JIvggnAJLdM~LTdreH z%Ua~gr7QHENoB&b$8r{WW*emrVJhHQqUTs4@ra>lsg`c7GI*BPZK1ovaUplZaUrF% zLxd7GbEe$LkQQ_wfLB%UVOOo5-{@_09Oy7TrDTk@WyZ2wZhczMK8?b)s_D#y;qN+* zv+fTPigu=5dRS*a=ytu`E@t0?OVa}8<%ZnK*lN+u2=FS<7^1%mBG6@v{b{R|SrsN8 zY~=aw8yL&3eG}}5fLXvi++kzU#*^ORf-HpidB+Kmz=%Qx*gXT5;A(T+>*|qp7Xia4 zu1nj#?yvDSi$CR$kB1cXDJ-APXn^{V9h+VC4jMy$_49z`5Cjl;@38g+Ka$=u=aCvv z8xJQR9#t2V$imZchlfYVwNQt9tZIk77Vh=|V0Y0LxO@4mRzdVm%!G%6=H@D2UhO+M zkn({!i;S!!Ys8Rt=MRKQWs{w7cAk`a&L$PS<_l`-97NV_85mQ~D$bitxX~{h_ZrV| zB5NMW6a=OpV@&v1yG{uhU8jI{)|<_I!TBAwbP+;k!$PO}(E#1DD(pAb5flv)%T1N{ug9uA7Qe5h9|9KD8#C+Fsks1Ox zj|=?Q>?$d37~5HF{FDyvf~*cp>UIwcr388L zT5rI`^RZ}n8YmGR6sJR78_51|Aj2-~;0)t`*(bx|t;YiF0faxA>bo7A5SbkgS*2TS zg*WiA<74jw&bPuFHODu1o2pmOtqgn@e@zRhryYn7(hwEpd_E%Zwa7A(Xr> zFotbznO&^Z<0_YekfrF&7T5OmS?JV4tygF(!fPJT`o2#8M5ku_X|(aKUiS2HAm1c% zHbR_UTJ~kF;Y*;q!`pJRlZChzVf~Y6iYsw)o^UB$H$pk6LRRU*&5ofG!YQ7tNn%4$ zy;*b`E1#%{gZ?ye#q;kTl*Mu0n8h(R-3~k~)v_SEkB}-NL7JXL9jnlqcI)e|W$QP) zc270u4{FEja9l$l6ox9rAS=8{@D#$Z8h!7&URpJM;rD?6-2CknjYVw3 z3}AJ@3e=6oCo2s0@jFS7xiy)9i>cPGVI?A*XHxczr&aOn*h=&a-1H=d0ebFF+~P^B zb~2TamenUdfk57K=L54TBc*vch5QlS?euW0 z?915f4@k+7jzx?E)ut>b-Qm9N%{q9ChOh;XVdqW!5y3K@+9Ie-2@I_Nb+w5&$KGPl zJdND!URbVPtAt*xjeu7Ckel^=rg(@0|FX(GTaGVe9=5ml2Hbe`21bDStM3ojdL1Jp z6lsGovXZ285`<{C?hE)|dB0iYEnWw$4h6B>9hw;K4uuD|v$)Dq<>RT3&+CWAvgFwy zLZCsiGqI&bp~-$q5M+pr+;1sIC>z@0TT{c6=bnel$HnFrWRr>L2iA zDt|YKvvJDjpcXRZ9#5yD8n!Ur8LC_uO^X_l>mle(4w^tx;yL113e$EPUDXiOcD3~FC$re~}+-*S) zHn;t_h=X-pz=piF!q!D;7=O@TJ>To^Fx~%aa{%i6du7c#>)>5Y-3)*2ptYhkMTKJ_ z>Ur={x9AP*&i?Bfi|i3@4qnqmqy3{xNkDl*%B(0707$?Br| zrhQ^@VQcjm)^waf7d@c(xP;Ast5!@jHhaY(lBeY5!~zx{J7FK8zB!PV2l~01R8dUC z!oPM)n_s|C95aY9nYb>ciOtl?hw>V8(sn}R-vSP-D>H_%?erxQ3|_{Q!4ZO#kGl5>dw;!Mlcl@Fah^A7l%-xY6nhDyPr zLN|D$LXClc+G-Bhimoh8w*CU-w*U1%*kaQ$4+fEanXeW;aw<65svJ@35$P#@67y{` zdI!1Mi$#AI10TnA0=2kNHO=+;I(v!F(BE5{d8J=}^ghrbo_Ku5$x*&G(q9pRynly~ zE9i1Cp0=wt_-5=@^lh|BQ^W0|10gvSF;d<3XQX=H0trN4HyzRwqd>RuDS5zz$UJt1 z!cm-p+I|hQ>HgB4#`uXi)47u$#RkeR{m}haOIuovRndCC^i(DOWtoBe$o@V>J6r(6 zxs3jdrT_&g)&O>cPMVzlX31s_QayNcY1)e;2mtrib3)!RPCuzwSgY1y|DPl4-0swR z%YMOni0u{wyg4;x&_Vj9uqfK>$js~}WGiF3s96y-7#QHVzPU#(26+daHd{sJin%&YHs*B3p|Y#}B=4N%*x0Hw+~AQg5RVvw3R|QzECoad#*+ z-M;v(n+#HK_qK0zjxF}A4nyVTZo9gsDtRH((Z*MzF_1KJ8$Ro}E*y&LmT;)!qR4CD(nx6ouf3N3ZHh&tcc%x<= zu<|d{ndg6|r&E%?ti|^xGiww;SpjcIG~<_wAXVH9afIwK~g z88-jioe*QTtP_Bo*^_H}jBxnn9QhychRd%U<_w9W*HMcf&@{2!RlJ&-!D&aYP=kwRCWtPN4kw7KN#3 zLWWpHX*X7<%dRG|tvGlfjFfuRCpaXj*%p)ITlC#Z^v6zz7PP>$46X^=0HO zHqs_~$Z}?n<+b`WY&#<;;#z&CcgCNSPfD%|s}Zg<9#~g}m4O&dyZYc%=r6j=z^3(( zV$K4qi*KXHa80)~n?vi)5PZ7&^@&D;=ZM^RG}3=}?#t<>oi)0H%a*Hy?!qbD^CXpq z=6I4#??uv&RSU8mGqA>f%K=*EG0R21^ohe^P5G0ui=E8ulB_nuoDm;4UoO^GZGTfF zzq^N=H77B5_KjE{AzEz2d|?NouhT=2G<1ns{48L4?LvM4103PN^=c5 z;d0?)Usj^i&$9NPO&9wfu*gw1rSW9cCghT*fh;!5q2x+&~5~Nbu4~*wIhIObNYY&C zblO8GIq)gAIoeI@fHg&zfo8TqDNw7aewW5q0GT_$%_U{N$?(~y&C|xGhf(H1AuqyV zvU1@?7lZ4+_pT;t4(ULgzgAs166OA#{Yz3M(awsA{59G?jFHtlc5A@Q!J=`K?#dGM z8iFY>*(lzc(~bM^COI8{NS6ptQp$}d&wk$!pO5A z4>c<0OS_gb-}`05gw2ZOp+3^2I9zRW&n8bQm^=Qv#o}iiXQoiC#zm65_6c3!jc*sy zy4?%lp~h)HI7RLGA#{}cm(scdrC>b}=t&G*&$CS{2GGTRd|~NWAm*L(Ho?aL+1yZx6PyRjGb*9h!2<16R$%^3rCuN+LqrZ|6%+sDJ5Vn zB?ZiU9O`w;_9T#TP|lHYHZabZq_c)GPE%fqdZqB_##-BE^yww@@$;mh;8;fqo|2D2 z8hc1?Izl#Tfd;5^Yd_CrA;0$Z;VZaVBl*W3X-R>D@W0Kbx1sA(o+U92qYm3=VBu0| zIplj?PQBmtB??W|V9RT+|M5KfIMvJ|n`2Lx)1naHsIqAlk(U*ZM(;EmnY|9>eGo;| z=9B24l@Iw+@xPscGmljro5P@mg)qeF^u1NNfsditVv2AVtcHiW>X%3Ty-DIp{g3+u z{9E@D*>>;OZo3Sk@`B+OC8LE3n){*E2zZ%ei|#S+v*&l@A1|`gh{0~tylLG?6xqz= zOd42=?v2B%Ij^8>LqV4fMjzCl-5=E2h#)j`xt@>lPMo@uA_CBthk*_BrshkeH&p4r zSkfb1y#x`!M+Fs6FWW1}%s4z$0AG$SI&-6GrPZcco!wC;74L}b7(s6k3iPQ*~v~db7ulK3Rl7g3C`|}Bu-IBR~ z6AM#G@yniv_cfwSfa3tzJgfXTQ`X&+NRY0oM1+|9%PuaDI(D(HPqoU^AiAQ{zmDDu z8OQ7=PBr3-p(+84o3Gb<*W=+KFXD^5>ikkHvX2V{dF2d$MhNs2H)sE^KPhdbh%n@6 zynNdQb#dW1qOzgy!V%=7|3BGkunT+4X!(kul-S>C!6W0OtUO0DpK14*A+R)Y?ej51 z%V7m5VYe*RDyUvGj% z8-Z_@O;k?h9f?U27DK6G9yEt}K8MA&%I2wAM)i2=;@b)Ry^@l@Nvn?q{!#Q+dMH!_ zi|e1-X_YHq&LoWl9dC^U#j-rDAG)K#W}XN+-sd`)$qudy(!4aE?u;q%q;A2b(t!PT zrY>QzXE*tXZ^mM)+AfX6?Dl~*JJUV8W!G32V=j^|+A2ftH_10zQgdVexpWbU=egPsmMea#V5B&@^oz>m#XHgOfXM|?i9f97d5 zcZ`L_mJk6H>;3aMhHapz36PZjBdbj$_C$}*%%tF>w6wv)P2T`Brecp4MEm8uzXzrU z#gudMno|d|LJ{vl8=r2mi5T>ooL5!&MdilJ43LycR7uquO7Z8)LFY*NeLV4H*CV|z zmur;4DA1bHqP)$}r>wutCyfzKD9%g+c8;vP&xk=U{ZzV(c1E(~*S4(WD_c*x$LxMM$L!?g$fl47RA`~Nx)kw)a zdM#BVQS7$e?PZ9!o>DhDZNYOB<;ep&gOIAAL)q+5@?3_<>q#DQhV_INgzn}`BghlQ zW~N@WiKg zOfd4Nechi;r0dw|GYiTZA#vU&r20Sdvbjv+g$RfW*g;N`PSQG>l{1q-4*&5K1GW7o zwu)ULrGA4Kp?mUekk}rgu|DDiqq*?;?PsFIYfKs9t|R{3i%!DC-e)5s$XUh&@NOGt zqEMLalk);%)^s639@qE@Vuh#9E5Y^w=5-Z(M<>(#BLYFdHS;-uoquO*Q}*>qL%i7g zOC##MzgPL^ga^KAY@K+_v3Gc$Y%AjtG zr%NjvL9IQtzOyY-elh&L7_g&A-YI@=^fAoA=6z&w8DOwie=(}!hsAqkY$2985RJ0D z4SDA~@vr#HOB4fbuKQr>6V5agK7LWgTem}HE0e@iiY=ZR8P_5k+t&-M54OQ(i+C$w zUHXHm&k2b$I|dza>UJ$|JQ0Unp$zJ9Co2EwZ8qRs5pR8`VI9@KstruLg`gbH?@cu# zg+Ma&ZUB+@5{mZL^%SM7n@vI@k&S<7Z=E0i6uv@kwdwySKh!e#_x+pZn?+9E`l)Th zjv1+%wY2*~Ua>Wt{38aAEao?agPlA;VXnEa%dwye(v|Y5S1PZAUM~*|=<>b27K>2> z!4I1|stQ?HH0bI5g?t4|B8|T1P9iivjd3ZyPh&s#J5+2k`d*9|^M?2fM{Nq}G>%P% zU`xG6%cs}pWylyD{Fld8FZnWqB6dczFS_-?<;{^(RK0hG)CT22^1_c|bFLh8dgZM} z78((U6p?R57UDyWcH2U9&aC~-*n)W>w!Y^7XgUwW4b&f$Rdfjl^U_TH@2XWZNDftz zBXm?p7an`n_Z(V6IH;DIBoUZzes5Z27S-ButZlb)gfAusrLgIYx6W@y!h-h_%@ZgX zJ2Uj=1ta{-j}vzGP0|)F5k6*(bolFkY5JyVYGov#8sUo1Be-FYBKnndcpYrdD{&X3 z&l-2j@k!q z_8~rgRV^yNe5C$gmc3%FoWI2lq=?lvoyMgnY{H!5@3WBeAyZ(@aXo5mcnsi_#Nf!{ z#3)AdEv?IA?KhX&n!JNRa8`St=8BzxlEdQS%pt#F7^$fsq0w1304&y}5~b;pH_~S^ zf^1qIcYW)Y6XU5fu(JP$=`jn;)eAW_To+?1WZf5Pt=Wobv?FC*SQ|kVMRIsOL~nn^ zlC0#G-TkRtwJOIk(`7ky>31n!UoT(m7Qak`*e{yBs|sdsG*>Z}KjR&8@W^XC3PK%- z@?>cYeTB@vrko!VxD^{C1W)l;&#~a4p6j8GUnUaH)vyJZc$NXJV&erk@NnSFkmn) z!icV&Y3iArl)c94+0XOXZf7$SP`A6T( zjNe*}ngMtf1WIUB!ZM7Hj5O=`7YA!O)INCJWN3{B`S&>lnWzj4d|ZfG-dW$ z`8b!ZH|1QtUMBXlR$?yUti0*jnHk8M9Is=n%m=8LKHYkqu+x2U$4#_YQ9h+vVBG%X zarL*QDG!q~Fc_1$$Na)u#7CLajv>dfEKpqqy^M5xxhy>NQ;CzPqa&s*5ej!Li!aka zDV|ne@?Ow@A2<-LOAE8#Q*eiy1E+5iKhrKWs#m$2+lO}P9(J|;GAllLIDE7mHq2hG z9kgXY81ueCfGJn?yoxi{JJ3Y}$<{PF~IUduK+zYLlxDD7U_4{>pa^+83 zUhQiu*^_fcmrEl$!~*F-L!AV9s-isD<>YZ2PpRrBucYLz-LM?1DG_EPx15552lA9g z*O`-0*TfS1fsyt}|H%k@2jpA+1f$ELY zrfDaMUg!O^Mt4RB_7-Y_kmZC)up__8jK?ZwyLI)=htD6>1hT|6N>PN~5>_;;cf2l_ z5u05~Ef*i5RiE(_({_EF9@}XFjp;|ewEtd`&Wkreo62o; zEBe0`T(B@yQW6djTPN3I*Zv}6SWQC&&Wps{EEk>2L_foGClY5$>pH8!8MwGc30sB( z|3x8lzKKq+&re2jnjD;I_hpc@A8&>H9iO5;Ey$ppAv$(m4?UB-Wj$_3y@PUU?J5Zg&ry`TQp-fZLFv~RV9qDLQ=9y#CYYV~LAk9kGAsV83ZZt`ktzSyT zRr!X%j%u)s-9-D2pBrXkv$qKdO9orV(~M^>yaKQ14!_y8k$*VyWvY|kIohE>@a5S5 zmUB5-=A%n`;TcDx9SgL$R_61xg?o$l)50`fyw(d)9zTgS)(n=&v@zMcb*X5|se})6 zJ95&rs_iDYwYHkcRVWtMm+GmGTu=VI>QG{0_Ey%n?S2#ef%ppp<3lAsEsML_Bfvjq zI+8z&R}GE7gx2YfI7`#!1Dco#njw#HYl20-`?m7v&7+a(%S28{ynk@iyU%6s3opBi z${UO)NMFj@+KSPyPB$ZUS!1DMG*!qjv<@R*#6}2!6+5MqX-FCn51NYffX`YWC0X3l6-8IKsO2R^!o~36m z1WxY!xwa(woJT_b)WeXVGW|41M8tnTQj71Zj3@YqIR}~{HM~Y+=QGTHYPK}U?VWA0 zjvtFDazvao>+4%}z!DN2^w=}nJfpX9Z$WR--H2#!20h^koZk@%3rcgK^Bs4i>D*0mDap6bCAwEh-D8(@5;Io7C>Z{%D?6L9@ z*m=o+`Q}}{gQW3~pc9EGdQCqvfr(aUSUp$)w(C6|x-x6gLRrwhtY)TTJuI#iStWAP z-f@Tu^w~YJPHY@#3rsc;vs8Kcw!`viBZz%IV$5dC{f<3!WrUI?6PMJTR5xu5+g$D+;XJkU3FK;?yf@( zb1ZpXV_2@ueNUX{kqY9Tb&&T)>(rgalrIvZcji@x&y2Dvi`KKPsE{QEB;7H!b|WFD zmi=CUMt!SKUEAg#4_)=ebN($!^2MVrrfqG<@J0!6PYX>Bi?`c0Er#3d{Qb#D-3^+P zx1YE2nhg!vFQ9c| z>Elb8i^SBx>q-+A+=xPXX@8=lQ8x$(ie#>A8JvXEwBNEvk^G!mMTz1Y(^CWc#(BA( za|({^c3&13%$0Ni+O8EYv5?B^F#mxF-ozn^$G+BW*g!y{W06N+$6V1!QNyK;jf8JO zqRHYI{3@a^80x`+$*WD_UISkP5Z#xzss9|kJ*rm z@5Agi<4=6V1YKsFZ!L4U4z^O3G`Iagyw*hhp(_?-oKnzG?W6r9bEzA7A)2A=#scJ!WZDbDk~~toZ`&Ru5If76-ZqhTeop_17uFfdyuQbdDLI*G6(hcKbYi( zmWwq^lb9r#J+QTBa2AA1+n?Iw{${_UXV%fvN@CZb(IV+%3LQ)$mEkV;zN$Ac%R8ma zLPEQ@9vN_{NVaSr9)#=4EP`B$5|s7=Z^uZZaXPYk-GH>NcWv!BRAt90%ob1+o$&0| zOQ}-puuGtQl-HjgPr4_6pQBF(eplhXfS*H)js_OKa~(IyTK=Hqi-_W~K`5;1@Q}4kT1reThvAG5(e>deM3b$~ z`V0hy4`RHXZ{xt!Zk31};BMmH6@cbJWqoVIFO;2aUO9}z-=y7G`xo&!Mx3QUC|fdN zhv^nw-QJ@B&UU^Qh)_f*&!~=F<(IA}BW8a)eE9RbK7kw>wD2n9&(h6%dhfkxDi`$( z>{Z?^3H9*JJKfVy1t`R$D$ZVYOpyUi&Cn6|~dkWUPd`$8Q=DS1-Py4Yp0>5_fr4`^IA%Yq~ki@mU%B@)R|pxzw@MN+jTB z&PDK@C%R!KmWF`~JRCja2630;LJv#!OWE__1>gb$W3k0%1Pf;op_uNGaj{=}t<`hE zi#Za*0~LvT1_`VQb6S`$y1l%TnLnr-i-s5~AG9(s@0j0|L$4GEbqo zArk}H9*6s*W6}53Ia~b$F3pjYXOBQ;4^6J$nmYqADuv|i(+~||(&dW5{LQZy3u;h_ z%p*a4**mYd4osZKn-&qSN>)`LD!{$!p@6Hb{Fcnq0nM4kwgn4cYJU+VQ)HEC{^w+3 zAdv}daBCAv-z@!6egD;RVAoOb(MIru8Rk4|btAuajcC8vd+q944xeUHEnD!A!5RZ^ z^ljiWyRGjfLD|t@=)a4*gW>3b4@Rf$k@8SoTIu0_MoFcn3OP%`FQ?JnvAB!>7MAuUHf>*r)?D@@)6`zZQ`$qFrIxkqg zD2gmn?RjY???Qj%2?;ox;ka#)DBKu%YIi|>;DIYs{+H%alsT{6v3D$%p#;~ys}yX_P$lBZ)tfh25wLTFIvVyP8RNE!8YK}xjqRU zwMZh98c!>RwKPzGJ(@G(*w8hkZZq3jy#dEMMalbED!sAwJ0{%TMrijLk!#;Ua-zg$ z^WWUftQPy&Nro>k<#o*;K|?`;O+Kw$Vs3@`0dw(()?cq$%RRU?skxC%d*60$%M!G? zS;i-}O!(aMvkDc`N1ktl-mK1j84wni%mC!vtnNZj6*B*cSXMFtx%6KfCztrI^fH%y ztJyNU7g%O{N2D21v#Oc)+L)i|&{|mZ2P%@o`=)@J_AR*!pJUD^Wnm~&`<1Vhcm}M6 z(mKP~Q#6gNDVaFPbH|qj6xhn^`Q4W%PsfQ+Nny?}x60RrGvvH@q~u?G znv!i+B%29m^F3%ewfWHXdr{)aCRr+FU@V&7|7j#{{jJ8J(XdE= z>b_MrM^r0HdWbcgJ*>TL;5X(4rl|hRK5J|EO^$Q;S%QcXPiBZZ;IX-UkB{qMG5b5! z(l;y?p(WLFHt?m->bVqli3ue!j92h={+*xr7RA~4!>{YAug#9cQq4lcGom2tEFj^ z9(EQ2yl$(??Y$~rd>PB=(cN`(C4`Q}z4)VV#CMyg^`pZOe;|{C#GtlH{>1|C#cQh& zPUOTvyZ>^jF)gCeZNqB&X0k`Z$D1Ee?x=FF!Ep}nC(*uGBlSpB>5tH-g`F)EZXa55 z#R}7LCEM1PnJ=3q9u#a^KMI-mlW~(l{p04@`L1^N(X_o6Qg6kQD?OhN&L}Um!zMZW zx7y;WLT%P65u+t2ilUt{3@C zHlg@s?dT6_RR4>NN9 zOa$sk$`j#*>{@26W2bo4_U*g4UBjxIQ@5R5S%6KC1hQ{rsr%~V8edWOGy!8Wz!fd$2VTh$1K;!Yzxbk<=s1N+1 zF4xR#8y<-o+6Ho`oF@kGl&Sv9kYD~;!_yeCea-?N4}=jaM1s9nn10zLOsDm0S1^2_ z&HQimIiAkAec=*l^t^=Yq0Ao-8UQwZ&vd6zt(p&itmjVJ8#5& zMgQ`7tolc90YZ=y?@qtE{c9q5z16{7$ZCW>@y!xj8cX}Nf$7DJUjY`ZOIsG=CW9wt zVtSpfnq|WQ@$gfTHo6VzZQ>t8!}k*>)L2`(fm@pf$|rMYiJ=N0bire$irRRvJV1;=_cAx#ID=%E9Q+dTaB|4A!iF&7U0G)9HHxf+wbZ zG;7K7%ie;tzei*5WPeDl)-QI_k%$Fb+#KgL|8j5PXU@A8K||3 z{3WN>FsG61I84Lm-^-1L3s;2fUZx$JeS%aL(Q)DJ%qWQC`u&&hn9LQpq@HA71qXOA zE+~!%Ofp|NzVfD9q!x@eGcKvm`qCMvaVeT#4TAno1JIyE-B zs@-U#OprH|rkg36Msq{KOXNsn=7B(3M^#a1e>w6vK9r-rf9LLP(XxvRWu0SwbZk~} zUVd&bAAXG?Ai0HTFkzI;7OelWqs9GOgdZjqoPKN-CQaY7@2&9?aobSjeGO_{E1nyP zwd$vBik$cgx4h}ZTnjc8So$PN^m*YJ>X2dg{emW<$4ggspk3nLZY$2}pNyn*5TDe4 zuajMCnzwAL>?{ac@I^C$z09RM8cMGnbP4osPgd>*(AVY}X=>a&b>zs^6K?i zL&2^G6F*(%>i!gt)De5+mJOS0?G%>&#s1I+SzoMCokX)>cCE~<3RnffMi9U6rxtgN z)gxMA(I>dzy`YsbU_DQY#vgb0o#d~)cE|WO{7Y%G0ZNd`c-->SGWxseJ%{1WmGyRV zIpDY2*oB(6`(fY+cl`gXH2B?|5se1&T$9XlYR$j*L7bK9y$zyy0{-valmvbkG-pKOmdc&ol4YrkLYKIf{$ zCYyn)0KfA({!}RpsM5*ZpkrAk z*y38-&fJ`brNJRsfMrZ+K(N0eb@~+$gv0W`Sz*5y58(`n`5>bM4W@ajRE1KJ0 z!7#D40MV5Bgynk9T;j1}d9*#v3MG`Y7&x0Gt*A-+D%M(beJlA%!JoflbcS9ob&`d7rcmaie zk}@bOH2!Q{sX-{Hd_m&iI}4k6Yze*oZdW@z2$*Z&O%5vnPE22_RM!N%&>|G;HWSdDSX36X9`E>jCZK4 z0QoKM0%gZbo5y5w`g_bvGErp`SAnSR$Nu)7iQT!G3tg}C3;)POSWn9_KDlQ+UVebO z+}SU&V}B1Uy0H>aBF4_S`R&EEp?k<@GXl*r)kbGrO>h`)E?UXP1ofGm_f{%SVEdcc zw#enis2M*eohc?%;ko3qg{LKgfxA?nUP@jWxujpd^1>Th>-WBh+drqsDwY1zzpLX{ z1r2SNzY}6i>p>E4re1%1EG(&$eMgR(@cEE5yI7g7_+Vw2f$PZnzSte)iP@^Z7 zOdS6QHbBY0Jiw?WT(R)9q(!jSSDZ~rC>4mAdzN>V>yP%9*Yo!Dd*{0gT_)Ycn8<}A zvL-+^JUpIPS!n~9oy@}4&%3zBcOID%LY>b|ZU71>ae8*miND8#@5+f+_)GM-C}p(T zlX+!@q;@>ikvLGysS5e;y#KQB|6?TT@BXj)Qkt@>Efi|>VyNY`tJ;%*Mt7W^;J=!k zz5*SB+;;RYc5RWZfF4}4mi~bq*7=7T$=1HY@$>I6K)hGwU>;VhySX@aTxOT51uKN` zp-UyithWCF&Wki!Txw{e@tI4P;d}xx8KnWUBZ4P5F_Y|c-U~x;1tDL?@TGs`I#xai zR5u{%qZeJ;d41#nINMuzj{WFzQ-Q%aMMWs6&mGVVgWo1Afb1~v8`tU5q#NBsStBIf zYBwo#%omC)WY@4m4MbaKK*012o-k|X@|4D8pDk%4<8u<^MuU~HD^k0wv*e}0Z+1PL zcMzq_SSbnmCf?r)Fi>~lL(XB4m)n~0?Wf53j-2P96dY@ANF*fzl$BWNEY64Jr%2mJ_IiTH2{edJBCL31CW3@-yFT>+E>laNk@3fI(6PfJaJ~ z?-UeLc;{Owyey)@P(R@1wp)^oq&s8(NGoAA8+JCgMP zBSOBM+;8;g_SLKB_8g>ZQEr2C@upO2xo{lZZN^nJu5tT%e&co0tMODZV2lmx;K4|R zOPUAQx&R)4HwkyGdtjVv5&8PFJ^s zqLVe0GZC&H=}O4xJ=L(B*#3Sf^E6hi-P(E}nSBvzZyXp!Fds|5WgQy17t%x6E5f~1 z=LZXI4F;MfaAfjv)U{X#0j|N%zIox?8A$8=1 z2Fwc`wcWuy7si^_N3YAg7c=HgrE4=+l4Hig!~B=10{_m+qJ9X~1CBSZhU6QwLW--R zAEyp>b`r`8xNzx;tq@a!Lj8Lsd$0oF`cqH+Um6M+0y`N!5xIcE5}YyvudVDRJ3q9t zdgib@gT#|KDsGG!ORYNmZ&(?{E^Te?OgfuDsw7VDvGRqw4i(IU1Gy(udWF6GoGheV zB-}gZyx_Zdg6ie0-G=E9XNGpvs36F$m<7@%t?4XkaKD4r1_ zIEe*hcAf82%ne2a8uVI}47ijh!%jEhKo>rQQ&MB7r|GfF*rp&*y(r>I7 z6MDrdDy}d!+sz;9ECc=~C*P7Q&u%1NwK+Oej#_!B*cj!04$|u~6U-8EA9Xnm6qKAa zP$jN7=Z;lL>L)U}{j*ttnH|gut)lh}pf&s_mVNNj3yZP!v&czFNV(Dl@mx_oHe8sYLA~Ew5cd`aWj!@!Z>8+lW;Qs;r8mY8`?cZt{2WyA@Cj#Bz}TUYD}1$UjCcZzQO%=&bdM( z6N3a7rB0b;f`=yyP9AGx0b;Ivx)mO3yg9fh(W9_>P(jm|vT|}t*5!``xd-^0 zo%N`66>x|0k}tdRFFp0S|AwMCey?`=8+f&%l#`ag+$PvAt4+0DxYB04gY_tz9aNuu zb6108Po%Xn!imm8ZdX@-{9g3c<91W8C$M{}5M23a$HB91&BXzFl@G)iAX)#~aoVQY zB>TE2*(>)*V1{$AWc>iAmyJKg%qu8)4PwBouwxEM_u{-%t(lztfT|329(aX!S#${E z>YWP+I|n<1BNv=eD75>1)XZcgiUoC#FC`kv(#AZ!#%0JVv3E+Ba}>~sYiF{v8U_om zg)&88{75LX0~>FXsyd%jap*CbS@}}=2GZILUMOdys&=1gfm&DW2>697DBO*58ndRh zl)dNEyrc%_UZVX#*-I}9Vo!I+&IIG*lgP;9q=HgaW#oLh`;)OgbOC&<#H7k@Kltm4 zV?IbQLc=fbl)R?Tx~`d;(G!1?>aV|RTnSeo*rnMnvPw}ciXnAUlr&$SYjO>{M@K12 zw}c*z#4ULPoV}qcmfjrp!9P>y(~xK5^$netGD6H6lAWSGqU$r|>2X5dcglg0H{b8c zgIR2DyWRPB612^zDCOus=1$%u-)G#|++XgqKF2<1&j!wRd7!MxAiG5Pv;8En+&ue% zyRy@ou^`5q`(RzdueYb+ua~svi4Fs2ud`pppt{3U9&2UrYQOh9P!;NXA?DfHob1sk zAHmzWaNEm+<8i5K8Vq#ISde@4-#%eNNx+@2ZyfcXMn(M>>A*EhQ3jeNz==U(3rN10 z0*`R|0@Ns4(TlVHV#LQ6`PNd``GP}|-q&|bIsoed=pltdae9!m{6SRoHMO_*$h_t1 z_vOsv@%ngiUd6We`~;0G>Fddx2ZglGl$1|a2=y+kc2>-=GoGugHFiqn0$4vU?Ipfv zb`UcN^j%6Lr3W*DH>Xx{ooz_+4++5}qKrR9gO#>#@m-u55qQFaCMj9bhc@IX<6&OBe$Bgi#;X z6Q{kvJz8M}(x;u>?>V~=^iQw$n^MD;pZ4su=9wLPFP(@%R+sLNM;emy&=-nhsr;PY^D`aHN06`7NksS@DUrN3}CK3)!zI&or(MS&?Wx{2b9`+!0e zg{&|X96DeltZ<|`xniZ)Q@)}94xBny+~yN!U==vyv92t+s!jx=vT!E8G~8CVibu7} zPF`%X07+u8`jBK#5QCu=9r~aKjy29>Q z_rCi5jBDsUE?xH@ssfy)@$0bf1FJiZIVVU<|G=&Trbl9~ag9?)W@Z2%gBSvz_9XAK zRc2r-WGpJj!I@`r7SaRPr+v-XihCsV?aCt~`O0V6Nr~)AmI_ArBWS!+eYBg>i~BqW z?UZ(M!Q10q7ysNp)o(fDkGw;Ym&p^P^O$1D%ppu7c$n30V)L9SLT>!;a)2JPL*_?S z9eNceM*{v^X{`ReI$uxDO~AQKHh{*x*GE-xxK3FU=|b$?5VCEDvQTQL@`%wgjwP0d z3elYfTKCrK$&98Fo1IVBC#5Bczc4>klXOS=yfgTfkw^996k_mG&V>_WVTT?kq){cI zdm!Il*rah*?yYb&c0jAlIGu;8RV&kAGtiJY0Iy#uv5IF>y~(lg+u<6KIfS!2z<)VATg z(RSMc+0l!>00vT@8@>2*R`*$3g5p*c}UUu?C`;)tW%+lCd8p`Z@#RxkJrd+wUSwfecQouu15`%cCS}>s>hEjEW4;id5Uq za^XZV>0pe9qI5=6rFGJ*+rx7IqVb~YQ}t*zi_|5tgYG6&v4MraM*BNZ?Q*h@6*-{1TOOU?%cD8C-apoHc5SnGt~unO>-mNtH1mMNog`#f2jPrb;>CCxMXcdH{}sN+b!QJRZSc+kj0nSr{?d=d^tu>tmyGsDq3qv2IAMy`SrL^v^&j4OIcp@3#7 zlsV-yrqH^hDNvSG{qA?1Bt*YL-$vG&kyX$ytx!&rfVn_h&MF7**xDex!ZDg(eZjlB zky-j=+(mL4S@kDXoBiwfmW56_6@;?`l>*`uZ1nw|AWKatr;E=TVeC!;w4Y6;bG}{VkxjFjJh|%L!RjJ!c-B zcYV3j>^oNWgHPW;*;0!&f&8Vb3j&R4R+Qf}Zk86JiS$N8#W7qX= zPoBT7 zR{bw*BDsg?lu?g*S4TaCA{fUi8Wrsv#*yULFs7uw*stgil+oeNfNKHqqVG(dolJ-7 z95x}{W?$&rh)RJf&3HXuUzCIZfw6 zC4!^*mg?Dz7mjM>j5wb{egOYoU}~D+SsLS?#b+z=G@7sp3At zmkJ)-j_z0E3Fg*1F=f@5QP5s$PX5FT-68mlM!E;82Dx6X^`}C7wPpPl_vb%VFnFrg zB!4ZHCk{>4A6~UD(v6hQJ|!#ll@e|2jx%Ub?n`Ij>JUXg6`wR@uz_A*@Y#?HNyV?U zb65zv$)4gBssOxGD1p+7Eb$w5H|Q(;DlM!}brp(APhb}v>IM3GaZc2z)IoShhoAT; zr=K3aM)o?gx+p7(GqLR01HHMV%N(vzN%X#Qi=mC=V~`R&l1>dmIRqz?#J!MH$oW(H zeCH=8+6`jGB7GjtI*n}W0F=A$)LasD^=kbsYD2ku$m_S`hLcN3zCnS+rKCxp+pWIA zX*lQ|l!{6Jhsw{I2NYfSNm1u%CA8OFeE>P$c6E^C$^AnyBd*Dr_{?5*oN06-rB)oD z_6z5QqKF)S^-;ZxE}pq2@jiBXX?J$kJr$T%c%OQsip*{+)RJNaTALk}PbGyxP!HgW z2p16u59u!aZQ_HytK+L?B(aO>)u*JW#vOky8>=qPO!`-A^0gYd@=!nHa|k)zVTm7fOaISS@>b&cqY=p&8aw@#HGyNVGlCB=rx$9^T zxWsXD-<-~)5$hykycHxRDUS@#u-OxmiKu8)#lU6TbxP(qjve)!&N$Nu_o}g<{JflO z#OK6I+HopBH1)hsnNnqI72#?#xfgO$#!Rngtu46(I+24@I-Fpmgj938=wRdpci*_r zmsuNtVh~=IibsA2GVG{22UeUaSH{vM3nxvntDzDTuZm)e?`32j z{+!hUrx#V?OgcG+K1&*Vk5wvR#w(RE6|8*e9qTTUe5w%lOg*OkI^t>J$yqI}KcIpk zwMT4~-J$BS$@IkaXyqy64Am)f=LhBPoxh`5Dd%OP=_Wl%J&0AQtTf?|#PW;M&P2pd!8cYSJ@X5mY@S`{x3`UEb%<+k2 z>MG2;qc$TWj~oDT@OC4s$~4c5UuwKlQLC_(q%Z+u)|`sOE(`ldl}ZIIZMWw&Xff-a z=C|2HvQp3`jrY@NcS1)(1&0m+q!?wWPYmL3=050eHf|;78ElT{3~U$7mrmD(?J@WM z&AfS?RbR=+?nua(6eeirq+Xr#!9}Gi^)uL?(`jI~$a1*lxr7T2oC0kU&&b@eSQgTl zo&LuDFwQGsZ3meT%Be$PJuvy3ncLTTz_}8fO686X4Zgf8yc!wDp{TF%8KF2Ol_btH zuxE_5QQaMLlZwpi9c%g2VvCuRoKYe-H5IBIugW&`eP@5rUSVuNRjxZTG8c8L>d{Kxw`|-M3k1xXE5OL5^m%9gx*Z=t>9)}uG{b5XZ(KR zzf}y4*#O~-&OK#((b=JIkdx6#!C*hifw3<~`XA<4<&}d1Erxymr1IhQ=lmI5t4hl< zTZX4jNpRMWQJtXoVEl8+HrRxiJO4b^YgL=Z8Ksk_oIl;v@`t zvf3%7fl(-h>a1EcmLG(O&USMuNHir>Blf1jKIka_Xu(&9R4mL3(2D`~T!z|`1U zZBWvoKlbWNIuR#W(l6tBg7uJd(@GT-UO2p1W##xXzQNgM@Sfwn{b-r>c`L{mQ4_zD z*_PDEyRLR&qvWtUBvL#-Cie;8J92`#t|cjylZS;&+wCOghT0u<7(IsZ1UWa|%YlXq zHT3CCB+imDLeh&o&0bB|naEQVM`h$wAA1!=gQ0pD=Oav#FM49JKJ$ z&BdC1kT2mh*45HikYWlS51NYhUN!Of^l3Zt#ETiF#!C;Hh)W9RnXoU=XtEML7H(KL z>;?o0C*xd*u~T=1fp|Q^xyys=lY4GFBPzc=7WCiTD|2+hAj}zSoH$VzeQsl{*(2VQ zJ>Z>`ymxV-M(IK8#jE>VL~*L*N;)%x{#a)d;^?$$5KlF3P}HWXBAjr8GEj_zm5nef zR6{7%wBn*zvmc%+lzv$kE=ao<&(po5x)^_2h$;)vS0`AZDC0hBP8so)Z^Jw2el103-<30r~)=KY$T;W01`cp;0L4rqUrF3sA2E-T*S)Ix^^ooT# zO}#ul7CpLMbE{nBd*>ZYY1*7((;T&W^0}iM9%v@QyD5$;yeeAdho(F_>q9sbecp?* zG7PR&nj4=1$L}Qll9QWIQ;}#*HOAQylS>aFH?2>v)|n&5rSN*N8?d&ZqLNO)D?4fD zxXS2(sUp}-rE%~M1op}9Pwy?p&UmKWw^aU4o5j7xJW-$MQ)S2gV2zIdpPat_9%5g8 zzd3`k_CSyZOLl}a=s}8w4x>|X_Jy!6uor8RP+D`(NO{t(@0=E~Gw{+rV?tGr&xs@* zMhG~Skkg=WrjZE6nGk494WAdg=tN zmDm6C880Z-!FzRnMDZxfUd*$$?sMB<&2I2-^#rLImW(JeI6ET?Ry_}PYq7lE9RFPxGvNO|5`Y%THBPEH?JEy)wMa9a@$9T_` zrR{8Gw2(hfpL&k`IT_JJ-GHw8ALW04DXt11EISu{!j?-V_;;UVLUkdH0NHg$oY0n{ z1{tUHF&E1PGAr+@ecf}Pl$_9_KE!$H>DwOpQm58_qzXf1c(1z2btcNZ}GM z$H@}TY9vX9#0zRk)sObL3f*a!*azvJdT)geTsh83&hguOHP^pXJMJOi z#HoPDxB@X!U2nb@H`MXJbF#5eM-+tCXp=NVVh%WjvK~m9_L+-ZRF(45Q&&^_`Walkv3n$|sDuZY-vy(3 z_neH!XoRd{tsNGo%vp9^U{1KzV3ELJq|U4V zwMdZbZVsF|fy#@tU~=%qV37voR9~&G&F9{CP8r%5_we8lXqlxa^7!NVsjs!)UCCXZ zlR{;|IZ$f#eXeuq#oP?{u4aTCjSgCWE2wlH0~sT%v}O>L2xiYdV-Qy>S|hn9-5>Hh zqUhz6AoM?^5Nj?tQ*E5S9EQy5QdJ9P4p`}>IP{p46BR$Zxo=gx&Q_$i!mXxzKqW}k zVCD<$JBYpRqdoq(aB+6gv5-3;?FPxpR5A4H-fd_3Cuc>ld#lpvdhpAvXv32MTcKA@ zwsNkCE^_j5joIrn&BVDw<7f2t`ulsmAqx*p&iE8)sk)saCFB@yj!*WXLqnCRwUMgG z_eSdkPBx#jvbS8h+dP=e8nNr*ACexxIRnfIF;%Xt$VZ%1o(dADIDv`00O=)MJ<{>e ziz=yy%s`)H&b$YYAsI=G1W8i=X~VzU}$iR_EfbW^>Ub`cm@&*QrQWv*4 zAM0z_p-DKhUjmgV&yWjTQ$Bj zTC;zIxlWFs^b=OTlH*D@OGS+)Lw*1&Ii-T6S|b+|mch6!n4OZb{dKC^j8~F<3hTki zM_Nx&B3m_#MoMl|tE077(nY(o5oY6x<;V$KiMmULeWf}frHoVLQOXLlY0kG=8vO8- z)Z9jAs)&*0G;nv9+HaHyBsB@Oq7{80JR{%Ah=NupJl9s^h*v>v=QJN7qr>xVp8H#J zlu2HlM5nOo^XQM=JD@lKam6FAQ$)PdqRbN_>T~UM&X-3MDn0sr`YK*g22s!HYn;TX z{>taAa~csnDXR^;u3R!61w3&a9_ah&hsA5)UG*6Pj6nL=6QsMu=gaS}tSpke0y9u==sa|d>VbXV&FU1CV^Xw{3ot(S*_I$EDtzq?N|eTl#D{Jo4_lq-QX*!I@y<2+b9quT+gKw$G^~ z!eQJ;3N5mSRml)w3_cBoSDb4?r0wQIs}C!Hs!jZq?-|K3^kV9@P!l;lRj<9AQLohq zJ%9cAU-IgqSG(wR-hul|rCp$2Ll08f_f2cr&BxNuQ9*Im7H2V>b4)n1o>%JDWA4m3 zgX9yC7{u?!t2i$-#a04cx?Vo1Qzv$~N|Z*}?le|^)SjyHoNYjTOFzZ;Ac071hqVYe zvv8}SPXNOaO79eDcwMQ?K%%LijqFJa;F=wMyQ)KyjrhOZH&ze&Li<#-a*!)ve>Ylg zF)BPWzgG|JUPgV1`xV9)xd$h$<{rk*TJSCTYK$?FvmlRu_e45b6z3YMQ|XNDS&`=0 zRkErF=T*=*p{SEXihMvZnTq=mDN96O(jL86=b^5 zjPOjYK*Qrq!mjWbQdjJkM-8Yx%BOfsd+qbiSpml-58Jf}W`%Q3r$P4~_m9=G)|Bu} z8Dr{o%$ZAZ~gWwZ`3g>n_+2J?BI_(gR8 zoPw>-(cge#XIaGE9+FT!DQK`CuG+e2~`+Q z(*{FCe}X<jXv{A%0PA`JHNnRU<1d` z4c!Cq5Tn;(?#S6NH_h)NrKHo;xF+by*{S+jW_vid15cK)EN3FVylfux%it5# zAzaBnDC?fen4~qnp)mQpw;e%%vCCN-KX)_~a6-cjTTnj%D`T^3) zs6d3Hc=t*-xu-VA4yr9jF$zMu4HTSCWTPWhXX^ErGc2?gnQf9vWvycmsLXBY^Bab5xy zD=Zg@%yP8abkt%L?cgb7ss zhgbm3aQpCae&7VtDH7)A@Y$;W_L)t1_uG!isE}&YfobPAZ-eHyp!5azHbOA`UdlfA z)TrcG{WTN(UJ5qw1CJJ|IXG9(@1?G@BM3wdCkrY+Y4X7kjNZsA>C+l{0-)tmqqsJQ z=Y#Kf`(~qu=uGYG0L>G^WYvatChrQU+E4%q#cJ-kX4W)=n4@i0kBcVNs1&O)&Oy$& z!wDzV5!eDVg-m6fan3=|jAXU5V+ECl`dKugc-uWP;lR*DZ@y4Y9l{AR>1z4JU7a2S zuiZK0-FX!zWfCq4Rhg7sMjS_65!PBrmEtr^svuT>&M^Ra=Co5Gu||sU3()H+nb)Y& zvOT%qBrT{*4i=I~ry7;}YetORQ1%()j)ET*H$)HP)%vaxUcGrmqhUT*x?-QP#owe8 zIlQA$$?1$NFd+C>uS!z8IQNQs$|*srR!Ht~<;BXR>dPmmqKj2NKP$Lb8?7|fhtp4^ zm|A{L>w%py!S~YlEI01U+%XIi8ePz97zF))rvUU`r&2gMVTDXdcy1~^dRMK8 zK2e(fn%u~HrIDuRmZHksu9zt^B3dQTM7t8h#l#tt;tIVgXty`31p+Mngm(h197cw< zV>opR9Yk@tDInO@_TSsyMCCQ~yA2fSqf44_0R%0`n4fKnhc9wDB)% zwPl5M<~a=6JuR=c?n_wRNRS$rE(($gi|4ccl=6czlwH7cTa_7ElitnbYO^Bn|C+k@ z7=81yzT+{{R7I;GP|#@IN~=v3DxPB91M70w<*@AjSg1gtgoBDyKqV}#LN$O^}c5&{o~H_%)tI;=6=5S_j_HR z(>1g<^&1xKY{p?;4EhK^2lKM(Bd&R=C}p0RqAP=7j7k;UE`N9FKdY8fuVLEsM{*1( z`lf=jo@LgR__IC>W|Jb$&Nt?SQ_S4>@T1|I@T)x7a@hICx$fj^OiN#av*a_qs^6bZ z{uP^Cv&+6&@{H7(hb>=W&kVnhoQ@P-t5W-(DrYI>JU~tIftg~XbP4$3F;eZjKenEC z7-H+ry_<(Aylm)k{Js5VmAPuC$+_#d+mFVJr*xCkDY0uRP?cYw`E*NpCE&Z``%^AY zqc17jYbU#k;A{~{0&dO@hfhBViJ#hTKB)#-Sco*3@f};@*1X<7H0aF`~K=> zhc-tp&GaemO^#amkCgJNRiF7ZMJ3D0WE?(FwLaZK`1BI#%7EF_qL|OkrERv&Qm^D4 z-0*eJJSoFwS_rcOoU53XSVZU6aF0^Nqy^JO*{7(-DL&68kTu&@9>JK>6#4e&v6OY$ z!#S(WR;6A04a*UnXRjOw&imc5~-SwZIHnDX+j~JGzKJ)PM>Mw;dmD)Vvo#s?hPWGwS?(?&ornKGjzDs4%_tlG1 zHd3{W9csb~<;>DZsNd2t&^coqp82(DpJPZ&@RfB9Rld!OkN0CwN=Djbqmq$cw0U!! zhddoaw)O6*%W#ZiV$4wpHQmgHQk6~c89$q5m_8qtV@H=5-TmT7cL#DVnjdblUwI8~GE0=eFLG&%}Mw(qTXE9qxh zXi(eqFz6O_-fK&0M%tmYXrze6(reyQZPfX$$kCZ^t*%Y3p-n3C{feUi0QqdYQj2)s)+I zcv)4|dcZ0LZn-6`cvf@-Rp!$Im`_UYjulZMRGrQfgAr<8$C81h@-zc;X)kFtFj&if zP;FUXxv0j~V#}K9vy-%(MDAFuj$NL@@l`Kr*Kda46l=|an}gL3)rB`8maRND7gs?x;2ek|EX|H(=C z;?i@jdhoYYt?@D4Qv$WL=G*@E;bDhA4)fEmwgw-zc@};LKX;X@XQ9l0NjezVW!qup z%3rX*jin0_UTRSt1!jtlYODn8(; zC#C8tXBz*sPn%xO8LT_G*Lz=i!VQg$D{B&}zACK!O6m<3`F(SnPC=al5avsDiRDyb znMSt?pu{ApTRXw(qn*4{RmMEQE-*R2`F$t%e(7b$`*lcRbK3`H63{)+sW>NSDenC} zZ#mMSjiXWy$XMRL(}ROY6HS7InJgVNMsnbXnK)!N5~OV)+L8$IjmPRj}4 zOVPE-4_N0Wt2gi4_9*y!-T9rCCAhrW*4n=2@XUR?`}4FPg)I0gm(G+tW%KKsIBL|~ zQ!|A0?~b(F4T(xEv7IiCcDrqfn)`3+@3wQ4m!{UhEXaeG!m>0L`c_I72oh&mbC<`h zq#DyZ%<;gF*%L0=Ro{Uh`TS=ePtAITyxClNre|=oO<~hChVjlJ!`m(LY(pgVrb7(v z0L_BOeT!+hXV%HJ=_4VvcyYgjX%Bl+n^Zx?Go7=nGQLhdEzBdYCBG>5 z;9g@4B-89EwA;umS6iC;m}lRP#K%iV{q)tLMoO4fn*D-(W>7NUyNnNVBoF7AAXd@Y&) zZSVnUivoMWrR>kq>F8>dBjRbBok+CAhNGTqqXFDD%DqNG`dfEJ`wsAT2=;hw$PZI( zj~PLl?=5_}hyufeXCtv}ifcClkjQfKfq+uu5m4%kQG+P5=Rg12LkLL=-ukFO7W}Me71hSJJ{Ttm(W5AENrC;=HyvisFTG8zRIvkl zw340NalhNF{mFztc4Gg|V*v-{6;MN{DMS`@My#oIpS-5fn-xk~lLWrZpqtzRX{9Y? zx@E=oe&kiM@m+%VANH-$sR;b?Zvi)r<7=lx`Ik(P`JCAtJ@jAjs0flHO z=xJmHn|)Bz4yP-=d8W)6Fhz3&@*~HX$h=3eD1`PsgOa-x$BX|^2qXYyJsD9_G$nZz z*?{66i3T2-c(>xmxB@aHHjl-f2B+QQ*N>q>#=wXkAkZFRl00CxA0kIZZqXDA)J~dZ zWbC9$ZB2})YtGt!X7mhl7{Fa%^coycL9Y{Y_B}zuE(FF*wh6#T4uC!y6sOP-b?N=A zBcR5B_n)58S$rfo!21m`YDkd4wm}vNs1+2ckn70K2Ey%qX3yPQMGHGi1^(`n`(G=MkOPkk0GbUVQd_*l{=SNp zSsaAS7md5HtLXSyZz9ZB{EK;TLLs$jYQ{gNs4$ z1>m6r2u5cKBi=6V0yHtJN<)&@_+) z(ksWKV_LyZ8nq&~0Z-tXt>1_>c89b58fg^tzTdefP6H=8V$iFhbnJC~9wJJ>F5vi$ z&q$WY-~cvdaQN8w9mCXi-b6iYk12j+=R+aoU-j(M1|%@Fk}^%)8W8gFw;b_9J12?? z1^Xk)SzOMV$@W~)zY7(}p7DQQe-wH~(`cLqCwH8)=dBbA0wq})fnBioWH3IzhME&2 zJ1@Y@f}XwZOuGD2*PUGc;jcJM5c__{bdf~A{zb?5U!m9&^D4?UChpJ8w+2AkJrEXF z@7=#QO4LwQD`$!Gl3M zC4sVE_OuGEi^w^<#)?oPqzcNNj@SRSUp{spH!3iJ&S@!xV2#`K}jl| zmFMU3U_!-cP<8^eD3E@^Dq4f08h7#qBpV0rLOcmKCe^stKR+TSFMG-9ga7)1vk6QM z+Fr2O_XJS4< zzG(R9XEA|mTYEkbUD|7~+Jbd5fu4=s2tE+JI}(F%H~(@3ep=~)sAM5U?m{qM`&$F_w%b{48|e?o($dq z@&wDPepCmlg@#_P= z==A)36|L%fF(XAWzWRMNr;$YdCJpTNlf=QAEaqZy1`9XJ@VQ99Oa8qV{P^jmD{nc1 z>t|5UHbX`N*7GVJW9_(Dg`TNZr@`q9V`_{gvKtw&*TUotDqs8uDxVmZ7(B9jWVri> ze7G1hG0}Q^$jq_gGlCX0+=iqheQa zjOh2AGrt!0KGGZ<5^1z3p8l4R|JPl1&lA?6_PFhJm;U>+inC1wNd~uPt|*&Xs4hB2 z(L|F#5F-a=zIp@Q>?GBB%<^-E(F$98!EX=#KZlwBOz0b_jg$p^3_1gq zCRGvEG^rQ!dqd?ww(F#>`rvmQ1m0hM=0pTU1UlKBVeG4_veL{_sEaP1WtK+IU9<-2 z4Cu!ApEdkRk>G7vCQ!ck|f$k$csPkUT}qKlbon6dv# zPbC=Ge|;vUAY@XrL>=mn+;M%x8~Xz~#|n}5T9yS|42Cck<}Pg$0Y~9k%<|Ke)uZxj-8e=jgAi5uI#ov3GlCz;31|qX{<@wlQ5|+8N6h z!CsH&iJBltAS{_j;lq1(1entw0I|nj#5|OM?i<=Ix!{i_SC$JZzQI0xvU%eA&nI z7*(k40ln%~-+p*|&!!Lf`sz2YFtW)novw0vAD@&BTG^W}yI1@fmb=4$2W3;}ceqsO z;I1&AXIjI!zZFEZUL}AmKlT1U`S1g&+YBVd2d2H^L*~&Y?&Wyd%;c_Zc1rjbvm^yd z0E>MPK&|J0dc|>1&b69LuRH&DxQZ~gnBNpIa-+uR2MtUWLsP|1*9Nc+%|Y0#+*dv#d={y0Eduu8bs+Fc3IJC?$qc++1no}nIb8t$@ zk$1UAFa#K-d}xKALkqh-eEt5<`1_}&e}af8JI;8_()aLV3INjuy=rgPI8;Cl;<}hQ z>))G>Q*jkGzu(8>7EDJEJ9Ps^R7#>-mikCnz6K(d%kEj}96ZOH#8IX(6~!C}IU08V zQ#wFGr+(0^D0Jc)mTa4bQ^By}6LM|+)@nmu!JXrwZCR$kO6Z6?Sw^t;c*vkL2{Cxd zSKszTgVUrxTh~o@IM&ek!t4L>4<7~FIKT9aYgjZV>7m!2c2I7+Y=Xda>E+MkJ@~kB zwR0=QX|8zy;pC5AzLMgRU;1K}Ge?}}YZ~7iT99dn$yt(VsiV?jq0chM=Nl}0Y*&sQ z2`Zc|xHE1U{>?etO4Zn*iu*5plKa9GDWhvnC}DynR|Gd3*|V*%mPbb2 zAgOh;yMhEHR`B@W^1@S7S@^5J>!mx?DA&f(oGZ@K=zVcPdnj^$k{XhY)9L;W=DR!w7FwjD{+rfSwQ8>N+`Q= zT+)@l_~d-6SN_bWPhwGaVKsn?(0@1yUi zn69wiYbA(Om^Bht{N%?@YLqwoAw`iiDlRH1%3Bxd-*+5SD8`8|VMuYED#k=(YMFO$ z1|Mi|Tzp^z4t#}teBYc-DV;ceq=O8)?irwY;5J+F)Lsg{8w_|p-#k^HJ?d=Yb&L=W zbXtn5F8j)Zcaz6IIwM+OGZ@naXzm6bHocQ_O-iV1TKVR}J74T`(=pt9t0XfyOKAiD zzX@xmr7L!*MlRj`3r}k`ylP_~HFx~3P9{lc@_5e(2joLuv;kkG)7|nPkJDZ9nkfW> zJe0(#z|9%;$}i@r8F^yvNeTp{z@eFAN;Cv2+DQit zrlWGgQ>>bi}(q5E79}~{^_ycB@`~Vw>1TVn_jwUE_OM|^ZR{F3dj?zt)A*2wfmwHLVxFk`o2!x0fv}^~O%L?5pyD?21T0yYPG7 z7oV=3F9BfbndCmbTM(!rG~(AHM+MA92!bYb01i*13{uYZq+O-$h~c;%_}5UC{0zs%u_+ zhf@hJd0N3TG@{Zdq_6T2m~hGQLZ>=cOVyi{kbQQ+b^d#Q_aF|=4Hd?v*Szr5KCw+Sn{wK~MPO**A{z?an@c_mj5jI7chiP2hBah&HTj=ZKQvc7BX z*S+O0l|446i)74{E1((h()0`YTc_;1`sU{zlWi!VOm_XD?@HkYj-ocqeddzfF?cim zGd1V_Kj%5|AP|*KIo@CU6Nds3*-4Fvdf47|)jdy7`~2q_q1LHv#Iv_YPmw^bA`(BvWuGFe&!OdYTy&3+Wv)!eDlk!(l?XtDs@MlNj!D{X*89KMCpLzU^>n-a)B*XJoE{`}J6%Cbo73itQ zSe-Q2_wfY~@u&$VyfptMk2j>i50!tj_N*^_6rH|NUz393f@}n=>=dHcgp(1FWUEj||w) z>F*i{=&AQ=-c8zAI0^tQK+?apyPcl?_jR|N)>FN5Su@jonhOL%<6Wg0H>rX=sJfVL z@2)R8Tv~kB^NE&^{P^degaJ9x|ncI*9{+9B11_mqE7O-`ft z?CE3mB=oY|&01?NGLzw4%{nuoVO3A5h^q`rvEeK^ODl}@puwImHig(gS1G4*ruaft zKame4IIHTC3QX_KL18v8|6FQ&nair&RR;P!y$X1dHZNRe6C|x6lK|WCM}g-!7`0v% zb-u2jbmQ?{uE$Q333D^Y%WkoqB<~`L)QF`{&=Hw@qE5gICy{g)c#mt^RZ@yh4Xz^U zI92QX)v9NGWh9R0c$GD$Lb+(eKiZvFu~oILvdM#KcG}?6){P3wdRrC^-g*TssmkcY(2QhNY`id0WE4fcHlF8YEs$pjCH@C>%m)s*2 zly4inOjA5us4N4|%1CY6_WQp7wMh!{Z~v?3A0&g!InuIW#3fpcJZsO8e@=@|>IYle zbMTP6e02)Y6rt;J+oPtjZD!e0QOY%wIy@|+N_i=z<>)OrD*dsH0*X*tUTcfAVc#!x z7L6e3cyc5dCv1)~wA53%eo}!wcSFQGEhS~lw~_aoJc=}mw8+=bAO+>6W zQ$a}In=8JaLkfkyW0gz3;Pvq(bIe)anj;XCwj#Ba%m#*_Hug8m^p?CjHmXb|7ffJi zo1VC5WFo$h(@o4w$hX~~q!V+fC-A)H34e7;haHxp*ZjbG`No*FCc^A* z$`y6PGmhtMa;NLR{ZHG~l0VR=PnVu=a6&#=aSvGzek_K{2CG4pPc*p6rX?m$Y+WLX zs(0SUezzs3 zqNLX!g)9%J3&B>MIGIiURIyHal28o;qB4HL`;W@Nx|?9G-jdYfmL)&H8+R&(_!2l1 zEq>&hqu(Viv-^?%d}LELhuY>I+%%oUl_sK2j?r5B#BEFU9*{ptQ6JrYBSoA39z3325V4&9YVTSo)~f(*NFe%3_I zSAFj14kyW{A9ye+TW?81w)b)9lun5{U#iFjUbknr*#=7{#==6bIkQ65cg0Al##hOg z4>oUopUIi{vQu~oOR*3z4l}?=!@RCSZhg-|Qoh{QEh%1e={V7N)@6!$efy_QF~OZ` z*3s5AlKs+`KK?kN%p9DgrMUuIL$Sm2l0Wc&fAXZn3C!WDF`Y}Nu{E>`x4m)VMBT=N z9`v}RDa5y-{gcx)U(mBdhJw|F{|%o~(I--c!49J#CN1~ewW&S%5XRoYLHqx5TghUZ zOQ_0K#kAMtXxK87j7Uf3P%M*<>khu{;=IBXO^VSzYtOZ(sPH7EU}B}b4o`vuHiv*D za6Nnkzdd`zKDnMuywYW@ZT-G|V|{)&li!Qbn+Shm@4WqYn46A3`t{2`kfzBYHeoU` zG8syr2Ui<$=j3|{B~%@<(6G$6`Q&kr)-jbXI(beNkG%KSzpDzf2`BFH_NCmg`+7;8 z<4ew+9iSjD<_kdQ6eAZD|ZL{g%CM)?edS(5)g|0 z{Ia)jNV2(xoL|)YdQ-`jDupwIu3de+X+80f-Fv_BPY!__(T>b589`omkyO-Y;AgDj zt>%mAe~^>?T)V(IdVa;fyX6oLPVjQYqkpU2Bte}XaMluX`)Vjw6#Nj$$?7Ry&2p1= zwJIUy>>*MyU-`0LSwyb+r9bUb&*#LRw>Xq>7LsaIZkx$Q;<)6F?=p)QG8!&m+TalR z(n=)f?+NN2I>Fga_>aw58gf;p>)n>bv8&|8xrCZkW}@Iz9-7Cqgg3LO@Akw*vUcjS zm%jX=RXytHSOAp?3Cu8sdUhgNRoixL0>)AnKeAa{6URd9s|v=0IxR@{Lz`g@ds@-~ z{K_)gOo-d=9bY;mPF{^uRkm`wo!P(`8Im4$iATKW*m>T+x}5Ki6P;~fm65!p%RlkT z1YnT8{=WGiL)SKa0pfMPH@yAm$M`AtlqVW9cj`7zzmP9q@zsBDa`^{;zpSc+p?7%F zA*!7OytRod)OViEKet=Lcd9brJwNrb0~LlYIyvm7tD*G{s-99bJu4hma;G7c(qqQM z(j0^+&QI5TGV|A`Z6OSHzpuIC>znRKh!(rTFDN^oPz%el{;8}V{ecM#6VlqbBtCdw zLcyxilnbT=3c8WBWha_qlWThcU^wmYc z$)-O^#60PH&lfxgJGHNy?6|+3h}Qf3;e=@)d&+5gj4@=fJoFjF;jh{(u~iDC`sl~b zEJUhAT8d7pt$CuoP8`gaQ*$S!ejJ}4BMGA%pRa;S$}j!o7q351Sd(YAFPPd=9y9%e zG#H7d+ku;|J3Px&lyC-2)i9;9?o>c55kvnW_EIDy9@F^1KAVGZ#>narRI}PElo6o?LE)OBw9u3etK69{0PW4r1d0B_4^#9*}VcxgoT9sW^ zJhrT9K!_SWPq4Wvs#RWWe0jBHg_&$2$#>6vSl{E`c$$Y5Zao}{T_f?nIc}S9nNS&` zFB5ExMimDO5G&BwP-!$_M3QRY!{2oKgNI*)v&vmFdt!(_uX(uc?xTAbV*+bk#Yv=JN;9GJSx0qo zCYx!u97h7nGQXDsqY5(TJ;ua-pgZ=g{pVjxFdUnM2_BaSOIMKBT&D4MQhgTRZ3;p?rxV}npMBd&yXiQVCLZW5 zZrFqZxx9GHA*XtlCXM0wT-bpjWF+bJMkbaohbL}c29PwsM?UA^A#su4dCztzzu(oL z`@9o#U+WQ75-wccdf9qz(*#IdrFqN*b5YMyfSK5^oe2s92aw6j&{)mrqNO znJFwOEriRCy*jSl*Qdfz$HxmX8hX?eJGoB!3R4Ws22vwWS?P2y)eoY`thjnFwbU}g zniaIZ3LZi|RQ}PV;_*D@s_8{$_-O(LZgUDLJ55!IDURduGv>LrOHvG)Hu%un{?}8x zRdaCUSS`sfesQb^Gn-1O=>SVQ_4GHN3Ob9gTQ6=t+O>c2b;pyZs&TE(Z#p2^H7_@} z{}LUnOrK-ZN3b+I@lwGJ@$Jc@WLt7$*u;nbsH)Dqwv1G&x?N5WJEi+iQI%p6yR(#$ zSe+>;Gd9VB!e)#|V>IoQpNO|~lj%h+J>?Upm!A0ABd>Nfzhp>tey5U*rTL@HMr*dy z+kUzPgPVTt&-+_?Kei(yqq$kL7gehd*}td~;3QjrDMu%4&PKD!AWUQkvgs_;rddKZ zU7{pP9H;mDT_C-=|D|G@25JtSlVd!!6Dw49#xm1!&ohBiJuWl3)D!B>B~N6K{@vd{ zP}jXm9tdrhasU~HlrE#E-h%;f$rKk!6jP)ot(SvjDuU$6k1S6P!^D7nmMq&8k!OF+ zv7^ZalH0szDk=S}Fz5b$Y+Skq$m9C!a@~j3q~RDR8ZIE=7XCEsDZ??$PCAEIDq0sT z&iA15FW+QRSs{mIv+U-3+-W)gspR}b}<%jM@weq0qC z)`T4yr)nm1X$IFk=1_x80akxIJis0YBlB-(-iF%ytahR^0XI2_YK))l&bI~N>%FLxBk^}f^F@%cUebN#7?l`B~@a2 z^QS&1b#)cJYItkCy_jvyKD?YlsV()>n{~yoOxbL$GBL&o!5qepddAbOz&xkps@n4V zPETPCgzkUqCr{yY4|($&vi&{ew?2KOWl%nQFFj-0qPj0t;<*Z3C&qe#MWx)_@00~j zrD5X5=36&^0uGJUfD6LgV1i3-iYb}GwB%mXH)X`X`r0ZYItZx#mBXXT6{Cn{0A82X zn#UNANh}F!IFEme;hJ_M76*GNxeA_fo+7<{AyfA%#pYMQ^duim8SHr}y+v{|sb&G$;L^+TKE29w&oyYlhFAXH(yQp>JlYTg}n?HJP7E9ctG@i*Tp zBbgy9MV{RaXT=z1H9-&YR}?RdTT>q4Wa`uL);NbN9!gSH+?{c^G0wNNJ_z z68JJ=z3g8eI(P;=+^C+3LTY_jcHVOy+GK*q-2Rv`wuC#qS$(1tbiHD!f8(}eSoC6^ z{;j0Yt{+*Pg@1oINGP7mlSmW2tEq8kJ&V>GqcIZ+xhC_k8)|m(0i>AMEs_LUOwG zb7!J!4jD^}Olv3kY%L*ywBu7VH6M6tqKnL3Js$3SoH<=$Kwc``{H-K@it= zbovRb#CE18i>koJ*)<`9N`~*`SH9=K2)0A!GjQUtbi3%^$6cp!f7N#%(HcMV>67MF zopEW!H>^^*?waQu1~%Gm`Q{QJ^kJK8a25#My<$qm=NEi2FEAOHOH>#=7LZcqHIKTv1tLop}C zVbD7K#C;C7XuDa;BWm4MVC{%btetS7=R5iGcT8y;Z!hm@$Z!!8a2h#5f9snw!ru6( zmz|`s_!pbSL_xaB({#srTKXNY^KMj$%p0UaPNy>cPLp(DNmHf43DX;hYk?UIGiz?v zrgs#Fd$&VVX^t%)qKokhD`_xyXs zEH>NfOoOjqtg?8^hhwkq^MkiLsRt8%GJ98Zr|3-l%{&LwXEf7|4g?>qZ^lV4*(bgJ zdvBXMu55FElNcFz2ovx84gLL70-DGYUrbp{=dyb*y;T>KfffERG}!)|-$xeHs_d{y z(~mKeG&4+{V<;~?FsiY!d~9ekx-m>WAEFcN-$Vb~Z@**wGd#fFDH33IGp0e%s{>j(>hGMQCWGOr#DLD}k zbFC;5Q#5rT21#v6|%IbP3zc6 zoT`KyLYmif0nHZCXL(vE7g%+1^zXXQ@n5fIyx&`2m;)&X)aO214`NbY^Ls@f4CxG& zaf-81qrLTnEhoAJy|dE~!Ui$Qv9%b55J-0dg5!JtYe6MYo}CJ*Jn-Cn z3*(uKqo4=%6mDLMPwIxtf3|)M_R(zZPfe@&N3S{1x*?}9z*-eVxFqaJo5CYQ-gWBoS<2Xr~R+STu4*MpJ@S7Fu8tcK5*IN z!Q%Vb)$#ogTXwA~iQi!n4@Fp*Nj43#e6VWn@7^vt)Un?B2Y#)SGGi?}`ZN@w;5-#n zjXOz%hBbHBPX4ryrJZoK{`5e5ikm)1nI(2&ZXXj$o7rS8__VXFrD^eczgcv%IkoBFG=V`ID34Tt_ zK2aUpp-hM=bZw}z9yX+%8dVUG%Eg*rNK@X_{Qh`9@4YwfJ?Ff8uk}o8ZxHW9zj*7_bXS#*;*BP)X>q*c?JLdNkexI(9p2D>K%2Xc)NKe-esM{ZzxFQk)0CRyHp}C z=cRn0u>^nUq@f00#Xb1Z2n{ZnAi=lVOZ<5I`u0C0c+P4GUU$0$yZ4{I|$ccT41-z7jqDM~OUlop=^im)PFxHD0x&#um@k*v4HN z+fdNbmY|gSe5R$>TS{qh7x#Xy{X)u5Um)4FJvF;#k>;v@t=YdCX=3PRNp#pRiPQjz zZ@X6#*}9UrzrLh1-Q6|zc}btOQqq-(l-j?prMs3%>Dl>Gy5~eG&G}h!4UW<5g2C>0 zo_MNe$3`{#_!pYpUq#E^qEg1+UD#I2J;zEp+FHuXXGrDPmn7BdMom6_ye6ODt;zlS zH2KRtnw-B>lRNrpGP6^Y-I{87Wl}=FwboF-shXH|NWu+{)NpvUhR^;|%dH1V=prAV zL1;=nK954Z<&yvBMkz!$N$#fOGbrSJJf&C9!*VIsi`gjEnN}R(j{FKpG|yGEs1Y@O1uYGt5*kz zcjM*aeW4C7-?59dP~9*22~SFX**?jyc#8MJ$u~*<gwu;O{P}HC z3^&!{!jH9>xK4}F543p0D;ny)TtnA&(a`6gYiRo|8hpNq1gF0*!R@cQwfD+QiFLh5 zW1YuqY|(Y@{@Pm2{+!V4^yfAEz;nvK->sTddgi1upsBk)cv}5GOcwvkkQbX?@)X>in z4c&J@L(l!|UTcjLG=%Gv^_u#iJ3s4Gs*a{IuWRbnEgHMNs@`^pulB}_#(H0)vFLjm zyW>ucy;of`YkrW-zaLBHw~}N}Z>8Bb|7x}v*6cZNXg1kRv+0A})&6l))IYzk_^M|tJ_d=KV>90Z>Z$n?j^ZR6C|JRBl)+ROMdQT$v-$(l66;0(mz8I z-+m#9InPS|`~H&O)L9CJo5k~84e>lXOFZjdkeNQrhvE zl(yY3r4RPA8hXAsL&76^Ncg;l%ANS)uTuQt9xX*jYw?>;w7B?HEmigL?_>6s-BxL^!Icu&_qMQVj&97axl)JyCAoV~S3Ix$olx(tbNQ@`FP7YAVxLcO;tv{M zT$V(wPbE>sC-D#V@;Qg=U!mc%&e!k*Q3>DtpkSvYMh%q2$d!_K;Uq~6KUeTMqB^5n z#9Q{3O1nD+)q%US#-9>>VxdN-9--0MUubm2V2z!9sz!JA)O_3Ol5caJoNNZ_nbB~brw@z3Ze z{df$5gM-ZwniSD zC6Sx65}Ey%L?%oW)YVX%zcrNp$KB7F$BG=feT4$`q;A`!srTPfysy++2-)`K1#1y(9no$d=ubEpO4x_9un){mMbg z+JE~h)_!&lygofiLzmsBp_@0d%S5VvE#a+gC310EB1b*K8WoR>ZvkBJa?QXo&M}8o-sSbbJFjUs2q@F|Hs|F8Tz3nXZ@hbp4Vw|GomKJxOt;O$8(Bj=^YH|2S>Ywtd`YV&w|Iv?Je}T6DxOFjW1ZzSdUPl7Y zyeWb6FO$H%Z%ClWIg+2}k-`c0NukO=l7H?7DJ=^~>5FrubnssZ{rjA&KNlUPp^w5u z4dG6|X*k!D2&6J@y;hd`wK8d$hTlC(!iRpB@PdW{1QMP%SRxm`B;mKOBue&w{i+0x zTPA@9Q^Y^sr@r@hs&CF+3RD}s=4%OF)KiFjM|>lp&(3qGV6ClzeOGDjs&6%SL{f8g zR8zJ3YwGG1nriVWpTTX1QX_BH@KbL~cgfB~IvE>q0*pR1% z2KRd$Mn9{+1~;Fjm7%kx(z?A=IxLp*FE2`^{b^F(^Nob(ektK+j+F41r%3qmTlu|b z>En_;=2eH$J~^IM4tNcy0yxii=DX_s_A~LW42m!RsQRk))Y8edwNzT5`G0Fm_{Oag ze(rfe#|hu}oCM$At+-Ns7ij9tQJRV@cU`x*Uy}!)(yVW#LovXAKlhdF(X|BC4ZFD1 zalaO~Jf)>`TWWFJdM$3KxYe+0nfhu?S6{DV)#p7`eO-ns5Jqg~PK~{o)7Vp|Yoz9> z5((cSiPs{Mkk2J?>oi{MzNH>gp7fiPXDpJ`*Q>NN;~yzkdqARtFOyi+A+D3`K3YRJ z1vK=`DnC5cVXzrP&nmcEx=I*>$bI-jfxtm{h zPY-Cwf1jp0SOu{w(er*uRHgGg4`uk@B4BQmHjw_;+mO#O8L;*vfSpKew6yqsG3k!QYEJkWZzu zL@G_LkzDr)noGa0xzqkoAcx#7Sxv0FRTA4`4t+KCOMHu067dz1xS)}5%R+ejim9WRl8Yf5_i*OF<_OEP^=mrSHR|2-3&AeqzCQusctg)iRIVteB<#%)r~ zTPV55D_bPcw=26!wn0#HhjvKrpO+=KVY}r1%?iF}=RVgE~>e3ZQX>YULlcYwl}87By#S1+-LDBPiVZ`F!xMwKcSg=7oH>DH%5zh-Ob`1|FnPy z`A!G#6DtiWKVRMq7 zYx=iGB)7G_lbxE8cN!G&%81WmRs{Qb;umAaY#(ctv;7>z53*;`0Tj$j@+&BKvRLXP@G^ogCR`cA0Mi3)d%X^c!~PY zpR2xuliV|Sqqlkjjm6XL4u^=cY4Lc6iszo=B(?8=0NJE|F{HC!B)01#r&s)XDUof= zxM*e0C}r0%R9blCEh#p;OAB)iO*P!1xu@1D+1e`?N_k>j%I%F|m{ya!=1$EGwtolL zw|jQ`5Y64UgxUr>(s}bG5V?ta1AC6$@zS3qe8xkr+kw%X^1Kp_AN7k?s-<1e=rKzw zttLq2+KEz$?2t}kam?*9Ahx%z5#{kt)V(B&;P6sfJ~U4x&Z9yF!c(!U&=HWcE)bDiBmJkzR)=ZVwB zGjk#lPUO!YB-(kNL{Gj!a{aDlUCk|M?E1yy6V)@Qy?9cl{a8Ote)1!))RA6IwXUtH zi~eDcB%8Hb3ls*Ff7_w-`Q0Ra+-H({>vK(gdy7JK;uAHzlYRdv&DFn}*d5h8`|vML z(bzgp$v>e!jeo&i?|el*3f}4Wk8~F}U%Qg}=DR(}?aof}_QmJyR5Qn3)u2-oX*^AsPlDTB3mee$yU9%`@t8Xs|Ikrx{m3z~8h{G=JNY~+j!C35@45_vf1 z@R%VX?Ap?7Kc6(Y|2G53*QsP;ryRT4@u5eKlJb$$B!Bk^$q(8s`E#0*4|?BOA$gNQ!#YUj_MMXG zI8_pFJm|QvAqZetyeH@rs7Txa8~17Hgce$SN& zjI1Yt;#9}`faQnXObnm9?ljF^mLcEHUHp>frl%y=<~@F&pL!N`s{8{r$+Po|K5?JZ z4fQ0k;~|Ft|M*oxUmGvFWweHm8LQ>t=Sz9r?e4;w#ya`YP-Ca@5F} zxQu_JI>r0%Vuy;~aVAxNAm&)OvBnzyuCW2z|LZ!`ea?H6NEL{JeCURowKVxtDTOv_ z>Bf3mn)x(UbMF6MAF8J`UOMfXF4QQjlX>o1ES{EEOTK9S? z|7-fB#_cNIx?W1z0b1%NTDrfp(mS#^TZ1Hc#%~flIwrx> z;}SKt5PdJEv8G>Z^zoB5cEO7pef@1E8nV3`xo3$aTWphL$EzgS@Dxex`BKssJwqmu zPQFEFgWAGxE8P8ug_`>4kmKTacUSJS@8gQU0|e{c{4bSuc2)Zt>xNtts(sx#4ux19 z1t*8z>z~t1{A2$niY+|7RSGYkEQQ5iNMS;E*HNvDLaU-a2tDbO1f343*w~lh>R+g* z{iBjuFiA3Bca(5%b0M(aQ5UeL@wF-QKuWP~-)n5b#fp96ZQ4mZKNf^{o$c-monuqT zekO&R>r26VKnkbSk-|}zNhNi$LmTKF_Pk474rpn3f2Zl9i`)01$69IV%A2Uw*7{_{Fzs&flxyo+Qzv;C%z%+aWS&!kZDK3S2Gf7 zJ4J&pUE|z`&-!WzY~0pN;ICUHXqr1Tc&z>unlc^t@X?FSpU?wCDJ|Blq=| zNMWXAUg~2AAbVn0h034VH&U|I_fo~CN?__#a?fj;ZWMCw_uX(!Urrxc3&Z#4~w`P-8e$bXeW^dhEMCbYk$HE9<9fWo1jPY&%CQ z8v_bOzM_{YSHjmNg_0VI%~!81cf4`rr{bO2g;fW;DqH7(Qg1OnoO%8X$*dhAr305s z*}v6sl6G%NY2|2k{Bp-R#{Ua@p{u=e&w!r7fGhkvD5~@YkfZ$Za&FAw`n!W-xEhp5jq?%lq7%E za3Mm1XE$PjhLZM~O?|`pW7e&`txfSg`*gYtR0OXWrrr-0h_7W8ry%zIOOuh`1?nN* z8vx<8x+4N{h z{5@E_*%KU>#@(*KbqgfW?*R8g^xS(SvSE&JkJ)~W9o3c@&78BwX?G9);xsqx zCunb`L{f=gc8Pcv=f$(*Bk_#?PP~={;66rW#h%`llJW=dO8K>fl>Z(o`3>`=aB*Gw z#`!PSJKq)#pLxj40|3?kk89GyHfI40A$#*y$2N?CWJe#N*^hry&+AR8BYHYtgWD7F$F6F^`b+h_iI2!J{Ae2qKRlLh zD(Y9Z<^s)n&eiN~trePQHaJk}(mraQyE#>{lamgmb-Iw9Ai8C>Mi*Q6g*H<(EfRa< zt@;}I@D-thWC#WM2)K=SP2qi2h&qALgHMi^;O@P2SA$>9Rr=qiI^7ZVczA+QjH}ri)O@iBiHJ7K1fi{h$Qon*uio5A#_4RydXW2bmpo6dCQBf_i-vIFCerA8Ip`mqL&d&udR zpFJpnA9_gOos%T+?a6d3qS<8<4Ll{#oWLHbnYWtZP}5jfP%{2Hxu7Li80 z&npsNzt4TW(lQ3nfbG3fi>nrwWV<1T01_*sf=-q6R$=u>j+vy%T_9ZwT?_-JH@s{f#h6eKGg5)opE&1-3N^bYHl5V`& zrFd%GEU8rwDpfGUx=7a)?iu5hIpd;mFp8zWxT5LD_ZY3I8QnBB>lsZgk8mZ2j`>7` zOMcKqXtbp3|10V17CObx@GAHKFm-CaRd357q$Z!Msr5amdJ|upzfVkNcr1ADP*?rN zyyJ8cc%^-A6>1!&ndCowjlNNS$qCdNkbq$t=q7VDke#ShkznH&h~I*q1aTen1W3%_ zwF7w=hEI((U(Wt6&_u#GDDgk=>hkBDJ7Ld}-4h-EbMyC%tK>fXQ<#de@3AAYN8BK+ zAV7ci^Z&Xd{JnR(v13xp>2Z(xNA_30+(<fXlQn*mXq7$gncncJLyLw8qEDIkJnUqb5`*uE3TC^C&X4IrOr3#G&6!G7oG4}8 z0x2z>Dy0`}Js6@OKmXn5)`t1L`P(}>9~WI6?p5c<*|FiJq7uOkiG49li>5~uK5xu~ zPhrD&4PG|G`P^fB@v~3;wSdnbN;uaEK5uiEP>IpCnNYf*mQ$Sj&lB&Sm2O8`ddb)E zcJ-P1hzf`6us?cOtw&xVzVHU_#P~1kiB`PTMz{R+Z=*fp0Xkf1$|AatA->A2 z4a&QD``J$)W-5&7FuMZJcay;A%N@#GyHEnw-Pp0gL05k$p1sGr%!~Q)sF-+m(HDNv z=vP;2^!HPgYt#~9#V1B6FVlWX$B$y}B;EfA$Ki~>h3@UDA@dePuRK8h7TVXtAsf?L zz$6Z{2vkXU@%b(0xz^|}flL<(Ts?r@i*>+KX2@iK>(ODkzc$@JqF5;Sh7c=xQ+|Q! zs*RHP(|v`X+r$wPezE~k3cJYoJE@ja1Ardz+I@nG4DG>lph*6WTn~^6Y7KHfUjI)_ zjS-a`nos4Pi85UFv*^d?YjnabN=9Sr)!Sn|*B~k^eEYIxZ0OvLIa1R5ujXeHeC7`F zcw`W`S0;^O;>|zwUS@Ipcb!d#6K+l{Fq9fl^uve9=jp8A>t$am-zU8b}5J9TT}8@t##r6D%|L9ft-15$@$X=ywv^@h-?Pu}>l&@IMm9Q}sNT;R2j zs;6)W|8wQgIg!ZA|0S_+KBV3hH>5xp8L7T&-eL-q@Aoj zOki0`B(=J=>sZ!t=~F$nS&=A7JzHOy5(I9urwqQ0?w&t+mgLXqAcaA1GGPH8QRuaq z7t~KCTWVkAvIhe{5M(Q%=%sUqB#m6{Gw(?Ha#Ns81p}JlDn}jwyTX}mfg%T_ z!BmW0zVm-qx(244`l83sH$t)`0qm}~NMillE=_1^d9ZkuQ-CZb6*L}zE`sbZ)NrPA z?mOKm0|_u=0_H%Fe=G6#i#IH%;aE>0n9xsjC*KZhq8 z**Ss^2oh0D<5*&dX&kFy^x!QG2|o*XI8(fqNFuqklO|W5rAef2;P)D*!81ZNqQ0}^ zQP+hG<-yZJ(uxS}!9h}Z@fLwg;2*ZkB?awoA~9HauC){fccT&-*!l~10aIw^mLSbU zkFD++CKQp_L#M%2IQ4ZY*n6U~imz?mFk@@UGWe8qq)?qLSL$K-@-wRn9E`i)3IBK_sLA;V{l1fQqb zsRAE3aN;}(9DT1#utD20Rh9_YF06Z|Y0maYs<_ z^u5qaA*c1>4n1=BLgr->*M7!*m>3dr87|wWL|*2cuxDTn16hB|;xR)8ZPV&GQ{Zso zErDAOJ?I>1XlmxBV~(VFsIzn4#%~ayl2Ra`WM!eGBugDM+2LJ70O0=Et-zR6aIET~ z67X)yi?3-Tmxu$m!LDna)ZqIa-Re4wS1=U;Wyn?=cpQ|M+@}{w?$MkA$1$_pyqV^! ze&?`_JX89ms5ItdAu-APCOlYF3-fK5+%j~7t`)qxox7X&56_6J z+)`GD@64%V3I)D4GWP$Q1>n>Zkk#DBCGzVNL|ybx_gcpLyM78an7(||r>T&!a~Y02 zte`=iOMf(7GBv-Ux}JXJH(|CM3O8mjTDO(Z^l2J;`%MjfyhV$%pLAZjdHeV~rFu8I z6$IZO3X`pi@b0T5a`KA;g)nl05(T2F2hMCE{v9Vc7XkU*StIBXdM(k#t_d9dL^Fwx zn!_0mDpZEC@i`ajw3Pg>Q=~9_ixl3pROGU^nT$j#2mWp0_@AY4wsDvX2e89~yX8Ok zQt8jXd5>godC*Nq**ZR)_M{qWosZq-y%*?W0I5QIqCWNOSCaevPRV^d-ucZ~Q&10W zT_O#~eDodPIsL*k3?d{`OUYSg#f8Xf-byzeWZogUF)|v4`3oQ3`v1wI>(8b0QJmX} zvkT>u-gGXExi83sLs_N&2sCIZwBVXO%*%J$xZ16+Q`@;a)*LuJ);2uz-c_A%et%Xt z|AMrWbwRion52T*iisaOG={dRmP0{ACcx$?;1u!e9S^vvGpnX0REpIrr1ZveYA$eG zm@9eA5{1Z>nE%Xt0Jt^W04STL&NGpD{9-M9)r-l@V%0m{97njRG7*MkvAH1RE0`Za zQqX?J^eQS}uU1sinesxFg_8mVjcyIy#`K%R`TXJwbKTmv*%5Si%nxCT4O$q{4O|f* z9Mi*?kG8Bd<`@Fozi>Gp%VQ8t*qX>b*2Zx;tQP&hp1?lwtgdGN^DLzxM|1lr+LG1g_- zYmP@N09v@tUi7JF_Xsb2E!&X-3>LkN@TJJaEyF@3ChChXCuk)NsD533!&% zY76oamHD4Cl?MH&^3HrF0+AVtoMyGl5Civ!y#^@P-c_5N9)q<3E{@q7(=NcDkyQuE z07B#cv)xj=RzJbKnC(KYe%WMdN612P&snye49gH#_W2{7tA)=UJQ@zy;rTL~L*VnU z`Jn6>pOK+4l~=AOg|3$i_7rtD%Ub{e9Ofy5S)jwSv!Bv=5B@s8an3apraf6%OIO& z?q}qxJ0Hau3;=bV2%P( zipV12x!^N`cA9$rP3Kq`6CXV0KcC({6-?wI2aVlB?7m>NOK`x4g%ULX0;Xk%1A+Y> zd5qjPHhi)2|1~e1{w*e&;JaA1AoIH^nufKYY_dyY_Hs>6L7tx--6r0m7xoj*i{8K0 zc`o3Ze|2)6HJ;{1kzC`H=4=-Ba890>+mslpGCxh%+)=-4?vxi5nNnabsEC+|@+6KW z3ShzyewWSra+1pO+vz>00l5%H?@Fd}96jaCO}!;^bjp45=G{?=1J7ILVR8d`YwkA7 zRx`VflwISOTxB$GDgEfL0(FIX0qY=8-Pv7E$x+WhdJQwGoJs@^w`nyb6u}FqGDA(W z+$Ov~bYm)l;F!pI+jJW;v(c(gN_5bC;Y>4>b<2ffuGZhNlOm~t#0jSDGKRDHtW3F0 zPWs5sJT3;c9&_Y+Sz;!sMO5$!aqim3X?B3In@9X z;cbUVrE6WU{w4+lO|T+OH`vUeAO>q--C>j#Kmwq`gf14&_U&Bpw5Tqq9HO_y>}&-M z;TS7ySQ`|>@E>8W#dX<0YTl{;YK(Z8SWbz^qUFH{EOy(`Ptg%0*`#=geZcOuP!qq#r~oWp02Blw z5&Xl*fDKe}(_5&CLjRry6aph3+eo|>3Kks$Cipxk*@61wYmAD3n%BKtAb7$s12~5G zkKil{H?|<(x6vO`6_^a8E)j-yBOj~(lab=j+rS-`7mOf_h6HMzNEwg@1q=%_DQF0x zh$CiVG8+IG)&<4nr;l@&4K2knylcVM!iu3<_`Qi14s546huAeux_hptIa5A!i;D!ji-?VX{#G*lojaRsMl)YxIIgAt`KwJM94h$mpgD3`RHP(g=(NTVFdr)Y^ zAbw?An!&f`BLy)Qlwk_$r(suEW}8|62iGPJjzgz?@(f45+Yb@ z@*a#U#{Lm8xBoZ$`&SCN191hIUpP{L(Q^wzA}+$w9N6pthW0($1*HR%SUF$~^T}cx zk-P6Akmh)g#ZO4Q)||;B4mc_VB8R^VTOWkZEPWCHiyIM0niR<9x4WtqOr~IfcW5h? zBaR(dnIc(jc^|SK9yWa(=ZAcUMy@49gbH;fLoM3Ec$6 zlV~C;HQ0c-eb2}YWC=*wh%ni@2g#)wZozsGU+lfZWS%iTtVdg$BzSnH2#N#D-2b_F zwp{8=x`Vsj*NDk2s?LFAg#Du}?X}zSAC)}|Ts#9n7STA1Ay#%7!Ljf>&GG-;A)WB( ziy2R7@_jpY0&UK|r%4HE3pK&sYrOR=;b84)?VMz^;bD$B@%Rce=1OD!rp#1o{<8vD zbMPM^)56Rk@;Ls2M~zrChGQhCX(r=hObu(dSlvSA80iCSfpmyilMUZ@ZZ%LV@BgEf zZ|cy9g@lH`hu0Ca$PNep%@~3iPI=Pc3*||WJ{x|`-#AC|{f?)JLl|J;2}-jd#CQC_ zph0F}H9^=P$BfX*K`7YnAZsw)cxZ#!7Qxt80;wH+3IPNFTtX-u0$@@WL&6|q7K3B7 z$X=fvA#uY@jSZUM0&@`laUn7?XM{VX0;8{>#J}4>B#u6?$UE#pM8GM*m}o`Q=i)pK zAb_;awQhEHs73#2#o`+Bk9f9O>>uFCY)V!vR3B6x$QrntxX7WIjgtg}R59;2v|0HY z5uILhltUiay+9!ljLkSFxw}ILrn@WVsOEhwQ+m^Scu99k_Pup zR!V9ZzO^a?@f?+wgP!OrIL^XEe2%x+Xg|eri}rFX#!NW^Jd8j??89L&CiNgBSbCxb zEppO)*rS6kSJK1*Z%`nl34p^{@Nb}y2s;{;LT4^|*J{CN2g8#lN@Fl7cIqQ8=!h_( zeGM6a!OD}fZYMw*fFxjDlc8S-qBy2tNOfA&sqrCa7~udLj7K9(Xd^|4ZkM+lBRE!r zVYxp(;W#&TF*_8lOVcj6YeQ;@r9GSsfG7;A9=^wAsDPsfKX!G;ZWvud5DtgbaEu5t zk5xe!RTyHfG_5M-Z;iJ8Pw>YuGf^NejPu~KYRXPA?jKB;7uE|06fK^Ca301#fG&WN@L2;fu@jm}M4{Klk3dm5 ziiO}MMzHL&LWhZu*-g4iDsES@Po9VVBvb8#0cAwJaJMy2Ne@B0UCR_w2PGXt7l=wE zM-^yNFrJPh7Az9LKro^Whz5ahEIWeIP4Kpq1at4YL0ZUqAT5Zw;BFm?UChuMN5)oJ zu=amOVRGE}Je&xjxX^L^Xox`)W4y@F2}c>GSe)sHX2PQ?EV9I7S}Y`pz#ydo7{ibT zXImioj%OIi0}|$OCO2IzwA2yThZ$l_#WVa`S4pjkdM>f=A9~EmX*NVh*%jus9kK&S zmO(j$yA5?hI>5nGG#=xR2?OI6N(U*0w9O*~>{tqaajzROvm*^CZ{RW$1@bzN(}B!k zNEzXFhB*yQB}ea4MyJl4&gVmvV(0{sFe3i1ZxKYoP@NfzyWkT_L_~fV3@o2WgCYL+ zJc;k>FT5LRCGkbqz?hH8i2N*Utug?PpstOsg)N%EBMWRu6Jr>t+2Q0VWF;^l8J+rb zI1Mz8W+6O`SOH=(_P6nTSexva5{x_m1s3l+o{z?eOK-d_ z0pxke@{s3kY$_MK!v#n&M0s9gc4mzDAVLDeIB;!UVeN-;ZfuUS4%i_i9Uxh&udog{ zywI=-_a0a!29DWj4YhEP+J+?%8$_rAV?MqCW(a~|^!<~Fl3yF@BT~aLMKI}1ctN0< zqeV6ZjZqhT&t71%Fh>FH&@|Q>$j25vq|swB7a}ejj=+%n;kYE8P5Q0dBz^E0jt4Tl zi_ym^7r9`s)h!HW;P+wV2tAkm+AI?a69a{@7988bP%@4;sPZ|%z$FIb=m-G zuccJ%d?E`m%!?8OSBHdnOy-f3r)VZ2EDqUz$6~N?cPGe)`tP_fkB^Q_?#_0i|wdP~c-!R37 z3)`HvBxd>SjH}>T4@)pPgdPx2q%*h(RSmGUbAx}iGGhv9H?dg&P^knxC+gXKcRqic8>v!y87Om!F)VM4*AkKzEET(^8JXYTCp;AAy|j)Ch4 z0o4D{{eft4&^k{?FNBoWW(Y#+XjQxi&ma z4z$T;9+~$-#?pS4{&bu3F13MA>^Jx~EV?$A%~TeCzcoj= z$Z(0{vc1E`h%6({UOSx7GKz-0Bk3;QFQ!GlX(ybqMdPxJBY-G#yhpWnrx{q&2>v|M zO($5Em-B$2nK-A%VA&p(BGl8GgpNPq@!r=Sq5T4Ukgwo`5} z%Uj-fwVO#dpp|^+QU^#+x=oO#PY$o4Xx^A7CBVh+!E}s$F2xIT$}{C~VJc%;@Th^Uq<8aca!89gjAX@;iVt+D&4Zs3G0WxOw zhsrKIK%`Af1}Yu-xATWUX)mn9Q*n469M0;(G|A_ym8ly@qnPcsQ%jhTz#YjFOwQo6 z49oc=!N8M1Awvw9B=C6&HUaE#nb!<|7H6dKbWF@y8$rt2IRcoArIKxHgA;$2`zNr5 z6ob?oGT=B3WYW)kk5QxMn34f-hfL=G`6rhz##vrCC&?!LFtx^M7n65!h80g>fCB+H z0`u^8DnsauWlC}dpA{2yma5?1yy^_4M~nRmna}nqPyE1?qX8_=p(9s`9Hq^r^VCO3 zxuEPw?m$8SJw^%zwGIHu+zk8-L#jMd{DR3F=UK#!>FF{eWyU(_Z=4#M{fBN&FKM9g;$m*qHROTFu#B0~O3e!xP8d zeDVB!eI-A=7F#*bLO@&3)nMDFl!bE(v5&F0v9D||(LNY$XU?H!0r0;27Nr&DD=alQ zxChQYNz8c5iM-gs&ko`|6*EWpWQ^L9mKO3FZPQ$MBH9-wj$l3&Pt9N};pifJ;yBL} zZaWMI(0n+x05E>#X?Q$0(@rNuK9wi4nXpt|dX;;%m>jorRw zA(Lfiv5=}-29~fFmC34b;f9qi--*mDCe1B_g|l0E9xT!x_8BEnt!x!hoZs_`LNMZM zH2;n>Xt-=uN9W$kT-yHu7ytk}_@9B{;84Kt85jU)0084+<6skDkWl0T(Tox)K$-!_ z=K<4u8U?m2kz4#;=j(oXyN<*Px p%btxtOad4jCirW@ literal 0 HcmV?d00001 From 73df774f6cb263303fab2af76024dfeef5f24d0e Mon Sep 17 00:00:00 2001 From: Ramakrishna Prabhu Date: Fri, 9 Jul 2021 14:13:10 -0500 Subject: [PATCH 03/26] review changes and doc --- cpp/src/io/orc/stripe_data.cu | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/cpp/src/io/orc/stripe_data.cu b/cpp/src/io/orc/stripe_data.cu index d6631ddf394..903f9475e2a 100644 --- a/cpp/src/io/orc/stripe_data.cu +++ b/cpp/src/io/orc/stripe_data.cu @@ -651,6 +651,7 @@ static const __device__ __constant__ uint8_t ClosestFixedBitsMap[65] = { * @param[in] vals buffer for output values (uint32_t, int32_t, uint64_t or int64_t) * @param[in] maxvals maximum number of values to decode * @param[in] t thread id + * @param[in] has_buffered_values If true, means there are already buffered values * * @return number of values decoded */ @@ -660,7 +661,7 @@ static __device__ uint32_t Integer_RLEv2(orc_bytestream_s* bs, volatile T* vals, uint32_t maxvals, int t, - bool has_buffer = false) + bool has_buffered_values = false) { uint32_t numvals, numruns; int r, tr; @@ -709,8 +710,9 @@ static __device__ uint32_t Integer_RLEv2(orc_bytestream_s* bs, } if ((numvals != 0) and (numvals + n > maxvals)) break; // case where there are buffered values and can't consume a whole chunk - // from decoded values, so work on buffered values and then start fresh in next iteration. - if ((numvals == 0) and (n > maxvals) and (has_buffer)) break; + // from decoded values, so skip adding any more to buffer, work on buffered values and then + // start fresh in next iteration with empty buffer. + if ((numvals == 0) and (n > maxvals) and (has_buffered_values)) break; pos += l; if (pos > maxpos) break; From 896fc8efafabc69c331510b297281684ba0dda1c Mon Sep 17 00:00:00 2001 From: Ramakrishna Prabhu Date: Sat, 17 Jul 2021 17:57:07 -0500 Subject: [PATCH 04/26] primary changes --- cpp/src/io/orc/orc_gpu.h | 6 + cpp/src/io/orc/reader_impl.cu | 149 +++++++++++++++++--- cpp/src/io/orc/reader_impl.hpp | 11 +- cpp/src/io/orc/stripe_data.cu | 7 +- cpp/src/io/orc/stripe_init.cu | 6 +- cpp/src/structs/structs_column_factories.cu | 6 +- python/cudf/cudf/tests/test_orc.py | 73 ++++------ 7 files changed, 185 insertions(+), 73 deletions(-) diff --git a/cpp/src/io/orc/orc_gpu.h b/cpp/src/io/orc/orc_gpu.h index fa91dd13755..a64d265917e 100644 --- a/cpp/src/io/orc/orc_gpu.h +++ b/cpp/src/io/orc/orc_gpu.h @@ -85,6 +85,11 @@ struct DictionaryEntry { */ constexpr int orc_decimal2float64_scale = 0x80; +struct ParentColumnData { + uint32_t* valid_map_base; + uint32_t null_count; +}; + /** * @brief Struct to describe per stripe's column information */ @@ -109,6 +114,7 @@ struct ColumnDesc { uint8_t dtype_len; // data type length (for types that can be mapped to different sizes) int32_t decimal_scale; // number of fractional decimal digits for decimal type int32_t ts_clock_rate; // output timestamp clock frequency (0=default, 1000=ms, 1000000000=ns) + ParentColumnData parent_data; // consists of parent column valid_map and null count }; /** diff --git a/cpp/src/io/orc/reader_impl.cu b/cpp/src/io/orc/reader_impl.cu index a5465090c2c..c6f3add5c14 100644 --- a/cpp/src/io/orc/reader_impl.cu +++ b/cpp/src/io/orc/reader_impl.cu @@ -28,6 +28,7 @@ #include #include +#include #include #include @@ -503,13 +504,14 @@ class aggregate_orc_metadata { selection[level][col_id].num_children = lvl_cols; } break; - case orc::STRUCT: + case orc::STRUCT: { + uint32_t lvl_cols = 0; for (const auto child_id : types[id].subtypes) { - num_lvl_child_columns += - add_column(selection, types, level, child_id, has_timestamp_column, has_list_column); + lvl_cols += add_column( + selection, types, level + 1, child_id, has_timestamp_column, has_list_column); } - selection[level][col_id].num_children = num_lvl_child_columns; - break; + selection[level][col_id].num_children = lvl_cols; + } break; default: break; } @@ -683,6 +685,70 @@ rmm::device_buffer reader::impl::decompress_stripe_data( return decomp_data; } +void update_null_mask(cudf::detail::hostdevice_2dvector& chunks, + std::vector& out_buffers, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr, + size_t level) +{ + auto num_columns = chunks.size().second; + if (level == 2) printf("RGSL :Coming to update mask and number of columns is %lu\n", num_columns); + + for (size_t j = 0; j < num_columns; ++j) { + if (chunks[0][j].parent_data.valid_map_base) { + auto parent_valid_map_base = chunks[0][j].parent_data.valid_map_base; + auto child_valid_map_base = out_buffers[j].null_mask(); + if (level == 2) + printf("RGSL : chunks[0][j].parent_data.null_count %u\n", + chunks[0][j].parent_data.null_count); + if (level == 2) printf("RGSL : null count is %u\n", chunks[0][j].null_count); + auto child_mask_len = chunks[0][j].column_num_rows - chunks[0][j].parent_data.null_count; + auto parent_mask_len = chunks[0][j].column_num_rows; + if (level == 2) + printf( + "RGSL : child mask len %u and parent mask len %u \n", child_mask_len, parent_mask_len); + if (child_valid_map_base) { + rmm::device_uvector dst_idx(child_mask_len, stream); + thrust::copy_if(rmm::exec_policy(stream), + thrust::make_counting_iterator(0), + thrust::make_counting_iterator(0) + parent_mask_len, + dst_idx.begin(), + [parent_valid_map_base] __device__(auto idx) { + auto val = bit_is_set(parent_valid_map_base, idx); + // printf("RGSL : Returning %u for idx %u\n", val, idx); + return val; + }); + + // thrust::for_each(rmm::exec_policy(stream), dst_idx.begin(), dst_idx.end(), [] + // __device__(auto idx) {printf("RGSL :idx has value %u \n", idx);}); + auto merged_null_mask = cudf::detail::create_null_mask( + parent_mask_len, mask_state::ALL_NULL, rmm::cuda_stream_view(stream), mr); + bitmask_type* merged_mask = static_cast(merged_null_mask.data()); + uint32_t* dst_idx_ptr = dst_idx.data(); + if (level == 2) printf("RGSL : dst_idx.size() is %lu \n", dst_idx.size()); + thrust::for_each( + rmm::exec_policy(stream), + thrust::make_counting_iterator(0), + thrust::make_counting_iterator(0) + dst_idx.size(), + [child_valid_map_base, dst_idx_ptr, merged_mask, level] __device__(auto idx) { + if (bit_is_set(child_valid_map_base, idx)) { + if (level == 2) + printf("RGSL: Setting bit dst_idx_ptr[idx] %u at idx %u\n", dst_idx_ptr[idx], idx); + cudf::set_bit(merged_mask, dst_idx_ptr[idx]); + }; + }); + + out_buffers[j]._null_mask = std::move(merged_null_mask); + } else { + if (level == 2) printf("RGSL: Coming to copy bitmask \n"); + auto mask_size = bitmask_allocation_size_bytes(parent_mask_len); + out_buffers[j]._null_mask = + rmm::device_buffer(parent_valid_map_base, mask_size, stream, mr); + } + } + } +} + void reader::impl::decode_stream_data(cudf::detail::hostdevice_2dvector& chunks, size_t num_dicts, size_t skip_rows, @@ -711,6 +777,26 @@ void reader::impl::decode_stream_data(cudf::detail::hostdevice_2dvector 0) { + chunks.device_to_host(stream, true); + // Update nullmask if parent was a struct + update_null_mask(chunks, out_buffers, stream, _mr, level); + // Update chunks with pointers to column data + for (size_t i = 0; i < num_stripes; ++i) { + for (size_t j = 0; j < num_columns; ++j) { + auto& chunk = chunks[i][j]; + chunk.valid_map_base = out_buffers[j].null_mask(); + if (level == 2) + printf("RGSL : number of values is %u and null count is %u \n", + chunk.column_num_rows, + chunk.parent_data.null_count); + } + } + + chunks.host_to_device(stream, true); + } + // Update the null map for child columns gpu::DecodeOrcColumnData(chunks.base_device_ptr(), global_dict.data(), row_groups, @@ -724,16 +810,18 @@ void reader::impl::decode_stream_data(cudf::detail::hostdevice_2dvector chunks, cudf::detail::host_2dspan row_groups, + std::vector& out_buffers, std::vector const& list_col, const int32_t level) { @@ -743,10 +831,12 @@ void reader::impl::aggregate_child_meta(cudf::detail::host_2dspan(out_buffers[parent_col_idx].null_count()); + auto parent_valid_map = static_cast(out_buffers[parent_col_idx].null_mask()); + auto num_rows = out_buffers[parent_col_idx].size; + for (uint32_t id = 0; id < p_col.num_children; id++) { + const auto child_col_idx = index + id; + if (type == type_id::STRUCT) { + parent_column_data[child_col_idx] = {parent_valid_map, parent_null_count}; + num_child_rows[child_col_idx] = num_rows; + } else { + parent_column_data[child_col_idx] = {nullptr, 0}; + } + } index += p_col.num_children; }); } @@ -864,7 +967,7 @@ column_buffer&& reader::impl::assemble_buffer(const int32_t orc_col_id, case type_id::STRUCT: for (auto const& col : _metadata->get_col_type(orc_col_id).subtypes) { - col_buffer.children.emplace_back(assemble_buffer(col, col_buffers, level)); + col_buffer.children.emplace_back(assemble_buffer(col, col_buffers, level + 1)); } break; @@ -881,15 +984,12 @@ void reader::impl::create_columns(std::vector>&& col_ std::vector& schema_info, rmm::cuda_stream_view stream) { - for (size_t i = 0; i < _selected_columns[0].size();) { + for (size_t i = 0; i < _selected_columns[0].size(); i++) { auto const& col_meta = _selected_columns[0][i]; schema_info.emplace_back(""); auto col_buffer = assemble_buffer(col_meta.id, col_buffers, 0); out_columns.emplace_back(make_column(col_buffer, &schema_info.back(), stream, _mr)); - - // Need to skip child columns of struct which are at the same level and have been processed - i += (col_buffers[0][i].type.id() == type_id::STRUCT) ? col_meta.num_children + 1 : 1; } } @@ -948,7 +1048,7 @@ table_with_metadata reader::impl::read(size_type skip_rows, auto& selected_columns = _selected_columns[level]; // Association between each ORC column and its cudf::column _col_meta.orc_col_map.emplace_back(_metadata->get_num_cols(), -1); - std::vector list_col; + std::vector nested_col; // Get a list of column data types std::vector column_types; @@ -976,19 +1076,17 @@ table_with_metadata reader::impl::read(size_type skip_rows, // Map each ORC column to its column _col_meta.orc_col_map[level][col.id] = column_types.size() - 1; - if (col_type == type_id::LIST) list_col.emplace_back(col); + if (col_type == type_id::LIST or col_type == type_id::STRUCT) nested_col.emplace_back(col); } // If no rows or stripes to read, return empty columns if (num_rows <= 0 || selected_stripes.empty()) { - for (size_t i = 0; i < _selected_columns[0].size();) { + for (size_t i = 0; i < _selected_columns[0].size(); i++) { auto const& col_meta = _selected_columns[0][i]; auto const schema = _metadata->get_schema(col_meta.id); schema_info.emplace_back(""); out_columns.push_back( std::move(create_empty_column(col_meta.id, schema_info.back(), stream))); - // Since struct children will be in the same level, have to skip them. - i += (schema.kind == orc::STRUCT) ? col_meta.num_children + 1 : 1; } break; } else { @@ -1097,9 +1195,14 @@ table_with_metadata reader::impl::read(size_type skip_rows, (level == 0) ? stripe_info->numberOfRows : _col_meta.num_child_rows_per_stripe[stripe_idx * num_columns + col_idx]; + // chunk.num_child_rows = chunk.num_rows; chunk.column_num_rows = (level == 0) ? num_rows : _col_meta.num_child_rows[col_idx]; - chunk.encoding_kind = stripe_footer->columns[selected_columns[col_idx].id].kind; - chunk.type_kind = _metadata->per_file_metadata[stripe_source_mapping.source_idx] + chunk.parent_data.valid_map_base = + (level == 0) ? nullptr : _col_meta.parent_column_data[col_idx].valid_map_base; + chunk.parent_data.null_count = + (level == 0) ? 0 : _col_meta.parent_column_data[col_idx].null_count; + chunk.encoding_kind = stripe_footer->columns[selected_columns[col_idx].id].kind; + chunk.type_kind = _metadata->per_file_metadata[stripe_source_mapping.source_idx] .ff.types[selected_columns[col_idx].id] .kind; auto const decimal_as_float64 = @@ -1216,13 +1319,13 @@ table_with_metadata reader::impl::read(size_type skip_rows, stream); // Extract information to process list child columns - if (list_col.size()) { + if (nested_col.size()) { row_groups.device_to_host(stream, true); - aggregate_child_meta(chunks, row_groups, list_col, level); + aggregate_child_meta(chunks, row_groups, out_buffers[level], nested_col, level); } // ORC stores number of elements at each row, so we need to generate offsets from that - if (list_col.size()) { + if (nested_col.size()) { std::vector buff_data; std::for_each( out_buffers[level].begin(), out_buffers[level].end(), [&buff_data](auto& out_buffer) { diff --git a/cpp/src/io/orc/reader_impl.hpp b/cpp/src/io/orc/reader_impl.hpp index 1769fb6f193..0392d6de44e 100644 --- a/cpp/src/io/orc/reader_impl.hpp +++ b/cpp/src/io/orc/reader_impl.hpp @@ -53,8 +53,14 @@ class aggregate_orc_metadata; */ struct reader_column_meta { std::vector> - orc_col_map; // Mapping between column id in orc to processing order. - std::vector num_child_rows; // number of rows in child columns + orc_col_map; // Mapping between column id in orc to processing order. + std::vector num_child_rows; // number of rows in child columns + struct ParentColumnData { + cudf::bitmask_type* valid_map_base = nullptr; + uint32_t null_count; + }; + std::vector + parent_column_data; // consists of parent column valid_map and null count std::vector child_start_row; // start row of child columns [stripe][column] std::vector num_child_rows_per_stripe; // number of rows of child columns [stripe][column] @@ -157,6 +163,7 @@ class reader::impl { */ void aggregate_child_meta(cudf::detail::host_2dspan chunks, cudf::detail::host_2dspan row_groups, + std::vector& out_buffers, std::vector const& list_col, const int32_t level); diff --git a/cpp/src/io/orc/stripe_data.cu b/cpp/src/io/orc/stripe_data.cu index 903f9475e2a..6c5132bc418 100644 --- a/cpp/src/io/orc/stripe_data.cu +++ b/cpp/src/io/orc/stripe_data.cu @@ -1152,7 +1152,7 @@ __global__ void __launch_bounds__(block_size) if (t == 0) s->chunk = chunks[chunk_id]; __syncthreads(); - const size_t max_num_rows = s->chunk.column_num_rows; + const size_t max_num_rows = s->chunk.column_num_rows - s->chunk.parent_data.null_count; if (is_nulldec) { uint32_t null_count = 0; @@ -1422,13 +1422,14 @@ __global__ void __launch_bounds__(block_size) } else { chunk_id = blockIdx.x; } + // Struct doesn't have any data in itself, so skip + const bool is_valid = s->chunk.type_kind != STRUCT; if (t == 0) { s->chunk = chunks[chunk_id]; s->num_child_rows = 0; + if (not is_valid) chunks[chunk_id].num_child_rows = chunks[chunk_id].num_rows; } __syncthreads(); - // Struct doesn't have any data in itself, so skip - const bool is_valid = s->chunk.type_kind != STRUCT; const size_t max_num_rows = s->chunk.column_num_rows; if (t == 0 and is_valid) { // If we have an index, seek to the initial run and update row positions diff --git a/cpp/src/io/orc/stripe_init.cu b/cpp/src/io/orc/stripe_init.cu index 317b7255718..3ee725c6e2f 100644 --- a/cpp/src/io/orc/stripe_init.cu +++ b/cpp/src/io/orc/stripe_init.cu @@ -456,8 +456,10 @@ extern "C" __global__ void __launch_bounds__(128, 8) ((uint32_t*)&row_groups[(s->rowgroup_start + i) * num_columns + blockIdx.x])[j] = ((volatile uint32_t*)&s->rowgroups[i])[j]; } - row_groups[(s->rowgroup_start + i) * num_columns + blockIdx.x].num_rows = num_rows; - row_groups[(s->rowgroup_start + i) * num_columns + blockIdx.x].start_row = start_row; + row_groups[(s->rowgroup_start + i) * num_columns + blockIdx.x].num_rows = num_rows; + // Updating in case of struct + row_groups[(s->rowgroup_start + i) * num_columns + blockIdx.x].num_child_rows = num_rows; + row_groups[(s->rowgroup_start + i) * num_columns + blockIdx.x].start_row = start_row; } __syncthreads(); if (t == 0) { s->rowgroup_start += num_rowgroups; } diff --git a/cpp/src/structs/structs_column_factories.cu b/cpp/src/structs/structs_column_factories.cu index 833ceab7518..4b06b87e1e7 100644 --- a/cpp/src/structs/structs_column_factories.cu +++ b/cpp/src/structs/structs_column_factories.cu @@ -40,7 +40,11 @@ std::unique_ptr make_structs_column( CUDF_EXPECTS(std::all_of(child_columns.begin(), child_columns.end(), - [&](auto const& child_col) { return num_rows == child_col->size(); }), + [&](auto const& child_col) { + // printf("RGSL : num_rows is %d and child_col->size() is %d \n", + // num_rows, child_col->size()); + return num_rows == child_col->size(); + }), "Child columns must have the same number of rows as the Struct column."); if (!null_mask.is_empty()) { diff --git a/python/cudf/cudf/tests/test_orc.py b/python/cudf/cudf/tests/test_orc.py index 1a785d28b48..ad980546284 100644 --- a/python/cudf/cudf/tests/test_orc.py +++ b/python/cudf/cudf/tests/test_orc.py @@ -895,15 +895,21 @@ def generate_list_struct_buff(size=28000): for x in range(size) ] lvl1_struct = [ - (np.random.randint(0, 3), np.random.randint(0, 3)) for x in range(size) + rd.choice([None, (np.random.randint(0, 3), np.random.randint(0, 3))]) + for x in range(size) ] lvl2_struct = [ - ( - rd.choice([None, np.random.randint(0, 3)]), - ( - rd.choice([None, np.random.randint(0, 3)]), - np.random.randint(0, 3), - ), + rd.choice( + [ + None, + ( + rd.choice([None, np.random.randint(0, 3)]), + ( + rd.choice([None, np.random.randint(0, 3)]), + np.random.randint(0, 3), + ), + ), + ] ) for x in range(size) ] @@ -949,51 +955,34 @@ def generate_list_struct_buff(size=28000): None, ["lvl3_list", "list_nests_struct", "lvl2_struct", "struct_nests_list"], ["lvl2_struct", "lvl1_struct"], + ["lvl2_struct"], ], ) -@pytest.mark.parametrize("num_rows", [0, 15, 1005, 10561, 28000]) +@pytest.mark.parametrize("num_rows", [0, 15, 1005, 10003, 28000]) @pytest.mark.parametrize("use_index", [True, False]) -@pytest.mark.parametrize("skip_rows", [0, 101, 1007, 27000]) def test_lists_struct_nests( - columns, num_rows, use_index, skip_rows, + columns, num_rows, use_index, ): - has_lists = ( - any("list" in col_name for col_name in columns) if columns else True + gdf = cudf.read_orc( + list_struct_buff, + columns=columns, + num_rows=num_rows, + use_index=use_index, ) - if has_lists and skip_rows > 0: - with pytest.raises( - RuntimeError, match="skip_rows is not supported by list column" - ): - cudf.read_orc( - list_struct_buff, - columns=columns, - num_rows=num_rows, - use_index=use_index, - skiprows=skip_rows, - ) - else: - gdf = cudf.read_orc( - list_struct_buff, - columns=columns, - num_rows=num_rows, - use_index=use_index, - skiprows=skip_rows, - ) - - pyarrow_tbl = pyarrow.orc.ORCFile(list_struct_buff).read() + pyarrow_tbl = pyarrow.orc.ORCFile(list_struct_buff).read() - pyarrow_tbl = ( - pyarrow_tbl[skip_rows : skip_rows + num_rows] - if columns is None - else pyarrow_tbl.select(columns)[skip_rows : skip_rows + num_rows] - ) + pyarrow_tbl = ( + pyarrow_tbl[:num_rows] + if columns is None + else pyarrow_tbl.select(columns)[:num_rows] + ) - if num_rows > 0: - assert_eq(True, pyarrow_tbl.equals(gdf.to_arrow())) - else: - assert_eq(pyarrow_tbl.to_pandas(), gdf) + if num_rows > 0: + assert_eq(True, pyarrow_tbl.equals(gdf.to_arrow())) + else: + assert_eq(pyarrow_tbl.to_pandas(), gdf) @pytest.mark.parametrize( From 6618164094b87883395cd730e8595f7e6de89a64 Mon Sep 17 00:00:00 2001 From: Ramakrishna Prabhu Date: Sun, 18 Jul 2021 14:10:05 -0500 Subject: [PATCH 05/26] primary changes --- cpp/src/io/orc/orc_gpu.h | 1 + cpp/src/io/orc/reader_impl.cu | 57 +++++++++++++++++++++--------- cpp/src/io/orc/stripe_data.cu | 14 +++++--- python/cudf/cudf/tests/test_orc.py | 12 ++----- 4 files changed, 54 insertions(+), 30 deletions(-) diff --git a/cpp/src/io/orc/orc_gpu.h b/cpp/src/io/orc/orc_gpu.h index a64d265917e..fac794de80e 100644 --- a/cpp/src/io/orc/orc_gpu.h +++ b/cpp/src/io/orc/orc_gpu.h @@ -266,6 +266,7 @@ void DecodeNullsAndStringDictionaries(ColumnDesc* chunks, uint32_t num_columns, uint32_t num_stripes, size_t first_row = 0, + size_t level = 0, rmm::cuda_stream_view stream = rmm::cuda_stream_default); /** diff --git a/cpp/src/io/orc/reader_impl.cu b/cpp/src/io/orc/reader_impl.cu index c6f3add5c14..9e3e3f15c69 100644 --- a/cpp/src/io/orc/reader_impl.cu +++ b/cpp/src/io/orc/reader_impl.cu @@ -692,19 +692,20 @@ void update_null_mask(cudf::detail::hostdevice_2dvector& chunks size_t level) { auto num_columns = chunks.size().second; - if (level == 2) printf("RGSL :Coming to update mask and number of columns is %lu\n", num_columns); + // if (level == 2) printf("RGSL :Coming to update mask and number of columns is %lu\n", + // num_columns); for (size_t j = 0; j < num_columns; ++j) { if (chunks[0][j].parent_data.valid_map_base) { auto parent_valid_map_base = chunks[0][j].parent_data.valid_map_base; auto child_valid_map_base = out_buffers[j].null_mask(); - if (level == 2) + if (level == 2 and false) printf("RGSL : chunks[0][j].parent_data.null_count %u\n", chunks[0][j].parent_data.null_count); - if (level == 2) printf("RGSL : null count is %u\n", chunks[0][j].null_count); + if (level == 2 and false) printf("RGSL : null count is %u\n", chunks[0][j].null_count); auto child_mask_len = chunks[0][j].column_num_rows - chunks[0][j].parent_data.null_count; auto parent_mask_len = chunks[0][j].column_num_rows; - if (level == 2) + if (level == 2 and false) printf( "RGSL : child mask len %u and parent mask len %u \n", child_mask_len, parent_mask_len); if (child_valid_map_base) { @@ -715,32 +716,38 @@ void update_null_mask(cudf::detail::hostdevice_2dvector& chunks dst_idx.begin(), [parent_valid_map_base] __device__(auto idx) { auto val = bit_is_set(parent_valid_map_base, idx); - // printf("RGSL : Returning %u for idx %u\n", val, idx); + // if(val) printf("RGSL : Returning %u for idx %u\n", val, idx); return val; }); // thrust::for_each(rmm::exec_policy(stream), dst_idx.begin(), dst_idx.end(), [] // __device__(auto idx) {printf("RGSL :idx has value %u \n", idx);}); +#if 0 + thrust::for_each(rmm::exec_policy(stream), + thrust::make_counting_iterator(0), + thrust::make_counting_iterator(0) + dst_idx.size(), [child_valid_map_base, level] + __device__(auto idx) {if(level == 2 and bit_is_set(child_valid_map_base, idx)) printf("RGSL :bit is set idx at %u \n", idx);}); +#endif auto merged_null_mask = cudf::detail::create_null_mask( parent_mask_len, mask_state::ALL_NULL, rmm::cuda_stream_view(stream), mr); bitmask_type* merged_mask = static_cast(merged_null_mask.data()); uint32_t* dst_idx_ptr = dst_idx.data(); - if (level == 2) printf("RGSL : dst_idx.size() is %lu \n", dst_idx.size()); + // if (level == 2) printf("RGSL : dst_idx.size() is %lu \n", dst_idx.size()); thrust::for_each( rmm::exec_policy(stream), thrust::make_counting_iterator(0), thrust::make_counting_iterator(0) + dst_idx.size(), [child_valid_map_base, dst_idx_ptr, merged_mask, level] __device__(auto idx) { if (bit_is_set(child_valid_map_base, idx)) { - if (level == 2) - printf("RGSL: Setting bit dst_idx_ptr[idx] %u at idx %u\n", dst_idx_ptr[idx], idx); + // if (level == 2) printf("RGSL: Setting bit dst_idx_ptr[idx] %u at idx %u\n", + // dst_idx_ptr[idx], idx); cudf::set_bit(merged_mask, dst_idx_ptr[idx]); }; }); out_buffers[j]._null_mask = std::move(merged_null_mask); } else { - if (level == 2) printf("RGSL: Coming to copy bitmask \n"); + // if (level == 2) printf("RGSL: Coming to copy bitmask \n"); auto mask_size = bitmask_allocation_size_bytes(parent_mask_len); out_buffers[j]._null_mask = rmm::device_buffer(parent_valid_map_base, mask_size, stream, mr); @@ -765,7 +772,9 @@ void reader::impl::decode_stream_data(cudf::detail::hostdevice_2dvector global_dict(num_dicts, stream); chunks.host_to_device(stream); - gpu::DecodeNullsAndStringDictionaries( - chunks.base_device_ptr(), global_dict.data(), num_columns, num_stripes, skip_rows, stream); + gpu::DecodeNullsAndStringDictionaries(chunks.base_device_ptr(), + global_dict.data(), + num_columns, + num_stripes, + skip_rows, + level, + stream); if (level > 0) { chunks.device_to_host(stream, true); @@ -787,10 +801,11 @@ void reader::impl::decode_stream_data(cudf::detail::hostdevice_2dvector(chunk.type_kind)); } } @@ -812,6 +827,7 @@ void reader::impl::decode_stream_data(cudf::detail::hostdevice_2dvectorget_num_cols(), -1); std::vector nested_col; @@ -1195,7 +1216,6 @@ table_with_metadata reader::impl::read(size_type skip_rows, (level == 0) ? stripe_info->numberOfRows : _col_meta.num_child_rows_per_stripe[stripe_idx * num_columns + col_idx]; - // chunk.num_child_rows = chunk.num_rows; chunk.column_num_rows = (level == 0) ? num_rows : _col_meta.num_child_rows[col_idx]; chunk.parent_data.valid_map_base = (level == 0) ? nullptr : _col_meta.parent_column_data[col_idx].valid_map_base; @@ -1205,6 +1225,9 @@ table_with_metadata reader::impl::read(size_type skip_rows, chunk.type_kind = _metadata->per_file_metadata[stripe_source_mapping.source_idx] .ff.types[selected_columns[col_idx].id] .kind; + chunk.num_child_rows = (chunk.type_kind != orc::STRUCT) ? 0 : chunk.num_rows; + // if(level ==1 ) printf("RGSL : chunk.type_kind is %u \n", + // static_cast(chunk.type_kind)); auto const decimal_as_float64 = should_convert_decimal_column_to_float(_decimal_cols_as_float, _metadata->per_file_metadata[0], @@ -1240,6 +1263,7 @@ table_with_metadata reader::impl::read(size_type skip_rows, if (stripe_data.size() != 0) { auto row_groups = cudf::detail::hostdevice_2dvector(num_rowgroups, num_columns, stream); + memset(row_groups.base_host_ptr(), 0, row_groups.memory_size()); if (level > 0 and row_groups.size().first) { cudf::host_span row_groups_span(row_groups.base_host_ptr(), num_rowgroups * num_columns); @@ -1305,6 +1329,7 @@ table_with_metadata reader::impl::read(size_type skip_rows, auto n_rows = (level == 0) ? num_rows : _col_meta.num_child_rows[i]; // For list column, offset column will be always size + 1 if (is_list_type) n_rows++; + // if(level == 2) printf("RGSL : Is nullable ? %u \n", is_nullable); out_buffers[level].emplace_back(column_types[i], n_rows, is_nullable, stream, _mr); } diff --git a/cpp/src/io/orc/stripe_data.cu b/cpp/src/io/orc/stripe_data.cu index 6c5132bc418..62d8f5cdfe9 100644 --- a/cpp/src/io/orc/stripe_data.cu +++ b/cpp/src/io/orc/stripe_data.cu @@ -1133,7 +1133,8 @@ __global__ void __launch_bounds__(block_size) DictionaryEntry* global_dictionary, uint32_t num_columns, uint32_t num_stripes, - size_t first_row) + size_t first_row, + size_t level) { __shared__ __align__(16) orcdec_state_s state_g; using warp_reduce = cub::WarpReduce; @@ -1422,13 +1423,15 @@ __global__ void __launch_bounds__(block_size) } else { chunk_id = blockIdx.x; } - // Struct doesn't have any data in itself, so skip - const bool is_valid = s->chunk.type_kind != STRUCT; if (t == 0) { s->chunk = chunks[chunk_id]; s->num_child_rows = 0; - if (not is_valid) chunks[chunk_id].num_child_rows = chunks[chunk_id].num_rows; } + // Struct doesn't have any data in itself, so skip + const bool is_valid = (s->chunk.type_kind != STRUCT); + // if(t == 0 and not is_valid){ + // chunks[chunk_id].num_child_rows = chunks[chunk_id].num_rows; + //} __syncthreads(); const size_t max_num_rows = s->chunk.column_num_rows; if (t == 0 and is_valid) { @@ -1834,12 +1837,13 @@ void __host__ DecodeNullsAndStringDictionaries(ColumnDesc* chunks, uint32_t num_columns, uint32_t num_stripes, size_t first_row, + size_t level, rmm::cuda_stream_view stream) { dim3 dim_block(block_size, 1); dim3 dim_grid(num_columns, num_stripes * 2); // 1024 threads per chunk gpuDecodeNullsAndStringDictionaries<<>>( - chunks, global_dictionary, num_columns, num_stripes, first_row); + chunks, global_dictionary, num_columns, num_stripes, first_row, level); } /** diff --git a/python/cudf/cudf/tests/test_orc.py b/python/cudf/cudf/tests/test_orc.py index ad980546284..169d9ce9b61 100644 --- a/python/cudf/cudf/tests/test_orc.py +++ b/python/cudf/cudf/tests/test_orc.py @@ -950,16 +950,10 @@ def generate_list_struct_buff(size=28000): @pytest.mark.parametrize( - "columns", - [ - None, - ["lvl3_list", "list_nests_struct", "lvl2_struct", "struct_nests_list"], - ["lvl2_struct", "lvl1_struct"], - ["lvl2_struct"], - ], + "columns", [None, ["lvl1_struct"], ["lvl3_list", "lvl2_struct"]], ) -@pytest.mark.parametrize("num_rows", [0, 15, 1005, 10003, 28000]) -@pytest.mark.parametrize("use_index", [True, False]) +@pytest.mark.parametrize("num_rows", [10003, 28000]) +@pytest.mark.parametrize("use_index", [False, True]) def test_lists_struct_nests( columns, num_rows, use_index, ): From 3494a793c69419f9e98a517d1bf44cd7f58956af Mon Sep 17 00:00:00 2001 From: Ramakrishna Prabhu Date: Sun, 18 Jul 2021 20:58:50 -0500 Subject: [PATCH 06/26] test cases working --- cpp/src/io/orc/reader_impl.cu | 19 +++++++++++++++++-- cpp/src/io/orc/reader_impl.hpp | 1 + cpp/src/io/orc/stripe_data.cu | 2 +- python/cudf/cudf/tests/test_orc.py | 11 ++++++++--- 4 files changed, 27 insertions(+), 6 deletions(-) diff --git a/cpp/src/io/orc/reader_impl.cu b/cpp/src/io/orc/reader_impl.cu index 9e3e3f15c69..c1cc1b57271 100644 --- a/cpp/src/io/orc/reader_impl.cu +++ b/cpp/src/io/orc/reader_impl.cu @@ -783,7 +783,8 @@ void reader::impl::decode_stream_data(cudf::detail::hostdevice_2dvector global_dict(num_dicts, stream); - chunks.host_to_device(stream); + // printf("RGSL : before DecodeNullsAndStringDictionaries %lu \n", level); + chunks.host_to_device(stream, true); gpu::DecodeNullsAndStringDictionaries(chunks.base_device_ptr(), global_dict.data(), num_columns, @@ -791,6 +792,7 @@ void reader::impl::decode_stream_data(cudf::detail::hostdevice_2dvector 0) { chunks.device_to_host(stream, true); @@ -801,7 +803,7 @@ void reader::impl::decode_stream_data(cudf::detail::hostdevice_2dvector(), std::move(out_metadata)}; + // RGSL remove this + if (_selected_columns[0].size() == 2) check = true; + // Select only stripes required (aka row groups) const auto selected_stripes = _metadata->select_stripes(stripes, skip_rows, num_rows); @@ -1063,6 +1075,7 @@ table_with_metadata reader::impl::read(size_type skip_rows, // in the same level since child column also have same number of rows, // list column children will be 1 level down compared to parent. for (size_t level = 0; level < _selected_columns.size(); level++) { + // printf("RGSL : Started with level %lu ------------------------------------\n", level); auto& selected_columns = _selected_columns[level]; // for(auto col : selected_columns){ // printf("RGSL : At level %lu selected columns are %u \n", level, col.id); @@ -1333,6 +1346,7 @@ table_with_metadata reader::impl::read(size_type skip_rows, out_buffers[level].emplace_back(column_types[i], n_rows, is_nullable, stream, _mr); } + // printf("RGSL : Before decode at level %lu \n", level); decode_stream_data(chunks, num_dict_entries, skip_rows, @@ -1343,6 +1357,7 @@ table_with_metadata reader::impl::read(size_type skip_rows, level, stream); + // printf("RGSL : After decode at level %lu \n", level); // Extract information to process list child columns if (nested_col.size()) { row_groups.device_to_host(stream, true); diff --git a/cpp/src/io/orc/reader_impl.hpp b/cpp/src/io/orc/reader_impl.hpp index 0392d6de44e..cd8597353f7 100644 --- a/cpp/src/io/orc/reader_impl.hpp +++ b/cpp/src/io/orc/reader_impl.hpp @@ -218,6 +218,7 @@ class reader::impl { std::vector _decimal_cols_as_float; data_type _timestamp_type{type_id::EMPTY}; reader_column_meta _col_meta; + bool check = false; }; } // namespace orc diff --git a/cpp/src/io/orc/stripe_data.cu b/cpp/src/io/orc/stripe_data.cu index 62d8f5cdfe9..3f92dce38aa 100644 --- a/cpp/src/io/orc/stripe_data.cu +++ b/cpp/src/io/orc/stripe_data.cu @@ -1428,11 +1428,11 @@ __global__ void __launch_bounds__(block_size) s->num_child_rows = 0; } // Struct doesn't have any data in itself, so skip - const bool is_valid = (s->chunk.type_kind != STRUCT); // if(t == 0 and not is_valid){ // chunks[chunk_id].num_child_rows = chunks[chunk_id].num_rows; //} __syncthreads(); + const bool is_valid = (s->chunk.type_kind != STRUCT); const size_t max_num_rows = s->chunk.column_num_rows; if (t == 0 and is_valid) { // If we have an index, seek to the initial run and update row positions diff --git a/python/cudf/cudf/tests/test_orc.py b/python/cudf/cudf/tests/test_orc.py index 169d9ce9b61..c16be73b733 100644 --- a/python/cudf/cudf/tests/test_orc.py +++ b/python/cudf/cudf/tests/test_orc.py @@ -950,10 +950,15 @@ def generate_list_struct_buff(size=28000): @pytest.mark.parametrize( - "columns", [None, ["lvl1_struct"], ["lvl3_list", "lvl2_struct"]], + "columns", + [ + None, + ["lvl3_list", "list_nests_struct", "lvl2_struct", "struct_nests_list"], + ["lvl2_struct", "lvl1_struct"], + ], ) -@pytest.mark.parametrize("num_rows", [10003, 28000]) -@pytest.mark.parametrize("use_index", [False, True]) +@pytest.mark.parametrize("num_rows", [0, 15, 1005, 10561, 28000]) +@pytest.mark.parametrize("use_index", [True, False]) def test_lists_struct_nests( columns, num_rows, use_index, ): From 97ad267b25bb147d2e08855cd102cbb304bed6e9 Mon Sep 17 00:00:00 2001 From: Ramakrishna Prabhu Date: Mon, 19 Jul 2021 18:42:28 -0500 Subject: [PATCH 07/26] clean-up --- cpp/src/io/orc/reader_impl.cu | 163 ++++++++------------ cpp/src/io/orc/reader_impl.hpp | 3 +- cpp/src/structs/structs_column_factories.cu | 6 +- python/cudf/cudf/tests/test_orc.py | 23 +-- 4 files changed, 79 insertions(+), 116 deletions(-) diff --git a/cpp/src/io/orc/reader_impl.cu b/cpp/src/io/orc/reader_impl.cu index c1cc1b57271..9401f4fa304 100644 --- a/cpp/src/io/orc/reader_impl.cu +++ b/cpp/src/io/orc/reader_impl.cu @@ -465,14 +465,14 @@ class aggregate_orc_metadata { /** * @brief Adds column as per the request and saves metadata about children. - * Struct children are in the same level as struct, only list column - * children are pushed to next level. + * Children of a column will be in the next level. * * @param selection A vector that saves list of columns as per levels of nesting. * @param types A vector of schema types of columns. * @param level current level of nesting. * @param id current column id that needs to be added. * @param has_timestamp_column True if timestamp column present and false otherwise. + * @param has_nested_column True if any of the selected column is a nested type. * * @return returns number of child columns at same level in case of struct and next level in case * of list @@ -482,7 +482,7 @@ class aggregate_orc_metadata { const size_t level, const uint32_t id, bool& has_timestamp_column, - bool& has_list_column) + bool& has_nested_column) { uint32_t num_lvl_child_columns = 0; if (level == selection.size()) { selection.emplace_back(); } @@ -494,21 +494,23 @@ class aggregate_orc_metadata { case orc::LIST: { uint32_t lvl_cols = 0; if (not types[id].subtypes.empty()) { - has_list_column = true; + has_nested_column = true; // Since list column needs to be processed before its child can be processed, // child column is being added to next level - lvl_cols = - add_column(selection, types, level + 1, id + 1, has_timestamp_column, has_list_column); + lvl_cols = add_column( + selection, types, level + 1, id + 1, has_timestamp_column, has_nested_column); } - // The list child column may be a struct in which case lvl_cols will be > 1 selection[level][col_id].num_children = lvl_cols; } break; case orc::STRUCT: { + has_nested_column = true; uint32_t lvl_cols = 0; for (const auto child_id : types[id].subtypes) { + // Since struct column needs to be processed before its child can be processed, + // child column is being added to next level lvl_cols += add_column( - selection, types, level + 1, child_id, has_timestamp_column, has_list_column); + selection, types, level + 1, child_id, has_timestamp_column, has_nested_column); } selection[level][col_id].num_children = lvl_cols; } break; @@ -524,11 +526,12 @@ class aggregate_orc_metadata { * * @param use_names List of column names to select * @param has_timestamp_column True if timestamp column present and false otherwise + * @param has_nested_column True if any of the selected column is a nested type. * * @return Vector of list of ORC column meta-data */ std::vector> select_columns( - std::vector const& use_names, bool& has_timestamp_column, bool& has_list_column) + std::vector const& use_names, bool& has_timestamp_column, bool& has_nested_column) { auto const& pfm = per_file_metadata[0]; std::vector> selection; @@ -545,7 +548,7 @@ class aggregate_orc_metadata { auto col_id = pfm.ff.types[0].subtypes[index]; if (pfm.get_column_name(col_id) == use_name) { name_found = true; - add_column(selection, pfm.ff.types, 0, col_id, has_timestamp_column, has_list_column); + add_column(selection, pfm.ff.types, 0, col_id, has_timestamp_column, has_nested_column); // Should start with next index index = i + 1; break; @@ -555,7 +558,7 @@ class aggregate_orc_metadata { } } else { for (auto const& col_id : pfm.ff.types[0].subtypes) { - add_column(selection, pfm.ff.types, 0, col_id, has_timestamp_column, has_list_column); + add_column(selection, pfm.ff.types, 0, col_id, has_timestamp_column, has_nested_column); } } @@ -685,29 +688,33 @@ rmm::device_buffer reader::impl::decompress_stripe_data( return decomp_data; } +/** + * @brief Updates null mask of a column whose parent is a struct column. + * In case of structure, if the structure has null element that would have be + * skipped while writing child column in ORC, so we need to insert the missing null + * elements in child column. + * + * @param chunks Vector of list of column chunk descriptors + * @param out_buffers Output columns' device buffers + * @param stream CUDA stream used for device memory operations and kernel launches. + * @param mr Device memory resource to use for device memory allocation + */ void update_null_mask(cudf::detail::hostdevice_2dvector& chunks, std::vector& out_buffers, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr, - size_t level) + rmm::mr::device_memory_resource* mr) { - auto num_columns = chunks.size().second; - // if (level == 2) printf("RGSL :Coming to update mask and number of columns is %lu\n", - // num_columns); + const auto num_stripes = chunks.size().first; + const auto num_columns = chunks.size().second; + bool is_mask_updated = false; for (size_t j = 0; j < num_columns; ++j) { if (chunks[0][j].parent_data.valid_map_base) { + if (not is_mask_updated) chunks.device_to_host(stream); auto parent_valid_map_base = chunks[0][j].parent_data.valid_map_base; auto child_valid_map_base = out_buffers[j].null_mask(); - if (level == 2 and false) - printf("RGSL : chunks[0][j].parent_data.null_count %u\n", - chunks[0][j].parent_data.null_count); - if (level == 2 and false) printf("RGSL : null count is %u\n", chunks[0][j].null_count); auto child_mask_len = chunks[0][j].column_num_rows - chunks[0][j].parent_data.null_count; auto parent_mask_len = chunks[0][j].column_num_rows; - if (level == 2 and false) - printf( - "RGSL : child mask len %u and parent mask len %u \n", child_mask_len, parent_mask_len); if (child_valid_map_base) { rmm::device_uvector dst_idx(child_mask_len, stream); thrust::copy_if(rmm::exec_policy(stream), @@ -716,43 +723,40 @@ void update_null_mask(cudf::detail::hostdevice_2dvector& chunks dst_idx.begin(), [parent_valid_map_base] __device__(auto idx) { auto val = bit_is_set(parent_valid_map_base, idx); - // if(val) printf("RGSL : Returning %u for idx %u\n", val, idx); return val; }); - // thrust::for_each(rmm::exec_policy(stream), dst_idx.begin(), dst_idx.end(), [] - // __device__(auto idx) {printf("RGSL :idx has value %u \n", idx);}); -#if 0 - thrust::for_each(rmm::exec_policy(stream), - thrust::make_counting_iterator(0), - thrust::make_counting_iterator(0) + dst_idx.size(), [child_valid_map_base, level] - __device__(auto idx) {if(level == 2 and bit_is_set(child_valid_map_base, idx)) printf("RGSL :bit is set idx at %u \n", idx);}); -#endif auto merged_null_mask = cudf::detail::create_null_mask( parent_mask_len, mask_state::ALL_NULL, rmm::cuda_stream_view(stream), mr); bitmask_type* merged_mask = static_cast(merged_null_mask.data()); uint32_t* dst_idx_ptr = dst_idx.data(); - // if (level == 2) printf("RGSL : dst_idx.size() is %lu \n", dst_idx.size()); - thrust::for_each( - rmm::exec_policy(stream), - thrust::make_counting_iterator(0), - thrust::make_counting_iterator(0) + dst_idx.size(), - [child_valid_map_base, dst_idx_ptr, merged_mask, level] __device__(auto idx) { - if (bit_is_set(child_valid_map_base, idx)) { - // if (level == 2) printf("RGSL: Setting bit dst_idx_ptr[idx] %u at idx %u\n", - // dst_idx_ptr[idx], idx); - cudf::set_bit(merged_mask, dst_idx_ptr[idx]); - }; - }); + thrust::for_each(rmm::exec_policy(stream), + thrust::make_counting_iterator(0), + thrust::make_counting_iterator(0) + dst_idx.size(), + [child_valid_map_base, dst_idx_ptr, merged_mask] __device__(auto idx) { + if (bit_is_set(child_valid_map_base, idx)) { + cudf::set_bit(merged_mask, dst_idx_ptr[idx]); + }; + }); out_buffers[j]._null_mask = std::move(merged_null_mask); } else { - // if (level == 2) printf("RGSL: Coming to copy bitmask \n"); auto mask_size = bitmask_allocation_size_bytes(parent_mask_len); out_buffers[j]._null_mask = rmm::device_buffer(parent_valid_map_base, mask_size, stream, mr); } + if (not is_mask_updated) is_mask_updated = true; + } + } + if (is_mask_updated) { + // Update chunks with pointers to column data which might have been changed. + for (size_t i = 0; i < num_stripes; ++i) { + for (size_t j = 0; j < num_columns; ++j) { + auto& chunk = chunks[i][j]; + chunk.valid_map_base = out_buffers[j].null_mask(); + } } + chunks.host_to_device(stream); } } @@ -772,9 +776,7 @@ void reader::impl::decode_stream_data(cudf::detail::hostdevice_2dvector global_dict(num_dicts, stream); - // printf("RGSL : before DecodeNullsAndStringDictionaries %lu \n", level); chunks.host_to_device(stream, true); gpu::DecodeNullsAndStringDictionaries(chunks.base_device_ptr(), global_dict.data(), @@ -792,28 +793,12 @@ void reader::impl::decode_stream_data(cudf::detail::hostdevice_2dvector 0) { - chunks.device_to_host(stream, true); - // Update nullmask if parent was a struct - update_null_mask(chunks, out_buffers, stream, _mr, level); - // Update chunks with pointers to column data - for (size_t i = 0; i < num_stripes; ++i) { - for (size_t j = 0; j < num_columns; ++j) { - auto& chunk = chunks[i][j]; - chunk.valid_map_base = out_buffers[j].null_mask(); - if (level == 1 and false) - printf("RGSL : number of values is %u and null count is %u and chunk kind is %u\n", - chunk.column_num_rows, - chunk.parent_data.null_count, - static_cast(chunk.type_kind)); - } - } - - chunks.host_to_device(stream, true); + // Update nullmasks if parent was a struct and had null mask + update_null_mask(chunks, out_buffers, stream, _mr); } - // printf("RGSL : after update_null_mask %lu \n", level); + // Update the null map for child columns gpu::DecodeOrcColumnData(chunks.base_device_ptr(), global_dict.data(), @@ -827,11 +812,9 @@ void reader::impl::decode_stream_data(cudf::detail::hostdevice_2dvector(out_buffers[parent_col_idx].null_count()); auto parent_valid_map = static_cast(out_buffers[parent_col_idx].null_mask()); @@ -988,10 +967,6 @@ column_buffer&& reader::impl::assemble_buffer(const int32_t orc_col_id, col_buffer.name = _metadata->get_column_name(0, orc_col_id); switch (col_buffer.type.id()) { case type_id::LIST: - col_buffer.children.emplace_back( - assemble_buffer(_metadata->get_col_type(orc_col_id).subtypes[0], col_buffers, level + 1)); - break; - case type_id::STRUCT: for (auto const& col : _metadata->get_col_type(orc_col_id).subtypes) { col_buffer.children.emplace_back(assemble_buffer(col, col_buffers, level + 1)); @@ -1030,7 +1005,7 @@ reader::impl::impl(std::vector>&& sources, // Select only columns required by the options _selected_columns = - _metadata->select_columns(options.get_columns(), _has_timestamp_column, _has_list_column); + _metadata->select_columns(options.get_columns(), _has_timestamp_column, _has_nested_column); // Override output timestamp resolution if requested if (options.get_timestamp_type().id() != type_id::EMPTY) { @@ -1052,8 +1027,8 @@ table_with_metadata reader::impl::read(size_type skip_rows, const std::vector>& stripes, rmm::cuda_stream_view stream) { - CUDF_EXPECTS(skip_rows == 0 or (not _has_list_column), - "skip_rows is not supported by list column"); + CUDF_EXPECTS(skip_rows == 0 or (not _has_nested_column), + "skip_rows is not supported by nested columns"); std::vector> out_columns; // buffer and stripe data are stored as per nesting level @@ -1065,21 +1040,13 @@ table_with_metadata reader::impl::read(size_type skip_rows, // There are no columns in the table if (_selected_columns.size() == 0) return {std::make_unique(), std::move(out_metadata)}; - // RGSL remove this - if (_selected_columns[0].size() == 2) check = true; - // Select only stripes required (aka row groups) const auto selected_stripes = _metadata->select_stripes(stripes, skip_rows, num_rows); - // Iterates through levels of nested columns, struct columns and its children will be - // in the same level since child column also have same number of rows, - // list column children will be 1 level down compared to parent. + // Iterates through levels of nested columns, child column will be one level down + // compared to parent column. for (size_t level = 0; level < _selected_columns.size(); level++) { - // printf("RGSL : Started with level %lu ------------------------------------\n", level); auto& selected_columns = _selected_columns[level]; - // for(auto col : selected_columns){ - // printf("RGSL : At level %lu selected columns are %u \n", level, col.id); - //} // Association between each ORC column and its cudf::column _col_meta.orc_col_map.emplace_back(_metadata->get_num_cols(), -1); std::vector nested_col; @@ -1238,9 +1205,9 @@ table_with_metadata reader::impl::read(size_type skip_rows, chunk.type_kind = _metadata->per_file_metadata[stripe_source_mapping.source_idx] .ff.types[selected_columns[col_idx].id] .kind; + // num_child_rows for a struct column will be same, for other nested types it will be + // calculated. chunk.num_child_rows = (chunk.type_kind != orc::STRUCT) ? 0 : chunk.num_rows; - // if(level ==1 ) printf("RGSL : chunk.type_kind is %u \n", - // static_cast(chunk.type_kind)); auto const decimal_as_float64 = should_convert_decimal_column_to_float(_decimal_cols_as_float, _metadata->per_file_metadata[0], @@ -1276,7 +1243,6 @@ table_with_metadata reader::impl::read(size_type skip_rows, if (stripe_data.size() != 0) { auto row_groups = cudf::detail::hostdevice_2dvector(num_rowgroups, num_columns, stream); - memset(row_groups.base_host_ptr(), 0, row_groups.memory_size()); if (level > 0 and row_groups.size().first) { cudf::host_span row_groups_span(row_groups.base_host_ptr(), num_rowgroups * num_columns); @@ -1342,11 +1308,9 @@ table_with_metadata reader::impl::read(size_type skip_rows, auto n_rows = (level == 0) ? num_rows : _col_meta.num_child_rows[i]; // For list column, offset column will be always size + 1 if (is_list_type) n_rows++; - // if(level == 2) printf("RGSL : Is nullable ? %u \n", is_nullable); out_buffers[level].emplace_back(column_types[i], n_rows, is_nullable, stream, _mr); } - // printf("RGSL : Before decode at level %lu \n", level); decode_stream_data(chunks, num_dict_entries, skip_rows, @@ -1357,7 +1321,6 @@ table_with_metadata reader::impl::read(size_type skip_rows, level, stream); - // printf("RGSL : After decode at level %lu \n", level); // Extract information to process list child columns if (nested_col.size()) { row_groups.device_to_host(stream, true); diff --git a/cpp/src/io/orc/reader_impl.hpp b/cpp/src/io/orc/reader_impl.hpp index cd8597353f7..55dcdb40c2c 100644 --- a/cpp/src/io/orc/reader_impl.hpp +++ b/cpp/src/io/orc/reader_impl.hpp @@ -214,11 +214,10 @@ class reader::impl { bool _use_index = true; bool _use_np_dtypes = true; bool _has_timestamp_column = false; - bool _has_list_column = false; + bool _has_nested_column = false; std::vector _decimal_cols_as_float; data_type _timestamp_type{type_id::EMPTY}; reader_column_meta _col_meta; - bool check = false; }; } // namespace orc diff --git a/cpp/src/structs/structs_column_factories.cu b/cpp/src/structs/structs_column_factories.cu index 4b06b87e1e7..833ceab7518 100644 --- a/cpp/src/structs/structs_column_factories.cu +++ b/cpp/src/structs/structs_column_factories.cu @@ -40,11 +40,7 @@ std::unique_ptr make_structs_column( CUDF_EXPECTS(std::all_of(child_columns.begin(), child_columns.end(), - [&](auto const& child_col) { - // printf("RGSL : num_rows is %d and child_col->size() is %d \n", - // num_rows, child_col->size()); - return num_rows == child_col->size(); - }), + [&](auto const& child_col) { return num_rows == child_col->size(); }), "Child columns must have the same number of rows as the Struct column."); if (!null_mask.is_empty()) { diff --git a/python/cudf/cudf/tests/test_orc.py b/python/cudf/cudf/tests/test_orc.py index c16be73b733..91e108b11ea 100644 --- a/python/cudf/cudf/tests/test_orc.py +++ b/python/cudf/cudf/tests/test_orc.py @@ -844,8 +844,8 @@ def test_orc_string_stream_offset_issue(): def generate_list_struct_buff(size=28000): - rd = random.Random(0) - np.random.seed(seed=0) + rd = random.Random(1) + np.random.seed(seed=1) buff = BytesIO() @@ -875,16 +875,21 @@ def generate_list_struct_buff(size=28000): schema = po.Struct(**schema) lvl3_list = [ - [ + rd.choice( [ + None, [ - rd.choice([None, np.random.randint(1, 3)]) - for z in range(np.random.randint(1, 3)) - ] - for z in range(np.random.randint(0, 3)) + [ + [ + rd.choice([None, np.random.randint(1, 3)]) + for z in range(np.random.randint(1, 3)) + ] + for z in range(np.random.randint(0, 3)) + ] + for y in range(np.random.randint(0, 3)) + ], ] - for y in range(np.random.randint(0, 3)) - ] + ) for x in range(size) ] lvl1_list = [ From 1bd9607bd83f5449fc11dc9837b25ae2fa059baa Mon Sep 17 00:00:00 2001 From: Ramakrishna Prabhu Date: Mon, 19 Jul 2021 18:57:31 -0500 Subject: [PATCH 08/26] added test case and docs --- cpp/src/io/orc/orc_gpu.h | 3 +++ cpp/src/io/orc/reader_impl.hpp | 4 +++- python/cudf/cudf/tests/test_orc.py | 10 ++++++++++ 3 files changed, 16 insertions(+), 1 deletion(-) diff --git a/cpp/src/io/orc/orc_gpu.h b/cpp/src/io/orc/orc_gpu.h index fac794de80e..ab46da53fd2 100644 --- a/cpp/src/io/orc/orc_gpu.h +++ b/cpp/src/io/orc/orc_gpu.h @@ -85,6 +85,9 @@ struct DictionaryEntry { */ constexpr int orc_decimal2float64_scale = 0x80; +/** + * Holds parent column mask and null count + */ struct ParentColumnData { uint32_t* valid_map_base; uint32_t null_count; diff --git a/cpp/src/io/orc/reader_impl.hpp b/cpp/src/io/orc/reader_impl.hpp index 55dcdb40c2c..7805e8e49cf 100644 --- a/cpp/src/io/orc/reader_impl.hpp +++ b/cpp/src/io/orc/reader_impl.hpp @@ -55,12 +55,14 @@ struct reader_column_meta { std::vector> orc_col_map; // Mapping between column id in orc to processing order. std::vector num_child_rows; // number of rows in child columns + struct ParentColumnData { cudf::bitmask_type* valid_map_base = nullptr; uint32_t null_count; }; std::vector - parent_column_data; // consists of parent column valid_map and null count + parent_column_data; // consists of parent column valid_map and null count + std::vector child_start_row; // start row of child columns [stripe][column] std::vector num_child_rows_per_stripe; // number of rows of child columns [stripe][column] diff --git a/python/cudf/cudf/tests/test_orc.py b/python/cudf/cudf/tests/test_orc.py index 91e108b11ea..48bbef71a8b 100644 --- a/python/cudf/cudf/tests/test_orc.py +++ b/python/cudf/cudf/tests/test_orc.py @@ -989,6 +989,16 @@ def test_lists_struct_nests( assert_eq(pyarrow_tbl.to_pandas(), gdf) +@pytest.mark.parametrize("columns", [None, ["lvl1_struct"], ["lvl1_list"]]) +def test_skip_rows_for_nested_types(columns): + with pytest.raises( + RuntimeError, match="skip_rows is not supported by nested column" + ): + cudf.read_orc( + list_struct_buff, columns=columns, use_index=True, skiprows=5, + ) + + @pytest.mark.parametrize( "data", [["_col0"], ["FakeName", "_col0", "TerriblyFakeColumnName"]] ) From 14994597b8422a92eda315ba9add7da2c71d7dd9 Mon Sep 17 00:00:00 2001 From: Ramakrishna Prabhu Date: Wed, 21 Jul 2021 15:38:36 -0500 Subject: [PATCH 09/26] Fixes for pyspark generated orc file reading --- cpp/src/io/orc/orc_gpu.h | 1 - cpp/src/io/orc/reader_impl.cu | 30 ++++++++++++++---------------- cpp/src/io/orc/reader_impl.hpp | 3 ++- cpp/src/io/orc/stripe_data.cu | 17 ++++++++--------- python/cudf/cudf/tests/test_orc.py | 9 +++++++++ 5 files changed, 33 insertions(+), 27 deletions(-) diff --git a/cpp/src/io/orc/orc_gpu.h b/cpp/src/io/orc/orc_gpu.h index ab46da53fd2..dfa9c9625be 100644 --- a/cpp/src/io/orc/orc_gpu.h +++ b/cpp/src/io/orc/orc_gpu.h @@ -269,7 +269,6 @@ void DecodeNullsAndStringDictionaries(ColumnDesc* chunks, uint32_t num_columns, uint32_t num_stripes, size_t first_row = 0, - size_t level = 0, rmm::cuda_stream_view stream = rmm::cuda_stream_default); /** diff --git a/cpp/src/io/orc/reader_impl.cu b/cpp/src/io/orc/reader_impl.cu index 438f805509d..8ae42b3aedb 100644 --- a/cpp/src/io/orc/reader_impl.cu +++ b/cpp/src/io/orc/reader_impl.cu @@ -181,7 +181,8 @@ size_t gather_stream_info(const size_t stripe_index, bool use_index, size_t* num_dictionary_entries, cudf::detail::hostdevice_2dvector& chunks, - std::vector& stream_info) + std::vector& stream_info, + size_t level) { uint64_t src_offset = 0; uint64_t dst_offset = 0; @@ -194,7 +195,7 @@ size_t gather_stream_info(const size_t stripe_index, auto const column_id = *stream.column_id; auto col = orc2gdf[column_id]; - if (col == -1) { + if (col == -1 and level == 0) { // A struct-type column has no data itself, but rather child columns // for each of its fields. There is only a PRESENT stream, which // needs to be included for the reader. @@ -710,7 +711,7 @@ void update_null_mask(cudf::detail::hostdevice_2dvector& chunks for (size_t j = 0; j < num_columns; ++j) { if (chunks[0][j].parent_data.valid_map_base) { - if (not is_mask_updated) chunks.device_to_host(stream); + if (not is_mask_updated) chunks.device_to_host(stream, true); auto parent_valid_map_base = chunks[0][j].parent_data.valid_map_base; auto child_valid_map_base = out_buffers[j].null_mask(); auto child_mask_len = chunks[0][j].column_num_rows - chunks[0][j].parent_data.null_count; @@ -740,10 +741,11 @@ void update_null_mask(cudf::detail::hostdevice_2dvector& chunks }); out_buffers[j]._null_mask = std::move(merged_null_mask); + } else { - auto mask_size = bitmask_allocation_size_bytes(parent_mask_len); - out_buffers[j]._null_mask = - rmm::device_buffer(parent_valid_map_base, mask_size, stream, mr); + auto mask_size = bitmask_allocation_size_bytes(parent_mask_len); + out_buffers[j]._null_mask = std::move( + rmm::device_buffer(static_cast(parent_valid_map_base), mask_size, stream, mr)); } if (not is_mask_updated) is_mask_updated = true; } @@ -756,7 +758,7 @@ void update_null_mask(cudf::detail::hostdevice_2dvector& chunks chunk.valid_map_base = out_buffers[j].null_mask(); } } - chunks.host_to_device(stream); + chunks.host_to_device(stream, true); } } @@ -786,17 +788,12 @@ void reader::impl::decode_stream_data(cudf::detail::hostdevice_2dvector global_dict(num_dicts, stream); chunks.host_to_device(stream, true); - gpu::DecodeNullsAndStringDictionaries(chunks.base_device_ptr(), - global_dict.data(), - num_columns, - num_stripes, - skip_rows, - level, - stream); + gpu::DecodeNullsAndStringDictionaries( + chunks.base_device_ptr(), global_dict.data(), num_columns, num_stripes, skip_rows, stream); if (level > 0) { // Update nullmasks if parent was a struct and had null mask - update_null_mask(chunks, out_buffers, stream, _mr); + update_null_mask(chunks, out_buffers, strr); } // Update the null map for child columns @@ -1144,7 +1141,8 @@ table_with_metadata reader::impl::read(size_type skip_rows, use_index, &num_dict_entries, chunks, - stream_info); + stream_info, + level); CUDF_EXPECTS(total_data_size > 0, "Expected streams data within stripe"); diff --git a/cpp/src/io/orc/reader_impl.hpp b/cpp/src/io/orc/reader_impl.hpp index 7805e8e49cf..58c38c07f09 100644 --- a/cpp/src/io/orc/reader_impl.hpp +++ b/cpp/src/io/orc/reader_impl.hpp @@ -159,7 +159,8 @@ class reader::impl { * @brief Aggregate child metadata from parent column chunks. * * @param chunks Vector of list of parent column chunks. - * @param chunks Vector of list of parent column row groups. + * @param row_groups Vector of list of row index descriptors + * @param out_buffers Column buffers for columns. * @param list_col Vector of column metadata of list type parent columns. * @param level Current nesting level being processed. */ diff --git a/cpp/src/io/orc/stripe_data.cu b/cpp/src/io/orc/stripe_data.cu index 3f92dce38aa..a18866f57c6 100644 --- a/cpp/src/io/orc/stripe_data.cu +++ b/cpp/src/io/orc/stripe_data.cu @@ -1133,8 +1133,7 @@ __global__ void __launch_bounds__(block_size) DictionaryEntry* global_dictionary, uint32_t num_columns, uint32_t num_stripes, - size_t first_row, - size_t level) + size_t first_row) { __shared__ __align__(16) orcdec_state_s state_g; using warp_reduce = cub::WarpReduce; @@ -1187,6 +1186,7 @@ __global__ void __launch_bounds__(block_size) nrows = nrows_max; } __syncthreads(); + row_in = s->chunk.start_row + s->top.nulls_desc_row; if (row_in + nrows > first_row && row_in < first_row + max_num_rows && s->chunk.valid_map_base != NULL) { @@ -1335,7 +1335,10 @@ static __device__ void DecodeRowPositions(orcdec_state_s* s, s->top.data.cur_row + s->top.data.nrows < s->top.data.end_row) { uint32_t nrows = min(s->top.data.end_row - (s->top.data.cur_row + s->top.data.nrows), min((row_decoder_buffer_size - s->u.rowdec.nz_count) * 2, blockDim.x)); - if (s->chunk.strm_len[CI_PRESENT] > 0) { + // Even though s->chunk.strm_len is zero, there is possibility that there is null mask. + // This happens in a struct column with nulls which has child column which doesn't have any + // nulls. + if (s->chunk.strm_len[CI_PRESENT] > 0 or s->chunk.valid_map_base) { // We have a present stream uint32_t rmax = s->top.data.end_row - min((uint32_t)first_row, s->top.data.end_row); uint32_t r = (uint32_t)(s->top.data.cur_row + s->top.data.nrows + t - first_row); @@ -1427,11 +1430,8 @@ __global__ void __launch_bounds__(block_size) s->chunk = chunks[chunk_id]; s->num_child_rows = 0; } - // Struct doesn't have any data in itself, so skip - // if(t == 0 and not is_valid){ - // chunks[chunk_id].num_child_rows = chunks[chunk_id].num_rows; - //} __syncthreads(); + // Struct doesn't have any data in itself, so skip const bool is_valid = (s->chunk.type_kind != STRUCT); const size_t max_num_rows = s->chunk.column_num_rows; if (t == 0 and is_valid) { @@ -1837,13 +1837,12 @@ void __host__ DecodeNullsAndStringDictionaries(ColumnDesc* chunks, uint32_t num_columns, uint32_t num_stripes, size_t first_row, - size_t level, rmm::cuda_stream_view stream) { dim3 dim_block(block_size, 1); dim3 dim_grid(num_columns, num_stripes * 2); // 1024 threads per chunk gpuDecodeNullsAndStringDictionaries<<>>( - chunks, global_dictionary, num_columns, num_stripes, first_row, level); + chunks, global_dictionary, num_columns, num_stripes, first_row); } /** diff --git a/python/cudf/cudf/tests/test_orc.py b/python/cudf/cudf/tests/test_orc.py index 48bbef71a8b..82f61b34442 100644 --- a/python/cudf/cudf/tests/test_orc.py +++ b/python/cudf/cudf/tests/test_orc.py @@ -999,6 +999,15 @@ def test_skip_rows_for_nested_types(columns): ) +def test_pyspark_struct(datadir): + path = datadir / "TestOrcFile.testPySparkStruct.orc" + + pdf = pa.orc.ORCFile(path).read().to_pandas() + gdf = cudf.read_orc(path) + + assert_eq(pdf, gdf) + + @pytest.mark.parametrize( "data", [["_col0"], ["FakeName", "_col0", "TerriblyFakeColumnName"]] ) From 5546b32cefc7defcaf4d7a975f1439276fdcad30 Mon Sep 17 00:00:00 2001 From: Ramakrishna Prabhu Date: Wed, 21 Jul 2021 16:41:09 -0500 Subject: [PATCH 10/26] changes --- cpp/src/io/orc/orc.h | 6 +- cpp/src/io/orc/orc_gpu.h | 2 +- cpp/src/io/orc/reader_impl.cu | 101 +++++++++--------- .../orc/TestOrcFile.testPySparkStruct.orc | Bin 0 -> 425 bytes 4 files changed, 51 insertions(+), 58 deletions(-) create mode 100644 python/cudf/cudf/tests/data/orc/TestOrcFile.testPySparkStruct.orc diff --git a/cpp/src/io/orc/orc.h b/cpp/src/io/orc/orc.h index 474f404be0f..ec272568891 100644 --- a/cpp/src/io/orc/orc.h +++ b/cpp/src/io/orc/orc.h @@ -539,10 +539,7 @@ class OrcDecompressor { }; /** - * @brief Stores orc id for each column and its adjacent number of children - * in case of struct or number of children in case of list column. - * If list column has struct column, then all child columns of that struct are treated as child - * column of list. + * @brief Stores orc id for each column and number of children in that column. * * @code{.pseudo} * Consider following data where a struct has two members and a list column @@ -559,7 +556,6 @@ class OrcDecompressor { * */ struct orc_column_meta { - // orc_column_meta(uint32_t _id, uint32_t _num_children) : id(_id), num_children(_num_children){}; uint32_t id; // orc id for the column uint32_t num_children; // number of children at the same level of nesting in case of struct }; diff --git a/cpp/src/io/orc/orc_gpu.h b/cpp/src/io/orc/orc_gpu.h index dfa9c9625be..d95f709244d 100644 --- a/cpp/src/io/orc/orc_gpu.h +++ b/cpp/src/io/orc/orc_gpu.h @@ -86,7 +86,7 @@ struct DictionaryEntry { constexpr int orc_decimal2float64_scale = 0x80; /** - * Holds parent column mask and null count + * Struct to store parent column mask and null count */ struct ParentColumnData { uint32_t* valid_map_base; diff --git a/cpp/src/io/orc/reader_impl.cu b/cpp/src/io/orc/reader_impl.cu index 8ae42b3aedb..1ddc0b5a2ec 100644 --- a/cpp/src/io/orc/reader_impl.cu +++ b/cpp/src/io/orc/reader_impl.cu @@ -466,7 +466,7 @@ class aggregate_orc_metadata { /** * @brief Adds column as per the request and saves metadata about children. - * Children of a column will be in the next level. + * Children of a column will be added to the next level. * * @param selection A vector that saves list of columns as per levels of nesting. * @param types A vector of schema types of columns. @@ -474,52 +474,34 @@ class aggregate_orc_metadata { * @param id current column id that needs to be added. * @param has_timestamp_column True if timestamp column present and false otherwise. * @param has_nested_column True if any of the selected column is a nested type. - * - * @return returns number of child columns at same level in case of struct and next level in case - * of list */ - uint32_t add_column(std::vector>& selection, - std::vector const& types, - const size_t level, - const uint32_t id, - bool& has_timestamp_column, - bool& has_nested_column) + void add_column(std::vector>& selection, + std::vector const& types, + const size_t level, + const uint32_t id, + bool& has_timestamp_column, + bool& has_nested_column) { - uint32_t num_lvl_child_columns = 0; if (level == selection.size()) { selection.emplace_back(); } selection[level].push_back({id, 0}); const int col_id = selection[level].size() - 1; if (types[id].kind == orc::TIMESTAMP) { has_timestamp_column = true; } switch (types[id].kind) { - case orc::LIST: { - uint32_t lvl_cols = 0; - if (not types[id].subtypes.empty()) { - has_nested_column = true; - // Since list column needs to be processed before its child can be processed, - // child column is being added to next level - lvl_cols = add_column( - selection, types, level + 1, id + 1, has_timestamp_column, has_nested_column); - } - selection[level][col_id].num_children = lvl_cols; - } break; - + case orc::LIST: case orc::STRUCT: { has_nested_column = true; - uint32_t lvl_cols = 0; for (const auto child_id : types[id].subtypes) { // Since struct column needs to be processed before its child can be processed, // child column is being added to next level - lvl_cols += add_column( + add_column( selection, types, level + 1, child_id, has_timestamp_column, has_nested_column); } - selection[level][col_id].num_children = lvl_cols; + selection[level][col_id].num_children = types[id].subtypes.size(); } break; default: break; } - - return num_lvl_child_columns + 1; } /** @@ -690,10 +672,13 @@ rmm::device_buffer reader::impl::decompress_stripe_data( } /** - * @brief Updates null mask of a column whose parent is a struct column. - * In case of structure, if the structure has null element that would have be + * @brief Updates null mask of columns whose parent is a struct column. + * If struct column has null element, that row would be * skipped while writing child column in ORC, so we need to insert the missing null * elements in child column. + * There is another behavior from pyspark, where if the child column doesn't have any null + * elements, it will not have present stream, so in that case parent null mask need to be copied to + * child column. * * @param chunks Vector of list of column chunk descriptors * @param out_buffers Output columns' device buffers @@ -712,12 +697,15 @@ void update_null_mask(cudf::detail::hostdevice_2dvector& chunks for (size_t j = 0; j < num_columns; ++j) { if (chunks[0][j].parent_data.valid_map_base) { if (not is_mask_updated) chunks.device_to_host(stream, true); + auto parent_valid_map_base = chunks[0][j].parent_data.valid_map_base; auto child_valid_map_base = out_buffers[j].null_mask(); auto child_mask_len = chunks[0][j].column_num_rows - chunks[0][j].parent_data.null_count; auto parent_mask_len = chunks[0][j].column_num_rows; + if (child_valid_map_base) { rmm::device_uvector dst_idx(child_mask_len, stream); + // Copy indexes at which the parent has valid value. thrust::copy_if(rmm::exec_policy(stream), thrust::make_counting_iterator(0), thrust::make_counting_iterator(0) + parent_mask_len, @@ -731,6 +719,8 @@ void update_null_mask(cudf::detail::hostdevice_2dvector& chunks parent_mask_len, mask_state::ALL_NULL, rmm::cuda_stream_view(stream), mr); bitmask_type* merged_mask = static_cast(merged_null_mask.data()); uint32_t* dst_idx_ptr = dst_idx.data(); + // Copy the child valid bits to valid indexes, this will merge both child and parent null + // masks thrust::for_each(rmm::exec_policy(stream), thrust::make_counting_iterator(0), thrust::make_counting_iterator(0) + dst_idx.size(), @@ -743,6 +733,7 @@ void update_null_mask(cudf::detail::hostdevice_2dvector& chunks out_buffers[j]._null_mask = std::move(merged_null_mask); } else { + // Since child column doesn't have mask, copy parent null mask auto mask_size = bitmask_allocation_size_bytes(parent_mask_len); out_buffers[j]._null_mask = std::move( rmm::device_buffer(static_cast(parent_valid_map_base), mask_size, stream, mr)); @@ -792,8 +783,8 @@ void reader::impl::decode_stream_data(cudf::detail::hostdevice_2dvector 0) { - // Update nullmasks if parent was a struct and had null mask - update_null_mask(chunks, out_buffers, strr); + // Update nullmasks for children if parent was a struct and had null mask + update_null_mask(chunks, out_buffers, stream, _mr); } // Update the null map for child columns @@ -814,6 +805,7 @@ void reader::impl::decode_stream_data(cudf::detail::hostdevice_2dvector(out_buffers[parent_col_idx].null_count()); auto parent_valid_map = static_cast(out_buffers[parent_col_idx].null_mask()); auto num_rows = out_buffers[parent_col_idx].size; + for (uint32_t id = 0; id < p_col.num_children; id++) { const auto child_col_idx = index + id; if (type == type_id::STRUCT) { parent_column_data[child_col_idx] = {parent_valid_map, parent_null_count}; - num_child_rows[child_col_idx] = num_rows; + // Number of rows in child will remain same as parent in case of struct column + num_child_rows[child_col_idx] = num_rows; } else { parent_column_data[child_col_idx] = {nullptr, 0}; } @@ -986,13 +980,14 @@ void reader::impl::create_columns(std::vector>&& col_ std::vector& schema_info, rmm::cuda_stream_view stream) { - for (size_t i = 0; i < _selected_columns[0].size(); i++) { - auto const& col_meta = _selected_columns[0][i]; - schema_info.emplace_back(""); - - auto col_buffer = assemble_buffer(col_meta.id, col_buffers, 0); - out_columns.emplace_back(make_column(col_buffer, &schema_info.back(), stream, _mr)); - } + std::transform(_selected_columns[0].begin(), + _selected_columns[0].end(), + std::back_inserter(out_columns), + [&](auto const col_meta) { + schema_info.emplace_back(""); + auto col_buffer = assemble_buffer(col_meta.id, col_buffers, 0); + return make_column(col_buffer, &schema_info.back(), stream, _mr); + }); } reader::impl::impl(std::vector>&& sources, @@ -1082,13 +1077,13 @@ table_with_metadata reader::impl::read(size_type skip_rows, // If no rows or stripes to read, return empty columns if (num_rows <= 0 || selected_stripes.empty()) { - for (size_t i = 0; i < _selected_columns[0].size(); i++) { - auto const& col_meta = _selected_columns[0][i]; - auto const schema = _metadata->get_schema(col_meta.id); - schema_info.emplace_back(""); - out_columns.push_back( - std::move(create_empty_column(col_meta.id, schema_info.back(), stream))); - } + std::transform(_selected_columns[0].begin(), + _selected_columns[0].end(), + std::back_inserter(out_columns), + [&](auto const col_meta) { + schema_info.emplace_back(""); + return create_empty_column(col_meta.id, schema_info.back(), stream); + }); break; } else { // Get the total number of stripes across all input files. @@ -1322,7 +1317,7 @@ table_with_metadata reader::impl::read(size_type skip_rows, level, stream); - // Extract information to process list child columns + // Extract information to process nested child columns if (nested_col.size()) { row_groups.device_to_host(stream, true); aggregate_child_meta(chunks, row_groups, out_buffers[level], nested_col, level); @@ -1339,8 +1334,10 @@ table_with_metadata reader::impl::read(size_type skip_rows, } }); - auto const dev_buff_data = cudf::detail::make_device_uvector_async(buff_data, stream); - generate_offsets_for_list(dev_buff_data, stream); + if (buff_data.size()) { + auto const dev_buff_data = cudf::detail::make_device_uvector_async(buff_data, stream); + generate_offsets_for_list(dev_buff_data, stream); + } } } } diff --git a/python/cudf/cudf/tests/data/orc/TestOrcFile.testPySparkStruct.orc b/python/cudf/cudf/tests/data/orc/TestOrcFile.testPySparkStruct.orc new file mode 100644 index 0000000000000000000000000000000000000000..7748e901bceb8f86c1371468dc4444a205fa6b72 GIT binary patch literal 425 zcmZXPK}!Nr5XWcU+qe6Ct?P3kdTt4}C+)DA1%Y%3yaa(RcG**~iIfB?t96c!ojlnW z>CnBe)KREe@nQ|k%$xuG8Gf(V@00v|ds5Iw=@mC`bVmUEjPvqY1n6>msk{ontpy50TtO3B3xN1Q*>nud zj*t>fnJP0?A?w@(v;mDq2nkHcDQad4q@m`fGz155+F3z_739vvz($?l_cnhqX@ z(QKM#kJ0lmn~f%8Uz|oKQLE{+*F5bFPx~JaDSvn_7u_pJP`>A%eW`^E-1>RbdHv?R I*YE880^3bFDgXcg literal 0 HcmV?d00001 From bcc7dcbed89c3ce14015b2cc48f1939df2097e1b Mon Sep 17 00:00:00 2001 From: Ramakrishna Prabhu Date: Wed, 21 Jul 2021 17:07:19 -0500 Subject: [PATCH 11/26] changes --- cpp/src/io/orc/reader_impl.cu | 12 ++++++------ cpp/src/io/orc/stripe_data.cu | 2 +- 2 files changed, 7 insertions(+), 7 deletions(-) diff --git a/cpp/src/io/orc/reader_impl.cu b/cpp/src/io/orc/reader_impl.cu index 1ddc0b5a2ec..3438b9f9257 100644 --- a/cpp/src/io/orc/reader_impl.cu +++ b/cpp/src/io/orc/reader_impl.cu @@ -492,7 +492,7 @@ class aggregate_orc_metadata { case orc::STRUCT: { has_nested_column = true; for (const auto child_id : types[id].subtypes) { - // Since struct column needs to be processed before its child can be processed, + // Since nested column needs to be processed before its child can be processed, // child column is being added to next level add_column( selection, types, level + 1, child_id, has_timestamp_column, has_nested_column); @@ -677,8 +677,8 @@ rmm::device_buffer reader::impl::decompress_stripe_data( * skipped while writing child column in ORC, so we need to insert the missing null * elements in child column. * There is another behavior from pyspark, where if the child column doesn't have any null - * elements, it will not have present stream, so in that case parent null mask need to be copied to - * child column. + * elements, it will not have present stream, so in that case parent null mask need to be + * copied to child column. * * @param chunks Vector of list of column chunk descriptors * @param out_buffers Output columns' device buffers @@ -719,8 +719,8 @@ void update_null_mask(cudf::detail::hostdevice_2dvector& chunks parent_mask_len, mask_state::ALL_NULL, rmm::cuda_stream_view(stream), mr); bitmask_type* merged_mask = static_cast(merged_null_mask.data()); uint32_t* dst_idx_ptr = dst_idx.data(); - // Copy the child valid bits to valid indexes, this will merge both child and parent null - // masks + // Copy child valid bits from child column to valid indexes, this will merge both child and + // parent null masks thrust::for_each(rmm::exec_policy(stream), thrust::make_counting_iterator(0), thrust::make_counting_iterator(0) + dst_idx.size(), @@ -733,7 +733,7 @@ void update_null_mask(cudf::detail::hostdevice_2dvector& chunks out_buffers[j]._null_mask = std::move(merged_null_mask); } else { - // Since child column doesn't have mask, copy parent null mask + // Since child column doesn't have a mask, copy parent null mask auto mask_size = bitmask_allocation_size_bytes(parent_mask_len); out_buffers[j]._null_mask = std::move( rmm::device_buffer(static_cast(parent_valid_map_base), mask_size, stream, mr)); diff --git a/cpp/src/io/orc/stripe_data.cu b/cpp/src/io/orc/stripe_data.cu index a18866f57c6..75586055892 100644 --- a/cpp/src/io/orc/stripe_data.cu +++ b/cpp/src/io/orc/stripe_data.cu @@ -1432,7 +1432,7 @@ __global__ void __launch_bounds__(block_size) } __syncthreads(); // Struct doesn't have any data in itself, so skip - const bool is_valid = (s->chunk.type_kind != STRUCT); + const bool is_valid = s->chunk.type_kind != STRUCT; const size_t max_num_rows = s->chunk.column_num_rows; if (t == 0 and is_valid) { // If we have an index, seek to the initial run and update row positions From 470b7c0b9cd9d5401a828dc6ea9d479f1a662238 Mon Sep 17 00:00:00 2001 From: Ramakrishna Prabhu Date: Wed, 21 Jul 2021 19:47:17 -0500 Subject: [PATCH 12/26] review changes --- cpp/src/io/orc/orc.h | 8 +++++ cpp/src/io/orc/orc_gpu.h | 3 +- cpp/src/io/orc/reader_impl.cu | 58 +++++++++++++++--------------- cpp/src/io/orc/reader_impl.hpp | 6 +--- cpp/src/io/orc/stripe_data.cu | 4 +-- python/cudf/cudf/tests/test_orc.py | 3 +- 6 files changed, 45 insertions(+), 37 deletions(-) diff --git a/cpp/src/io/orc/orc.h b/cpp/src/io/orc/orc.h index ec272568891..71be53243dc 100644 --- a/cpp/src/io/orc/orc.h +++ b/cpp/src/io/orc/orc.h @@ -560,6 +560,14 @@ struct orc_column_meta { uint32_t num_children; // number of children at the same level of nesting in case of struct }; +/** + * @brief Stores column's validity map and null count + */ +struct column_validity_info { + uint32_t* valid_map_base; + uint32_t null_count; +}; + /** * @brief A helper class for ORC file metadata. Provides some additional * convenience methods for initializing and accessing metadata. diff --git a/cpp/src/io/orc/orc_gpu.h b/cpp/src/io/orc/orc_gpu.h index d95f709244d..45149d56853 100644 --- a/cpp/src/io/orc/orc_gpu.h +++ b/cpp/src/io/orc/orc_gpu.h @@ -24,6 +24,7 @@ #include #include #include +#include "orc.h" #include "orc_common.h" #include @@ -117,7 +118,7 @@ struct ColumnDesc { uint8_t dtype_len; // data type length (for types that can be mapped to different sizes) int32_t decimal_scale; // number of fractional decimal digits for decimal type int32_t ts_clock_rate; // output timestamp clock frequency (0=default, 1000=ms, 1000000000=ns) - ParentColumnData parent_data; // consists of parent column valid_map and null count + column_validity_info parent_validity_info; // consists of parent column valid_map and null count }; /** diff --git a/cpp/src/io/orc/reader_impl.cu b/cpp/src/io/orc/reader_impl.cu index 3438b9f9257..873825295ea 100644 --- a/cpp/src/io/orc/reader_impl.cu +++ b/cpp/src/io/orc/reader_impl.cu @@ -694,16 +694,20 @@ void update_null_mask(cudf::detail::hostdevice_2dvector& chunks const auto num_columns = chunks.size().second; bool is_mask_updated = false; - for (size_t j = 0; j < num_columns; ++j) { - if (chunks[0][j].parent_data.valid_map_base) { - if (not is_mask_updated) chunks.device_to_host(stream, true); + for (size_t col_idx = 0; col_idx < num_columns; ++col_idx) { + if (chunks[0][col_idx].parent_validity_info.valid_map_base != nullptr) { + if (not is_mask_updated) { + chunks.device_to_host(stream, true); + is_mask_updated = true; + } - auto parent_valid_map_base = chunks[0][j].parent_data.valid_map_base; - auto child_valid_map_base = out_buffers[j].null_mask(); - auto child_mask_len = chunks[0][j].column_num_rows - chunks[0][j].parent_data.null_count; - auto parent_mask_len = chunks[0][j].column_num_rows; + auto parent_valid_map_base = chunks[0][col_idx].parent_validity_info.valid_map_base; + auto child_valid_map_base = out_buffers[col_idx].null_mask(); + auto child_mask_len = + chunks[0][col_idx].column_num_rows - chunks[0][col_idx].parent_validity_info.null_count; + auto parent_mask_len = chunks[0][col_idx].column_num_rows; - if (child_valid_map_base) { + if (child_valid_map_base != nullptr) { rmm::device_uvector dst_idx(child_mask_len, stream); // Copy indexes at which the parent has valid value. thrust::copy_if(rmm::exec_policy(stream), @@ -711,14 +715,13 @@ void update_null_mask(cudf::detail::hostdevice_2dvector& chunks thrust::make_counting_iterator(0) + parent_mask_len, dst_idx.begin(), [parent_valid_map_base] __device__(auto idx) { - auto val = bit_is_set(parent_valid_map_base, idx); - return val; + return bit_is_set(parent_valid_map_base, idx); }); auto merged_null_mask = cudf::detail::create_null_mask( parent_mask_len, mask_state::ALL_NULL, rmm::cuda_stream_view(stream), mr); - bitmask_type* merged_mask = static_cast(merged_null_mask.data()); - uint32_t* dst_idx_ptr = dst_idx.data(); + auto merged_mask = static_cast(merged_null_mask.data()); + uint32_t* dst_idx_ptr = dst_idx.data(); // Copy child valid bits from child column to valid indexes, this will merge both child and // parent null masks thrust::for_each(rmm::exec_policy(stream), @@ -730,23 +733,22 @@ void update_null_mask(cudf::detail::hostdevice_2dvector& chunks }; }); - out_buffers[j]._null_mask = std::move(merged_null_mask); + out_buffers[col_idx]._null_mask = std::move(merged_null_mask); } else { // Since child column doesn't have a mask, copy parent null mask - auto mask_size = bitmask_allocation_size_bytes(parent_mask_len); - out_buffers[j]._null_mask = std::move( - rmm::device_buffer(static_cast(parent_valid_map_base), mask_size, stream, mr)); + auto mask_size = bitmask_allocation_size_bytes(parent_mask_len); + out_buffers[col_idx]._null_mask = + rmm::device_buffer(static_cast(parent_valid_map_base), mask_size, stream, mr); } - if (not is_mask_updated) is_mask_updated = true; } } if (is_mask_updated) { // Update chunks with pointers to column data which might have been changed. - for (size_t i = 0; i < num_stripes; ++i) { - for (size_t j = 0; j < num_columns; ++j) { - auto& chunk = chunks[i][j]; - chunk.valid_map_base = out_buffers[j].null_mask(); + for (size_t stripe_idx = 0; stripe_idx < num_stripes; ++stripe_idx) { + for (size_t col_idx = 0; col_idx < num_columns; ++col_idx) { + auto& chunk = chunks[stripe_idx][col_idx]; + chunk.valid_map_base = out_buffers[col_idx].null_mask(); } } chunks.host_to_device(stream, true); @@ -801,12 +803,12 @@ void reader::impl::decode_stream_data(cudf::detail::hostdevice_2dvector(out_buffers[parent_col_idx].null_count()); - auto parent_valid_map = static_cast(out_buffers[parent_col_idx].null_mask()); + auto parent_valid_map = out_buffers[parent_col_idx].null_mask(); auto num_rows = out_buffers[parent_col_idx].size; for (uint32_t id = 0; id < p_col.num_children; id++) { @@ -1193,9 +1195,9 @@ table_with_metadata reader::impl::read(size_type skip_rows, ? stripe_info->numberOfRows : _col_meta.num_child_rows_per_stripe[stripe_idx * num_columns + col_idx]; chunk.column_num_rows = (level == 0) ? num_rows : _col_meta.num_child_rows[col_idx]; - chunk.parent_data.valid_map_base = + chunk.parent_validity_info.valid_map_base = (level == 0) ? nullptr : _col_meta.parent_column_data[col_idx].valid_map_base; - chunk.parent_data.null_count = + chunk.parent_validity_info.null_count = (level == 0) ? 0 : _col_meta.parent_column_data[col_idx].null_count; chunk.encoding_kind = stripe_footer->columns[selected_columns[col_idx].id].kind; chunk.type_kind = _metadata->per_file_metadata[stripe_source_mapping.source_idx] diff --git a/cpp/src/io/orc/reader_impl.hpp b/cpp/src/io/orc/reader_impl.hpp index 58c38c07f09..9f6f7b82b35 100644 --- a/cpp/src/io/orc/reader_impl.hpp +++ b/cpp/src/io/orc/reader_impl.hpp @@ -56,11 +56,7 @@ struct reader_column_meta { orc_col_map; // Mapping between column id in orc to processing order. std::vector num_child_rows; // number of rows in child columns - struct ParentColumnData { - cudf::bitmask_type* valid_map_base = nullptr; - uint32_t null_count; - }; - std::vector + std::vector parent_column_data; // consists of parent column valid_map and null count std::vector child_start_row; // start row of child columns [stripe][column] diff --git a/cpp/src/io/orc/stripe_data.cu b/cpp/src/io/orc/stripe_data.cu index 75586055892..75e96f1cc6a 100644 --- a/cpp/src/io/orc/stripe_data.cu +++ b/cpp/src/io/orc/stripe_data.cu @@ -1152,7 +1152,7 @@ __global__ void __launch_bounds__(block_size) if (t == 0) s->chunk = chunks[chunk_id]; __syncthreads(); - const size_t max_num_rows = s->chunk.column_num_rows - s->chunk.parent_data.null_count; + const size_t max_num_rows = s->chunk.column_num_rows - s->chunk.parent_validity_info.null_count; if (is_nulldec) { uint32_t null_count = 0; @@ -1338,7 +1338,7 @@ static __device__ void DecodeRowPositions(orcdec_state_s* s, // Even though s->chunk.strm_len is zero, there is possibility that there is null mask. // This happens in a struct column with nulls which has child column which doesn't have any // nulls. - if (s->chunk.strm_len[CI_PRESENT] > 0 or s->chunk.valid_map_base) { + if (s->chunk.strm_len[CI_PRESENT] > 0 or s->chunk.valid_map_base != NULL) { // We have a present stream uint32_t rmax = s->top.data.end_row - min((uint32_t)first_row, s->top.data.end_row); uint32_t r = (uint32_t)(s->top.data.cur_row + s->top.data.nrows + t - first_row); diff --git a/python/cudf/cudf/tests/test_orc.py b/python/cudf/cudf/tests/test_orc.py index 82f61b34442..41e96eb7782 100644 --- a/python/cudf/cudf/tests/test_orc.py +++ b/python/cudf/cudf/tests/test_orc.py @@ -843,6 +843,7 @@ def test_orc_string_stream_offset_issue(): assert_eq(df, cudf.read_orc(buffer)) +# Data is generated using pyorc module def generate_list_struct_buff(size=28000): rd = random.Random(1) np.random.seed(seed=1) @@ -984,7 +985,7 @@ def test_lists_struct_nests( ) if num_rows > 0: - assert_eq(True, pyarrow_tbl.equals(gdf.to_arrow())) + assert pyarrow_tbl.equals(gdf.to_arrow()) else: assert_eq(pyarrow_tbl.to_pandas(), gdf) From 40c81153bde093aaba85c2db56a76a5797c2537c Mon Sep 17 00:00:00 2001 From: Ramakrishna Prabhu Date: Wed, 21 Jul 2021 21:07:36 -0500 Subject: [PATCH 13/26] review changes --- cpp/src/io/orc/reader_impl.cu | 42 ++++++++++++++++++++++------------- 1 file changed, 26 insertions(+), 16 deletions(-) diff --git a/cpp/src/io/orc/reader_impl.cu b/cpp/src/io/orc/reader_impl.cu index 873825295ea..3ddd5a6a1d8 100644 --- a/cpp/src/io/orc/reader_impl.cu +++ b/cpp/src/io/orc/reader_impl.cu @@ -743,14 +743,18 @@ void update_null_mask(cudf::detail::hostdevice_2dvector& chunks } } } + + thrust::counting_iterator col_idx_it(0); + thrust::counting_iterator stripe_idx_it(0); + if (is_mask_updated) { // Update chunks with pointers to column data which might have been changed. - for (size_t stripe_idx = 0; stripe_idx < num_stripes; ++stripe_idx) { - for (size_t col_idx = 0; col_idx < num_columns; ++col_idx) { + std::for_each(stripe_idx_it + 0, stripe_idx_it + num_stripes, [&](auto stripe_idx) { + std::for_each(col_idx_it + 0, col_idx_it + num_columns, [&](auto col_idx) { auto& chunk = chunks[stripe_idx][col_idx]; chunk.valid_map_base = out_buffers[col_idx].null_mask(); - } - } + }); + }); chunks.host_to_device(stream, true); } } @@ -767,15 +771,17 @@ void reader::impl::decode_stream_data(cudf::detail::hostdevice_2dvector col_idx_it(0); + thrust::counting_iterator stripe_idx_it(0); // Update chunks with pointers to column data - for (size_t i = 0; i < num_stripes; ++i) { - for (size_t j = 0; j < num_columns; ++j) { - auto& chunk = chunks[i][j]; - chunk.column_data_base = out_buffers[j].data(); - chunk.valid_map_base = out_buffers[j].null_mask(); - } - } + std::for_each(stripe_idx_it + 0, stripe_idx_it + num_stripes, [&](auto stripe_idx) { + std::for_each(col_idx_it + 0, col_idx_it + num_columns, [&](auto col_idx) { + auto& chunk = chunks[stripe_idx][col_idx]; + chunk.column_data_base = out_buffers[col_idx].data(); + chunk.valid_map_base = out_buffers[col_idx].null_mask(); + }); + }); // Allocate global dictionary for deserializing rmm::device_uvector global_dict(num_dicts, stream); @@ -803,13 +809,17 @@ void reader::impl::decode_stream_data(cudf::detail::hostdevice_2dvector Date: Thu, 22 Jul 2021 12:14:47 -0500 Subject: [PATCH 14/26] review changes --- cpp/src/io/orc/orc_gpu.h | 8 -------- cpp/src/io/orc/reader_impl.cu | 6 +++--- 2 files changed, 3 insertions(+), 11 deletions(-) diff --git a/cpp/src/io/orc/orc_gpu.h b/cpp/src/io/orc/orc_gpu.h index 45149d56853..3ccc2c672c9 100644 --- a/cpp/src/io/orc/orc_gpu.h +++ b/cpp/src/io/orc/orc_gpu.h @@ -86,14 +86,6 @@ struct DictionaryEntry { */ constexpr int orc_decimal2float64_scale = 0x80; -/** - * Struct to store parent column mask and null count - */ -struct ParentColumnData { - uint32_t* valid_map_base; - uint32_t null_count; -}; - /** * @brief Struct to describe per stripe's column information */ diff --git a/cpp/src/io/orc/reader_impl.cu b/cpp/src/io/orc/reader_impl.cu index 3ddd5a6a1d8..10f6120d07b 100644 --- a/cpp/src/io/orc/reader_impl.cu +++ b/cpp/src/io/orc/reader_impl.cu @@ -182,7 +182,7 @@ size_t gather_stream_info(const size_t stripe_index, size_t* num_dictionary_entries, cudf::detail::hostdevice_2dvector& chunks, std::vector& stream_info, - size_t level) + bool apply_struct_map) { uint64_t src_offset = 0; uint64_t dst_offset = 0; @@ -195,7 +195,7 @@ size_t gather_stream_info(const size_t stripe_index, auto const column_id = *stream.column_id; auto col = orc2gdf[column_id]; - if (col == -1 and level == 0) { + if (col == -1 and apply_struct_map) { // A struct-type column has no data itself, but rather child columns // for each of its fields. There is only a PRESENT stream, which // needs to be included for the reader. @@ -1149,7 +1149,7 @@ table_with_metadata reader::impl::read(size_type skip_rows, &num_dict_entries, chunks, stream_info, - level); + level == 0); CUDF_EXPECTS(total_data_size > 0, "Expected streams data within stripe"); From 2d3842ad7de979a1b9872b649aa617d16285dc12 Mon Sep 17 00:00:00 2001 From: Ramakrishna Prabhu Date: Thu, 22 Jul 2021 13:16:45 -0500 Subject: [PATCH 15/26] review changes --- cpp/src/io/orc/reader_impl.cu | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/cpp/src/io/orc/reader_impl.cu b/cpp/src/io/orc/reader_impl.cu index 10f6120d07b..8e47da98a7c 100644 --- a/cpp/src/io/orc/reader_impl.cu +++ b/cpp/src/io/orc/reader_impl.cu @@ -749,8 +749,8 @@ void update_null_mask(cudf::detail::hostdevice_2dvector& chunks if (is_mask_updated) { // Update chunks with pointers to column data which might have been changed. - std::for_each(stripe_idx_it + 0, stripe_idx_it + num_stripes, [&](auto stripe_idx) { - std::for_each(col_idx_it + 0, col_idx_it + num_columns, [&](auto col_idx) { + std::for_each(stripe_idx_it, stripe_idx_it + num_stripes, [&](auto stripe_idx) { + std::for_each(col_idx_it, col_idx_it + num_columns, [&](auto col_idx) { auto& chunk = chunks[stripe_idx][col_idx]; chunk.valid_map_base = out_buffers[col_idx].null_mask(); }); @@ -775,8 +775,8 @@ void reader::impl::decode_stream_data(cudf::detail::hostdevice_2dvector stripe_idx_it(0); // Update chunks with pointers to column data - std::for_each(stripe_idx_it + 0, stripe_idx_it + num_stripes, [&](auto stripe_idx) { - std::for_each(col_idx_it + 0, col_idx_it + num_columns, [&](auto col_idx) { + std::for_each(stripe_idx_it, stripe_idx_it + num_stripes, [&](auto stripe_idx) { + std::for_each(col_idx_it, col_idx_it + num_columns, [&](auto col_idx) { auto& chunk = chunks[stripe_idx][col_idx]; chunk.column_data_base = out_buffers[col_idx].data(); chunk.valid_map_base = out_buffers[col_idx].null_mask(); From bf57b6b6dcce849cc17c33be1c5be9666d18208e Mon Sep 17 00:00:00 2001 From: Ramakrishna Prabhu Date: Thu, 22 Jul 2021 14:10:18 -0500 Subject: [PATCH 16/26] review changes --- cpp/src/io/orc/stripe_data.cu | 5 +---- 1 file changed, 1 insertion(+), 4 deletions(-) diff --git a/cpp/src/io/orc/stripe_data.cu b/cpp/src/io/orc/stripe_data.cu index 75e96f1cc6a..923dc366d74 100644 --- a/cpp/src/io/orc/stripe_data.cu +++ b/cpp/src/io/orc/stripe_data.cu @@ -1335,10 +1335,7 @@ static __device__ void DecodeRowPositions(orcdec_state_s* s, s->top.data.cur_row + s->top.data.nrows < s->top.data.end_row) { uint32_t nrows = min(s->top.data.end_row - (s->top.data.cur_row + s->top.data.nrows), min((row_decoder_buffer_size - s->u.rowdec.nz_count) * 2, blockDim.x)); - // Even though s->chunk.strm_len is zero, there is possibility that there is null mask. - // This happens in a struct column with nulls which has child column which doesn't have any - // nulls. - if (s->chunk.strm_len[CI_PRESENT] > 0 or s->chunk.valid_map_base != NULL) { + if (s->chunk.valid_map_base != NULL) { // We have a present stream uint32_t rmax = s->top.data.end_row - min((uint32_t)first_row, s->top.data.end_row); uint32_t r = (uint32_t)(s->top.data.cur_row + s->top.data.nrows + t - first_row); From 400a4a3e37f5e6c98262cd85bb5be558fedf57bf Mon Sep 17 00:00:00 2001 From: Ramakrishna Prabhu Date: Wed, 18 Aug 2021 01:20:35 -0500 Subject: [PATCH 17/26] changes --- cpp/src/io/orc/reader_impl.cu | 46 +++++++++++++----- .../TestOrcFile.NoIndStrm.IntWithNulls.orc | Bin 0 -> 101 bytes ...dStrm.StructAndIntWithNulls.TwoStripes.orc | Bin 0 -> 232 bytes ...rcFile.NoIndStrm.StructAndIntWithNulls.orc | Bin 0 -> 193 bytes .../TestOrcFile.NoIndStrm.StructWithNulls.orc | Bin 0 -> 193 bytes python/cudf/cudf/tests/test_orc.py | 18 +++++++ 6 files changed, 51 insertions(+), 13 deletions(-) create mode 100644 python/cudf/cudf/tests/data/orc/TestOrcFile.NoIndStrm.IntWithNulls.orc create mode 100644 python/cudf/cudf/tests/data/orc/TestOrcFile.NoIndStrm.StructAndIntWithNulls.TwoStripes.orc create mode 100644 python/cudf/cudf/tests/data/orc/TestOrcFile.NoIndStrm.StructAndIntWithNulls.orc create mode 100644 python/cudf/cudf/tests/data/orc/TestOrcFile.NoIndStrm.StructWithNulls.orc diff --git a/cpp/src/io/orc/reader_impl.cu b/cpp/src/io/orc/reader_impl.cu index 033a2d9aff5..3f623a8f9fb 100644 --- a/cpp/src/io/orc/reader_impl.cu +++ b/cpp/src/io/orc/reader_impl.cu @@ -269,6 +269,7 @@ class aggregate_orc_metadata { size_type const num_rows; size_type const num_columns; size_type const num_stripes; + bool row_grp_idx = true; /** * @brief Create a metadata object from each element in the source vector @@ -368,6 +369,8 @@ class aggregate_orc_metadata { return per_file_metadata[source_idx].get_column_name(column_idx); } + auto is_row_grp_idx_avbl() const { return row_grp_idx; } + std::vector select_stripes( std::vector> const& user_specified_stripes, size_type& row_start, @@ -457,6 +460,7 @@ class aggregate_orc_metadata { ProtobufReader(sf_data, sf_length) .read(per_file_metadata[mapping.source_idx].stripefooters[i]); mapping.stripe_info[i].second = &per_file_metadata[mapping.source_idx].stripefooters[i]; + if (stripe->indexLength == 0) row_grp_idx = false; } } } @@ -1057,6 +1061,7 @@ table_with_metadata reader::impl::read(size_type skip_rows, // Association between each ORC column and its cudf::column _col_meta.orc_col_map.emplace_back(_metadata->get_num_cols(), -1); std::vector nested_col; + bool is_stream_empty = false; // Get a list of column data types std::vector column_types; @@ -1113,6 +1118,8 @@ table_with_metadata reader::impl::read(size_type skip_rows, const bool use_index = (_use_index == true) && + // Do stripes have row group index + _metadata->is_row_grp_idx_avbl() && // Only use if we don't have much work with complete columns & stripes // TODO: Consider nrows, gpu, and tune the threshold (num_rows > _metadata->get_row_index_stride() && !(_metadata->get_row_index_stride() & 7) && @@ -1152,13 +1159,24 @@ table_with_metadata reader::impl::read(size_type skip_rows, stream_info, level == 0); - CUDF_EXPECTS(total_data_size > 0, "Expected streams data within stripe"); + if (total_data_size == 0) { + auto valid_stream_data = false; + // In case ROW GROUP INDEX is not present and all columns are structs with no null + // stream, there is nothing to read at this level, so we have to skip assertion. + if (stripe_info->indexLength == 0) { + auto fn_check_dtype = [](auto dtype) { return dtype.id() == type_id::STRUCT; }; + valid_stream_data = + std::all_of(column_types.begin(), column_types.end(), fn_check_dtype); + } + CUDF_EXPECTS(valid_stream_data, "Expected streams data within stripe"); + is_stream_empty = true; + } stripe_data.emplace_back(total_data_size, stream); auto dst_base = static_cast(stripe_data.back().data()); // Coalesce consecutive streams into one read - while (stream_count < stream_info.size()) { + while (not is_stream_empty and stream_count < stream_info.size()) { const auto d_dst = dst_base + stream_info[stream_count].dst_pos; const auto offset = stream_info[stream_count].offset; auto len = stream_info[stream_count].length; @@ -1238,7 +1256,7 @@ table_with_metadata reader::impl::read(size_type skip_rows, if (chunk.type_kind == orc::TIMESTAMP) { chunk.ts_clock_rate = to_clockrate(_timestamp_type.id()); } - for (int k = 0; k < gpu::CI_NUM_STREAMS; k++) { + for (int k = 0; k < gpu::CI_NUM_STREAMS and not is_stream_empty; k++) { chunk.streams[k] = dst_base + stream_info[chunk.strm_id[k]].dst_pos; } } @@ -1273,7 +1291,7 @@ table_with_metadata reader::impl::read(size_type skip_rows, }); } // Setup row group descriptors if using indexes - if (_metadata->per_file_metadata[0].ps.compression != orc::NONE) { + if (_metadata->per_file_metadata[0].ps.compression != orc::NONE and not is_stream_empty) { auto decomp_data = decompress_stripe_data(chunks, stripe_data, @@ -1324,15 +1342,17 @@ table_with_metadata reader::impl::read(size_type skip_rows, out_buffers[level].emplace_back(column_types[i], n_rows, is_nullable, stream, _mr); } - decode_stream_data(chunks, - num_dict_entries, - skip_rows, - tz_table.view(), - row_groups, - _metadata->get_row_index_stride(), - out_buffers[level], - level, - stream); + if (not is_stream_empty) { + decode_stream_data(chunks, + num_dict_entries, + skip_rows, + tz_table.view(), + row_groups, + _metadata->get_row_index_stride(), + out_buffers[level], + level, + stream); + } // Extract information to process nested child columns if (nested_col.size()) { diff --git a/python/cudf/cudf/tests/data/orc/TestOrcFile.NoIndStrm.IntWithNulls.orc b/python/cudf/cudf/tests/data/orc/TestOrcFile.NoIndStrm.IntWithNulls.orc new file mode 100644 index 0000000000000000000000000000000000000000..2103e0212fcdcc9a110e0dbe550d2fcb94bb640d GIT binary patch literal 101 zcmeYda%N><_ewK`ddygwP8Y z9x!aOO~_G|w<_diX#vDJf-glFyVRpXLuUW-OJMyhLL1)EO2t6qR(=1odbJ^{h+~)G&L%z`)@k Zz$nq6pv1($qrqUn+4O~(*+0lx3;@GaM5X`$ literal 0 HcmV?d00001 diff --git a/python/cudf/cudf/tests/data/orc/TestOrcFile.NoIndStrm.StructWithNulls.orc b/python/cudf/cudf/tests/data/orc/TestOrcFile.NoIndStrm.StructWithNulls.orc new file mode 100644 index 0000000000000000000000000000000000000000..32d0c85dd25b7cbe4c9e3c434b438f57cab52bc2 GIT binary patch literal 193 zcmeYda%N><_diX#vDJf-glFyVRpXLuUW-OJMyhLL1)EO2t6qR(=1odbJ^{h+~)G&L%z`)@k Zz$nq6pv1($qrqUn+4O~(*+0lx3;@GaM5X`$ literal 0 HcmV?d00001 diff --git a/python/cudf/cudf/tests/test_orc.py b/python/cudf/cudf/tests/test_orc.py index 635332d5c24..cf4601a2b37 100644 --- a/python/cudf/cudf/tests/test_orc.py +++ b/python/cudf/cudf/tests/test_orc.py @@ -1171,3 +1171,21 @@ def test_writer_timestamp_stream_size(datadir, tmpdir): got = pa.orc.ORCFile(gdf_fname).read().to_pandas() assert_eq(expect, got) + + +@pytest.mark.parametrize( + "fname", + [ + "TestOrcFile.NoIndStrm.StructWithNoNulls.orc", + "TestOrcFile.NoIndStrm.StructAndIntWithNulls.orc", + "TestOrcFile.NoIndStrm.StructAndIntWithNulls.TwoStripes.orc", + "TestOrcFile.NoIndStrm.IntWithNulls.orc", + ], +) +def test_no_row_group_index_orc_read(datadir, fname): + fpath = datadir / fname + + got = pa.orc.ORCFile(fpath).read() + expect = cudf.read_orc(fpath) + + assert got.equals(expect.to_arrow()) From 3c75f87fee6759d96388035eceebf21827a9865b Mon Sep 17 00:00:00 2001 From: Ramakrishna Prabhu Date: Wed, 18 Aug 2021 02:05:58 -0500 Subject: [PATCH 18/26] changes --- cpp/src/io/orc/reader_impl.cu | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/cpp/src/io/orc/reader_impl.cu b/cpp/src/io/orc/reader_impl.cu index d051ee31644..220d1231318 100644 --- a/cpp/src/io/orc/reader_impl.cu +++ b/cpp/src/io/orc/reader_impl.cu @@ -269,7 +269,7 @@ class aggregate_orc_metadata { size_type const num_rows; size_type const num_columns; size_type const num_stripes; - bool row_grp_idx = true; + bool row_grp_idx_present = true; /** * @brief Create a metadata object from each element in the source vector @@ -369,7 +369,7 @@ class aggregate_orc_metadata { return per_file_metadata[source_idx].get_column_name(column_idx); } - auto is_row_grp_idx_avbl() const { return row_grp_idx; } + auto is_row_grp_idx_present() const { return row_grp_idx_present; } std::vector select_stripes( std::vector> const& user_specified_stripes, @@ -460,7 +460,7 @@ class aggregate_orc_metadata { ProtobufReader(sf_data, sf_length) .read(per_file_metadata[mapping.source_idx].stripefooters[i]); mapping.stripe_info[i].second = &per_file_metadata[mapping.source_idx].stripefooters[i]; - if (stripe->indexLength == 0) row_grp_idx = false; + if (stripe->indexLength == 0) row_grp_idx_present = false; } } } @@ -1163,7 +1163,7 @@ table_with_metadata reader::impl::read(size_type skip_rows, const bool use_index = (_use_index == true) && // Do stripes have row group index - _metadata->is_row_grp_idx_avbl() && + _metadata->is_row_grp_idx_present() && // Only use if we don't have much work with complete columns & stripes // TODO: Consider nrows, gpu, and tune the threshold (num_rows > _metadata->get_row_index_stride() && !(_metadata->get_row_index_stride() & 7) && From 117347ec1f7700b3e59a5aba6494098d618e3ce8 Mon Sep 17 00:00:00 2001 From: Ramakrishna Prabhu Date: Wed, 18 Aug 2021 08:59:23 -0500 Subject: [PATCH 19/26] adding missing test file --- .../TestOrcFile.NoIndStrm.StructWithNoNulls.orc | Bin 0 -> 167 bytes .../orc/TestOrcFile.NoIndStrm.StructWithNulls.orc | Bin 193 -> 0 bytes 2 files changed, 0 insertions(+), 0 deletions(-) create mode 100644 python/cudf/cudf/tests/data/orc/TestOrcFile.NoIndStrm.StructWithNoNulls.orc delete mode 100644 python/cudf/cudf/tests/data/orc/TestOrcFile.NoIndStrm.StructWithNulls.orc diff --git a/python/cudf/cudf/tests/data/orc/TestOrcFile.NoIndStrm.StructWithNoNulls.orc b/python/cudf/cudf/tests/data/orc/TestOrcFile.NoIndStrm.StructWithNoNulls.orc new file mode 100644 index 0000000000000000000000000000000000000000..1c6e53a0b92f68b996b4a249b7050eae715331e9 GIT binary patch literal 167 zcmeYda%N{>U}a3=E$=&YnDZ-Sh)PaGSs>4kt-x4PlN{u7zh5rdTp?f|Qb!=@QCW-OgC`Ejpf%R(Vxc42|d(gh3*9A*NH5)BGUObpx_ P3<_diX#vDJf-glFyVRpXLuUW-OJMyhLL1)EO2t6qR(=1odbJ^{h+~)G&L%z`)@k Zz$nq6pv1($qrqUn+4O~(*+0lx3;@GaM5X`$ From 156c6594d8d7b71321d35c81ca9e8933dfe2ce30 Mon Sep 17 00:00:00 2001 From: Ramakrishna Prabhu Date: Thu, 19 Aug 2021 05:10:35 -0500 Subject: [PATCH 20/26] testing --- cpp/src/io/orc/reader_impl.cu | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/cpp/src/io/orc/reader_impl.cu b/cpp/src/io/orc/reader_impl.cu index 220d1231318..2aec6c650df 100644 --- a/cpp/src/io/orc/reader_impl.cu +++ b/cpp/src/io/orc/reader_impl.cu @@ -1440,6 +1440,10 @@ table_with_metadata reader::impl::read(size_type skip_rows, create_columns(std::move(out_buffers), out_columns, schema_info, stream); } + if (not out_columns.empty()) { + printf("RGSL: null count is %d \n", out_columns[0]->null_count()); + auto col = out_columns[0]->view(); + } // Return column names (must match order of returned columns) out_metadata.column_names.reserve(schema_info.size()); std::transform(schema_info.cbegin(), From 2bebc3e96eeb2c94c8268177b3e379791b518468 Mon Sep 17 00:00:00 2001 From: Ramakrishna Prabhu Date: Thu, 19 Aug 2021 05:51:05 -0500 Subject: [PATCH 21/26] testing --- cpp/src/column/column_view.cpp | 1 + cpp/src/io/orc/reader_impl.cu | 8 +++++--- 2 files changed, 6 insertions(+), 3 deletions(-) diff --git a/cpp/src/column/column_view.cpp b/cpp/src/column/column_view.cpp index 186669ae697..ae856e688f2 100644 --- a/cpp/src/column/column_view.cpp +++ b/cpp/src/column/column_view.cpp @@ -56,6 +56,7 @@ column_view_base::column_view_base(data_type type, CUDF_EXPECTS(offset >= 0, "Invalid offset."); if ((null_count > 0) and (type.id() != type_id::EMPTY)) { + if (nullptr == null_mask) printf("RGSL: The null count is %d \n", null_count); CUDF_EXPECTS(nullptr != null_mask, "Invalid null mask for non-zero null count."); } } diff --git a/cpp/src/io/orc/reader_impl.cu b/cpp/src/io/orc/reader_impl.cu index 2aec6c650df..29c8786bb68 100644 --- a/cpp/src/io/orc/reader_impl.cu +++ b/cpp/src/io/orc/reader_impl.cu @@ -1214,7 +1214,7 @@ table_with_metadata reader::impl::read(size_type skip_rows, if (total_data_size == 0) { auto valid_stream_data = false; // In case ROW GROUP INDEX is not present and all columns are structs with no null - // stream, there is nothing to read at this level, so we have to skip assertion. + // stream, there is nothing to read at this level. if (stripe_info->indexLength == 0) { auto fn_check_dtype = [](auto dtype) { return dtype.id() == type_id::STRUCT; }; valid_stream_data = @@ -1310,8 +1310,10 @@ table_with_metadata reader::impl::read(size_type skip_rows, if (chunk.type_kind == orc::TIMESTAMP) { chunk.ts_clock_rate = to_clockrate(_timestamp_type.id()); } - for (int k = 0; k < gpu::CI_NUM_STREAMS and not is_stream_empty; k++) { - chunk.streams[k] = dst_base + stream_info[chunk.strm_id[k]].dst_pos; + if (not is_stream_empty) { + for (int k = 0; k < gpu::CI_NUM_STREAMS; k++) { + chunk.streams[k] = dst_base + stream_info[chunk.strm_id[k]].dst_pos; + } } } stripe_start_row += num_rows_per_stripe; From d09c85524c046496e9c71e7b0d72619fd67b3847 Mon Sep 17 00:00:00 2001 From: Ramakrishna Prabhu Date: Fri, 20 Aug 2021 01:39:47 -0500 Subject: [PATCH 22/26] testing prints --- cpp/src/column/column_view.cpp | 6 +++++- cpp/src/io/orc/reader_impl.cu | 4 ++++ 2 files changed, 9 insertions(+), 1 deletion(-) diff --git a/cpp/src/column/column_view.cpp b/cpp/src/column/column_view.cpp index ae856e688f2..ff9b48d7fa3 100644 --- a/cpp/src/column/column_view.cpp +++ b/cpp/src/column/column_view.cpp @@ -56,7 +56,11 @@ column_view_base::column_view_base(data_type type, CUDF_EXPECTS(offset >= 0, "Invalid offset."); if ((null_count > 0) and (type.id() != type_id::EMPTY)) { - if (nullptr == null_mask) printf("RGSL: The null count is %d \n", null_count); + if (nullptr == null_mask) { + printf("RGSL: The null count is %d and column type is %d \n", + null_count, + static_cast(type.id())); + } CUDF_EXPECTS(nullptr != null_mask, "Invalid null mask for non-zero null count."); } } diff --git a/cpp/src/io/orc/reader_impl.cu b/cpp/src/io/orc/reader_impl.cu index 29c8786bb68..fa892824633 100644 --- a/cpp/src/io/orc/reader_impl.cu +++ b/cpp/src/io/orc/reader_impl.cu @@ -1444,6 +1444,10 @@ table_with_metadata reader::impl::read(size_type skip_rows, if (not out_columns.empty()) { printf("RGSL: null count is %d \n", out_columns[0]->null_count()); + if (out_columns[0]->num_children() == 2) { + printf("RGSL: null count of child 0 %d \n", out_columns[0]->child(0).null_count()); + printf("RGSL: null count of child 1 %d \n", out_columns[0]->child(1).null_count()); + } auto col = out_columns[0]->view(); } // Return column names (must match order of returned columns) From 860757ee57c5c98fb71dba9e88efd10b977fa21f Mon Sep 17 00:00:00 2001 From: Ramakrishna Prabhu Date: Fri, 20 Aug 2021 06:58:59 -0500 Subject: [PATCH 23/26] some more testing --- cpp/src/io/orc/reader_impl.cu | 23 +++++++++++++++-------- 1 file changed, 15 insertions(+), 8 deletions(-) diff --git a/cpp/src/io/orc/reader_impl.cu b/cpp/src/io/orc/reader_impl.cu index fa892824633..706f7d271e4 100644 --- a/cpp/src/io/orc/reader_impl.cu +++ b/cpp/src/io/orc/reader_impl.cu @@ -857,13 +857,16 @@ void reader::impl::decode_stream_data(cudf::detail::hostdevice_2dvector(out_buffers[col_idx].type.id())); + return null_count + chunks[stripe_idx][col_idx].null_count; + }); }); } @@ -1412,7 +1415,7 @@ table_with_metadata reader::impl::read(size_type skip_rows, // Extract information to process nested child columns if (nested_col.size()) { - scan_null_counts(chunks, null_count_prefix_sums[level], stream); + if (not is_stream_empty) scan_null_counts(chunks, null_count_prefix_sums[level], stream); row_groups.device_to_host(stream, true); aggregate_child_meta(chunks, row_groups, out_buffers[level], nested_col, level); } @@ -1447,6 +1450,10 @@ table_with_metadata reader::impl::read(size_type skip_rows, if (out_columns[0]->num_children() == 2) { printf("RGSL: null count of child 0 %d \n", out_columns[0]->child(0).null_count()); printf("RGSL: null count of child 1 %d \n", out_columns[0]->child(1).null_count()); + // printf("RGSL : number of children in string column is %d\n", + // out_columns[0]->child(1).num_children()); printf("RGSL: null count of child 1-0 %d \n", + // out_columns[0]->child(1).child(0).null_count()); printf("RGSL: null count of child 1-1 %d + // \n", out_columns[0]->child(1).child(1).null_count()); } auto col = out_columns[0]->view(); } From e15ac385ddca3230485326f31c0b215d04103a49 Mon Sep 17 00:00:00 2001 From: Ramakrishna Prabhu Date: Fri, 20 Aug 2021 08:48:16 -0500 Subject: [PATCH 24/26] changes and removing prints --- cpp/src/column/column_view.cpp | 5 ----- cpp/src/io/orc/reader_impl.cu | 29 +++++++---------------------- 2 files changed, 7 insertions(+), 27 deletions(-) diff --git a/cpp/src/column/column_view.cpp b/cpp/src/column/column_view.cpp index ff9b48d7fa3..186669ae697 100644 --- a/cpp/src/column/column_view.cpp +++ b/cpp/src/column/column_view.cpp @@ -56,11 +56,6 @@ column_view_base::column_view_base(data_type type, CUDF_EXPECTS(offset >= 0, "Invalid offset."); if ((null_count > 0) and (type.id() != type_id::EMPTY)) { - if (nullptr == null_mask) { - printf("RGSL: The null count is %d and column type is %d \n", - null_count, - static_cast(type.id())); - } CUDF_EXPECTS(nullptr != null_mask, "Invalid null mask for non-zero null count."); } } diff --git a/cpp/src/io/orc/reader_impl.cu b/cpp/src/io/orc/reader_impl.cu index 706f7d271e4..d78cf64335f 100644 --- a/cpp/src/io/orc/reader_impl.cu +++ b/cpp/src/io/orc/reader_impl.cu @@ -857,16 +857,13 @@ void reader::impl::decode_stream_data(cudf::detail::hostdevice_2dvector(out_buffers[col_idx].type.id())); - return null_count + chunks[stripe_idx][col_idx].null_count; - }); + out_buffers[col_idx].null_count() = + std::accumulate(stripe_idx_it + 0, + stripe_idx_it + num_stripes, + 0, + [&](auto null_count, auto const stripe_idx) { + return null_count + chunks[stripe_idx][col_idx].null_count; + }); }); } @@ -1445,18 +1442,6 @@ table_with_metadata reader::impl::read(size_type skip_rows, create_columns(std::move(out_buffers), out_columns, schema_info, stream); } - if (not out_columns.empty()) { - printf("RGSL: null count is %d \n", out_columns[0]->null_count()); - if (out_columns[0]->num_children() == 2) { - printf("RGSL: null count of child 0 %d \n", out_columns[0]->child(0).null_count()); - printf("RGSL: null count of child 1 %d \n", out_columns[0]->child(1).null_count()); - // printf("RGSL : number of children in string column is %d\n", - // out_columns[0]->child(1).num_children()); printf("RGSL: null count of child 1-0 %d \n", - // out_columns[0]->child(1).child(0).null_count()); printf("RGSL: null count of child 1-1 %d - // \n", out_columns[0]->child(1).child(1).null_count()); - } - auto col = out_columns[0]->view(); - } // Return column names (must match order of returned columns) out_metadata.column_names.reserve(schema_info.size()); std::transform(schema_info.cbegin(), From a21d9b826b04607aa760af82fabb5c785b090bb2 Mon Sep 17 00:00:00 2001 From: Ramakrishna Prabhu Date: Fri, 20 Aug 2021 08:50:45 -0500 Subject: [PATCH 25/26] review changes --- cpp/src/io/orc/reader_impl.cu | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/cpp/src/io/orc/reader_impl.cu b/cpp/src/io/orc/reader_impl.cu index d78cf64335f..8c98ab9849f 100644 --- a/cpp/src/io/orc/reader_impl.cu +++ b/cpp/src/io/orc/reader_impl.cu @@ -460,7 +460,7 @@ class aggregate_orc_metadata { ProtobufReader(sf_data, sf_length) .read(per_file_metadata[mapping.source_idx].stripefooters[i]); mapping.stripe_info[i].second = &per_file_metadata[mapping.source_idx].stripefooters[i]; - if (stripe->indexLength == 0) row_grp_idx_present = false; + if (stripe->indexLength == 0) { row_grp_idx_present = false; } } } } @@ -1412,7 +1412,9 @@ table_with_metadata reader::impl::read(size_type skip_rows, // Extract information to process nested child columns if (nested_col.size()) { - if (not is_stream_empty) scan_null_counts(chunks, null_count_prefix_sums[level], stream); + if (not is_stream_empty) { + scan_null_counts(chunks, null_count_prefix_sums[level], stream); + } row_groups.device_to_host(stream, true); aggregate_child_meta(chunks, row_groups, out_buffers[level], nested_col, level); } From 3c51111f98eae45624bff1884faa5c948ad1b4db Mon Sep 17 00:00:00 2001 From: Ramakrishna Prabhu Date: Fri, 20 Aug 2021 22:13:30 -0500 Subject: [PATCH 26/26] review changes --- cpp/src/io/orc/reader_impl.cu | 25 +++++++++++-------------- 1 file changed, 11 insertions(+), 14 deletions(-) diff --git a/cpp/src/io/orc/reader_impl.cu b/cpp/src/io/orc/reader_impl.cu index 8c98ab9849f..b7b8e03d400 100644 --- a/cpp/src/io/orc/reader_impl.cu +++ b/cpp/src/io/orc/reader_impl.cu @@ -1105,7 +1105,7 @@ table_with_metadata reader::impl::read(size_type skip_rows, // Association between each ORC column and its cudf::column _col_meta.orc_col_map.emplace_back(_metadata->get_num_cols(), -1); std::vector nested_col; - bool is_stream_empty = false; + bool is_data_empty = false; // Get a list of column data types std::vector column_types; @@ -1212,23 +1212,20 @@ table_with_metadata reader::impl::read(size_type skip_rows, level == 0); if (total_data_size == 0) { - auto valid_stream_data = false; + CUDF_EXPECTS(stripe_info->indexLength == 0, "Invalid index rowgroup stream data"); // In case ROW GROUP INDEX is not present and all columns are structs with no null // stream, there is nothing to read at this level. - if (stripe_info->indexLength == 0) { - auto fn_check_dtype = [](auto dtype) { return dtype.id() == type_id::STRUCT; }; - valid_stream_data = - std::all_of(column_types.begin(), column_types.end(), fn_check_dtype); - } - CUDF_EXPECTS(valid_stream_data, "Expected streams data within stripe"); - is_stream_empty = true; + auto fn_check_dtype = [](auto dtype) { return dtype.id() == type_id::STRUCT; }; + CUDF_EXPECTS(std::all_of(column_types.begin(), column_types.end(), fn_check_dtype), + "Expected streams data within stripe"); + is_data_empty = true; } stripe_data.emplace_back(total_data_size, stream); auto dst_base = static_cast(stripe_data.back().data()); // Coalesce consecutive streams into one read - while (not is_stream_empty and stream_count < stream_info.size()) { + while (not is_data_empty and stream_count < stream_info.size()) { const auto d_dst = dst_base + stream_info[stream_count].dst_pos; const auto offset = stream_info[stream_count].offset; auto len = stream_info[stream_count].length; @@ -1310,7 +1307,7 @@ table_with_metadata reader::impl::read(size_type skip_rows, if (chunk.type_kind == orc::TIMESTAMP) { chunk.ts_clock_rate = to_clockrate(_timestamp_type.id()); } - if (not is_stream_empty) { + if (not is_data_empty) { for (int k = 0; k < gpu::CI_NUM_STREAMS; k++) { chunk.streams[k] = dst_base + stream_info[chunk.strm_id[k]].dst_pos; } @@ -1347,7 +1344,7 @@ table_with_metadata reader::impl::read(size_type skip_rows, }); } // Setup row group descriptors if using indexes - if (_metadata->per_file_metadata[0].ps.compression != orc::NONE and not is_stream_empty) { + if (_metadata->per_file_metadata[0].ps.compression != orc::NONE and not is_data_empty) { auto decomp_data = decompress_stripe_data(chunks, stripe_data, @@ -1398,7 +1395,7 @@ table_with_metadata reader::impl::read(size_type skip_rows, out_buffers[level].emplace_back(column_types[i], n_rows, is_nullable, stream, _mr); } - if (not is_stream_empty) { + if (not is_data_empty) { decode_stream_data(chunks, num_dict_entries, skip_rows, @@ -1412,7 +1409,7 @@ table_with_metadata reader::impl::read(size_type skip_rows, // Extract information to process nested child columns if (nested_col.size()) { - if (not is_stream_empty) { + if (not is_data_empty) { scan_null_counts(chunks, null_count_prefix_sums[level], stream); } row_groups.device_to_host(stream, true);