From 7fe82e953f264f5318fb080885c599856c4740dd Mon Sep 17 00:00:00 2001 From: Ramakrishna Prabhu Date: Tue, 9 Feb 2021 16:33:52 -0600 Subject: [PATCH 1/5] Chages to skipcount calculation --- cpp/src/io/orc/stripe_data.cu | 5 +++-- python/cudf/cudf/tests/test_orc.py | 26 ++++++++++++++++++++++++++ 2 files changed, 29 insertions(+), 2 deletions(-) diff --git a/cpp/src/io/orc/stripe_data.cu b/cpp/src/io/orc/stripe_data.cu index 856c23c0f55..686cf2fd923 100644 --- a/cpp/src/io/orc/stripe_data.cu +++ b/cpp/src/io/orc/stripe_data.cu @@ -1211,8 +1211,9 @@ __global__ void __launch_bounds__(block_size) uint32_t skippedrows = min(static_cast(first_row - row_in), nrows); uint32_t skip_count = 0; for (uint32_t i = t * 32; i < skippedrows; i += 32 * 32) { - uint32_t bits = s->vals.u32[i >> 5]; - if (i + 32 > skippedrows) { bits &= (1 << (skippedrows - i)) - 1; } + uint32_t bits = (i + 32 <= skippedrows) + ? s->vals.u32[i >> 5] + : (rle8_read_bool32(s->vals.u32, i) & ((1 << (skippedrows - i)) - 1)); skip_count += __popc(bits); } skip_count = warp_reduce(temp_storage[t / 32]).Sum(skip_count); diff --git a/python/cudf/cudf/tests/test_orc.py b/python/cudf/cudf/tests/test_orc.py index 62696303ecb..9681eda6fda 100644 --- a/python/cudf/cudf/tests/test_orc.py +++ b/python/cudf/cudf/tests/test_orc.py @@ -8,6 +8,7 @@ import pandas as pd import pyarrow as pa import pyarrow.orc +import pyorc import pytest import cudf @@ -318,6 +319,31 @@ def test_orc_read_rows(datadir, skiprows, num_rows): np.testing.assert_allclose(pdf, gdf) +def test_orc_read_skiprows(tmpdir): + fname = tmpdir.join("TestOrcFile.skiprows.orc") + df = pd.DataFrame( + {"a": [1, 0, 1, 0, None, 1, 1, 1, 0, None, 0, 0, 1, 1, 1, 1]}, + dtype=pd.BooleanDtype(), + ) + output = open(fname, "wb") + writer = pyorc.Writer(output, pyorc.Struct(a=pyorc.Boolean())) + tuples = list( + map( + lambda x: (None,) if x[0] is pd.NA else x, + list(df.itertuples(index=False, name=None)), + ) + ) + writer.writerows(tuples) + writer.close() + + skiprows = 10 + + expected = cudf.read_orc(fname)[skiprows::].reset_index(drop=True) + got = cudf.read_orc(fname, skiprows=skiprows) + + assert_eq(expected, got) + + def test_orc_reader_uncompressed_block(datadir): path = datadir / "uncompressed_snappy.orc" try: From bb4a77c0466cbaf943a718c53b7060ea8e7fecf2 Mon Sep 17 00:00:00 2001 From: Ramakrishna Prabhu Date: Tue, 9 Feb 2021 17:01:38 -0600 Subject: [PATCH 2/5] review changes --- python/cudf/cudf/tests/test_orc.py | 9 ++++----- 1 file changed, 4 insertions(+), 5 deletions(-) diff --git a/python/cudf/cudf/tests/test_orc.py b/python/cudf/cudf/tests/test_orc.py index 9681eda6fda..ed91e909f25 100644 --- a/python/cudf/cudf/tests/test_orc.py +++ b/python/cudf/cudf/tests/test_orc.py @@ -320,13 +320,12 @@ def test_orc_read_rows(datadir, skiprows, num_rows): def test_orc_read_skiprows(tmpdir): - fname = tmpdir.join("TestOrcFile.skiprows.orc") + buff = BytesIO() df = pd.DataFrame( {"a": [1, 0, 1, 0, None, 1, 1, 1, 0, None, 0, 0, 1, 1, 1, 1]}, dtype=pd.BooleanDtype(), ) - output = open(fname, "wb") - writer = pyorc.Writer(output, pyorc.Struct(a=pyorc.Boolean())) + writer = pyorc.Writer(buff, pyorc.Struct(a=pyorc.Boolean())) tuples = list( map( lambda x: (None,) if x[0] is pd.NA else x, @@ -338,8 +337,8 @@ def test_orc_read_skiprows(tmpdir): skiprows = 10 - expected = cudf.read_orc(fname)[skiprows::].reset_index(drop=True) - got = cudf.read_orc(fname, skiprows=skiprows) + expected = cudf.read_orc(buff)[skiprows::].reset_index(drop=True) + got = cudf.read_orc(buff, skiprows=skiprows) assert_eq(expected, got) From 706e0081d1ccfcf82234c72637b15f3b0f194e66 Mon Sep 17 00:00:00 2001 From: Ramakrishna Prabhu Date: Wed, 10 Feb 2021 15:55:17 -0600 Subject: [PATCH 3/5] simplifying changes --- cpp/src/io/orc/stripe_data.cu | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/cpp/src/io/orc/stripe_data.cu b/cpp/src/io/orc/stripe_data.cu index 686cf2fd923..f4b1f577038 100644 --- a/cpp/src/io/orc/stripe_data.cu +++ b/cpp/src/io/orc/stripe_data.cu @@ -1213,7 +1213,9 @@ __global__ void __launch_bounds__(block_size) for (uint32_t i = t * 32; i < skippedrows; i += 32 * 32) { uint32_t bits = (i + 32 <= skippedrows) ? s->vals.u32[i >> 5] - : (rle8_read_bool32(s->vals.u32, i) & ((1 << (skippedrows - i)) - 1)); + : (__brev(__byte_perm(s->vals.u32[i >> 5], 0, 0x0123)) & + ((1 << (skippedrows - i)) - + 1)); // Need to arrange the bits to apply mask properly. skip_count += __popc(bits); } skip_count = warp_reduce(temp_storage[t / 32]).Sum(skip_count); From 70641c65917a7db9efdf430511835707630deb05 Mon Sep 17 00:00:00 2001 From: Ramakrishna Prabhu Date: Wed, 10 Feb 2021 17:39:14 -0600 Subject: [PATCH 4/5] review changes to remove __brev --- cpp/src/io/orc/stripe_data.cu | 12 +++++++----- 1 file changed, 7 insertions(+), 5 deletions(-) diff --git a/cpp/src/io/orc/stripe_data.cu b/cpp/src/io/orc/stripe_data.cu index f4b1f577038..3cb1cb52505 100644 --- a/cpp/src/io/orc/stripe_data.cu +++ b/cpp/src/io/orc/stripe_data.cu @@ -1211,11 +1211,13 @@ __global__ void __launch_bounds__(block_size) uint32_t skippedrows = min(static_cast(first_row - row_in), nrows); uint32_t skip_count = 0; for (uint32_t i = t * 32; i < skippedrows; i += 32 * 32) { - uint32_t bits = (i + 32 <= skippedrows) - ? s->vals.u32[i >> 5] - : (__brev(__byte_perm(s->vals.u32[i >> 5], 0, 0x0123)) & - ((1 << (skippedrows - i)) - - 1)); // Need to arrange the bits to apply mask properly. + uint32_t bits = + (i + 32 <= skippedrows) + ? s->vals.u32[i >> 5] + : (__byte_perm(s->vals.u32[i >> 5], + 0, + 0x0123) & // Need to arrange the bytes to apply mask properly. + (0xffffffffu << (0x20 - skippedrows + i))); skip_count += __popc(bits); } skip_count = warp_reduce(temp_storage[t / 32]).Sum(skip_count); From 92e3e0cabf067b2f62c6a6fe83df0df43f2a3428 Mon Sep 17 00:00:00 2001 From: Ramakrishna Prabhu Date: Thu, 11 Feb 2021 10:22:06 -0600 Subject: [PATCH 5/5] review changes to move comment --- cpp/src/io/orc/stripe_data.cu | 11 ++++------- 1 file changed, 4 insertions(+), 7 deletions(-) diff --git a/cpp/src/io/orc/stripe_data.cu b/cpp/src/io/orc/stripe_data.cu index 3cb1cb52505..cf28ee12af7 100644 --- a/cpp/src/io/orc/stripe_data.cu +++ b/cpp/src/io/orc/stripe_data.cu @@ -1211,13 +1211,10 @@ __global__ void __launch_bounds__(block_size) uint32_t skippedrows = min(static_cast(first_row - row_in), nrows); uint32_t skip_count = 0; for (uint32_t i = t * 32; i < skippedrows; i += 32 * 32) { - uint32_t bits = - (i + 32 <= skippedrows) - ? s->vals.u32[i >> 5] - : (__byte_perm(s->vals.u32[i >> 5], - 0, - 0x0123) & // Need to arrange the bytes to apply mask properly. - (0xffffffffu << (0x20 - skippedrows + i))); + // Need to arrange the bytes to apply mask properly. + uint32_t bits = (i + 32 <= skippedrows) ? s->vals.u32[i >> 5] + : (__byte_perm(s->vals.u32[i >> 5], 0, 0x0123) & + (0xffffffffu << (0x20 - skippedrows + i))); skip_count += __popc(bits); } skip_count = warp_reduce(temp_storage[t / 32]).Sum(skip_count);