8000 Fix debrotli issue on CUDA 11.5 by vuule · Pull Request #9632 · rapidsai/cudf · GitHub
[go: up one dir, main page]
More Web Proxy on the site http://driver.im/
Skip to content

Fix debrotli issue on CUDA 11.5 #9632

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 10 commits into from
Nov 11, 2021
Merged

Fix debrotli issue on CUDA 11.5 #9632

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
10 commits
Select commit Hold shift + click to select a range
d1ff133
the workarond
vuule Nov 8, 2021
bf25a00
re-enable test
vuule Nov 8, 2021
3af13af
fix via UB removal
vuule Nov 9, 2021
00234fb
Merge branch 'branch-21.12' of https://github.com/rapidsai/cudf into …
vuule Nov 9, 2021
ba1621c
Merge branch 'branch-21.12' of https://github.com/rapidsai/cudf into …
vuule Nov 10, 2021
2b0838d
fix thrust::sequence bounds
vuule Nov 10, 2021
eeb55ca
add test for sequece boundary issue
vuule Nov 10, 2021
81f8299
add missing test file
vuule Nov 10, 2021
9b0a8fe
revert to minimal UB removal + sync
vuule Nov 11, 2021
725f9d9
add alignment
vuule Nov 11, 2021
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
16 changes: 11 additions & 5 deletions cpp/src/io/comp/debrotli.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -193,7 +193,7 @@ struct debrotli_state_s {
uint16_t* block_type_vlc[3];
huff_scratch_s hs;
uint32_t mtf[65];
uint64_t heap[local_heap_size / 8];
__align__(8) char heap[local_heap_size];
};

inline __device__ uint32_t Log2Floor(uint32_t value) { return 32 - __clz(value); }
Expand Down Expand Up @@ -307,11 +307,13 @@ static __device__ uint32_t getvlc(debrotli_state_s* s, const uint16_t* lut)
return vlc;
}

static auto __device__ allocation_size(uint32_t bytes) { return (bytes + 7) & ~7; }

/// Alloc bytes from the local (shared mem) heap
static __device__ uint8_t* local_alloc(debrotli_state_s* s, uint32_t bytes)
{
int heap_used = s->heap_used;
int len = (bytes + 7) >> 3;
int heap_used = s->heap_used;
auto const len = allocation_size(bytes);
if (heap_used + len <= s->heap_limit) {
uint8_t* ptr = reinterpret_cast<uint8_t*>(&s->heap[heap_used]);
s->heap_used = (uint16_t)(heap_used + len);
Expand All @@ -327,7 +329,7 @@ static __device__ uint8_t* local_heap_shrink(debrotli_state_s* s, uint32_t bytes
{
int heap_used = s->heap_used;
int heap_limit = s->heap_limit;
int len = (bytes + 7) >> 3;
auto const len = allocation_size(bytes);
if (heap_limit - len >= heap_used) {
heap_limit -= len;
s->heap_limit = (uint16_t)heap_limit;
Expand All @@ -339,7 +341,7 @@ static __device__ uint8_t* local_heap_shrink(debrotli_state_s* s, uint32_t bytes

static __device__ void local_heap_grow(debrotli_state_s* s, uint32_t bytes)
{
int len = (bytes + 7) >> 3;
auto const len = allocation_size(bytes);
int heap_limit = s->heap_limit + len;
s->heap_limit = (uint16_t)heap_limit;
}
Expand Down Expand Up @@ -1885,6 +1887,10 @@ static __device__ void ProcessCommands(debrotli_state_s* s, const brotli_diction
pos += copy_length;
}
}

// Ensure all other threads have observed prior state of p1 & p2 before overwriting
__syncwarp();

if (!t) {
s->p1 = (uint8_t)p1;
s->p2 = (uint8_t)p2;
Expand Down
Binary file added python/cudf/cudf/tests/data/parquet/brotli_int16.parquet
View file
Open in desktop
Binary file not shown.
17 changes: 9 additions & 8 deletions python/cudf/cudf/tests/test_parquet.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -17,8 +17,6 @@
from packaging import version
from pyarrow import fs as pa_fs, parquet as pq

import rmm

import cudf
from cudf.io.parquet import ParquetWriter, merge_parquet_filemetadata
from cudf.testing import dataset_generator as dg
Expand Down Expand Up @@ -246,12 +244,6 @@ def _make_parquet_path_or_buf(src):
],
)
def test_parquet_reader_basic(parquet_file, columns, engine):
if rmm._cuda.gpu.runtimeGetVersion() == 11050 and "brotli_" in str(
parquet_file
):
pytest.xfail(
"Known issue: https://github.com/rapidsai/cudf/issues/9546"
)
expect = pd.read_parquet(parquet_file, columns=columns)
got = cudf.read_parquet(parquet_file, engine=engine, columns=columns)
if len(expect) == 0:
Expand Down Expand Up @@ -2139,3 +2131,12 @@ def test_parquet_decimal_precision_empty(tmpdir):
df.to_parquet(fname)
df = cudf.read_parquet(fname)
assert df.val.dtype.precision == 5


def test_parquet_reader_brotli(datadir):
fname = datadir / "brotli_int16.parquet"

expect = pd.read_parquet(fname)
got = cudf.read_parquet(fname).to_pandas(nullable=True)

assert_eq(expect, got)
0