Fix debrotli issue on CUDA 11.5 #9632

vuule · 2021-11-08T19:57:03Z

This PR fixes the issue likely through elimination of undefined behavior.
Modified local heap implementation to return void* instead on uint8_t*. This greatly reduces the number of reinterpret_casts. Also changed heap type to char*, presumably reducing/eliminating aliasing issues.
Some other clean up in related code included.

galipremsagar · 2021-11-08T21:24:03Z

I can confirm that this fixes issue in both cuda 11.4 & 11.5 on T4 gpus.

vuule · 2021-11-08T21:30:20Z

rerun tests

codecov · 2021-11-08T21:39:58Z

Codecov Report

Merging #9632 (e13c2ec) into branch-21.12 (ab4bfaa) will decrease coverage by 0.10%.
The diff coverage is n/a.

❗ Current head e13c2ec differs from pull request most recent head 725f9d9. Consider uploading reports for the commit 725f9d9 to get more accurate results

@@               Coverage Diff                @@
##           branch-21.12    #9632      +/-   ##
================================================
- Coverage         10.79%   10.68%   -0.11%     
================================================
  Files               116      117       +1     
  Lines             18869    19872    +1003     
================================================
+ Hits               2036     2123      +87     
- Misses            16833    17749     +916

Impacted Files	Coverage Δ
python/dask_cudf/dask_cudf/sorting.py	`92.90% <0.00%> (-1.21%)`	⬇️
python/cudf/cudf/io/csv.py	`0.00% <0.00%> (ø)`
python/cudf/cudf/io/hdf.py	`0.00% <0.00%> (ø)`
python/cudf/cudf/io/orc.py	`0.00% <0.00%> (ø)`
python/cudf/cudf/__init__.py	`0.00% <0.00%> (ø)`
python/cudf/cudf/_version.py	`0.00% <0.00%> (ø)`
python/cudf/cudf/core/abc.py	`0.00% <0.00%> (ø)`
python/cudf/cudf/api/types.py	`0.00% <0.00%> (ø)`
python/cudf/cudf/io/dlpack.py	`0.00% <0.00%> (ø)`
python/cudf/cudf/core/frame.py	`0.00% <0.00%> (ø)`
... and 67 more

Continue to review full report at Codecov.

Legend - Click here to learn more
Δ = absolute <relative> (impact), ø = not affected, ? = missing data
Powered by Codecov. Last update d60e2e6...725f9d9. Read the comment docs.

galipremsagar · 2021-11-09T05:09:36Z

rerun tests

…bug-debrotli-115

vuule · 2021-11-09T20:37:32Z

@galipremsagar could you run a round of fuzz tests with brotli to confirm the fix?

davidwendt · 2021-11-09T22:16:05Z

cpp/src/io/comp/debrotli.cu

-  volatile uint32_t* heap_ptr = reinterpret_cast<volatile uint32_t*>(ext_heap_base);
-  uint32_t first_free_block   = ~0;
+  auto const len            = (bytes + 0xf) & ~0xf;
+  auto const heap_ptr       = static_cast<volatile uint32_t*>(ext_heap_base);


This is const and volatile at the same time?

The pointer is const, and the data it points to is volatile.
I can avoid auto here so the type becomes uint32_t volatile* const, which reads fine (right to left).

Ok, that makes sense. I'd rather see the const on the left of the equal sign.

cpp/src/io/comp/debrotli.cu

galipremsagar · 2021-11-10T11:49:37Z

@vuule In fuzz-testing we ran into a data-corruption issue with brotli encoding in reader:
test_data.zip - Unzip this file, it will contain two parquet files. We will use only cpu_pdf.parquet which is brotli encoded. The gpu_pdf.parquet is not encoded.

>>> import cudf
>>> actual = cudf.read_parquet('cpu_pdf.parquet')
>>> expected = pd.read_parquet('cpu_pdf.parquet')
(Pdb) actual['21'] # See the last value, it should be 27195 instead of 0.
0       -28871
1         <NA>
2        18224
3       -13182
4         <NA>
         ...  
9426      6841
9427      3473
9428     31914
9429    -22488
9430         0
Name: 21, Length: 9431, dtype: int16
(Pdb) expected['21']
0       -28871
1         <NA>
2        18224
3       -13182
4         <NA>
         ...  
9426      6841
9427      3473
9428     31914
9429    -22488
9430     27195
Name: 21, Length: 9431, dtype: Int16

(Pdb) expected['21'][actual['21'].to_pandas(nullable=True) != expected['21']] # 1339 out of 9431 values are corrupted.
6571    -25740
6573      6433
6575    -21601
6576     19348
6577      4142
         ...  
9418     28837
9420      4155
9423      8501
9424     28381
9430     27195
Name: 21, Length: 1339, dtype: Int16

charlesbluca

Approving python changes

…bug-debrotli-115

ttnghia · 2021-11-11T03:50:01Z

Rerun tests.

galipremsagar · 2021-11-11T14:53:27Z

Verified with fuzz-testing, no issues found. 🎉

vuule · 2021-11-11T16:17:53Z

@gpucibot merge

vuule added 2 commits November 8, 2021 11:40

the workarond

d1ff133

re-enable test

bf25a00

vuule added bug Something isn't working cuIO cuIO issue non-breaking Non-breaking change labels Nov 8, 2021

vuule self-assigned this Nov 8, 2021

github-actions bot added Python Affects Python cuDF API. libcudf Affects libcudf (C++/CUDA) code. labels Nov 8, 2021

galipremsagar approved these changes Nov 8, 2021

View reviewed changes

vuule added 2 commits November 9, 2021 12:23

fix via UB removal

3af13af

Merge branch 'branch-21.12' of https://github.com/rapidsai/cudf into …

00234fb

…bug-debrotli-115

vuule marked this pull request as ready for review November 9, 2021 22:05

vuule requested review from a team as code owners November 9, 2021 22:05

vuule requested review from robertmaynard, davidwendt, isVoid and charlesbluca November 9, 2021 22:05

davidwendt reviewed Nov 9, 2021

View reviewed changes

robertmaynard reviewed Nov 9, 2021

View reviewed changes

cpp/src/io/comp/debrotli.cu Outdated Show resolved Hide resolved

davidwendt approved these changes Nov 9, 2021

View reviewed changes

robertmaynard approved these changes Nov 9, 2021

View reviewed changes

charlesbluca approved these changes Nov 10, 2021

View reviewed changes

vuule added the 5 - DO NOT MERGE Hold off on merging; see PR for details label Nov 10, 2021

vuule added 3 commits November 10, 2021 10:40

Merge branch 'branch-21.12' of https://github.com/rapidsai/cudf into …

ba1621c

…bug-debrotli-115

fix thrust::sequence bounds

2b0838d

add test for sequece boundary issue

eeb55ca

vuule removed the 5 - DO NOT MERGE Hold off on merging; see PR for details label Nov 10, 2021

vuule added 3 commits November 10, 2021 15:35

add missing test file

81f8299

revert to minimal UB removal + sync

9b0a8fe

add alignment

725f9d9

vuule added the 5 - DO NOT MERGE Hold off on merging; see PR for details label Nov 11, 2021

galipremsagar approved these changes Nov 11, 2021

View reviewed changes

vuule added 5 - Ready to Merge Testing and reviews complete, ready to merge and removed 5 - DO NOT MERGE Hold off on merging; see PR for details labels Nov 11, 2021

rapids-bot bot merged commit 507a3c5 into rapidsai:branch-21.12 Nov 11, 2021

vuule deleted the bug-debrotli-115 branch April 20, 2022 23:47

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Fix debrotli issue on CUDA 11.5 #9632

Fix debrotli issue on CUDA 11.5 #9632

vuule commented Nov 8, 2021 •

edited

Loading

galipremsagar commented Nov 8, 2021

vuule commented Nov 8, 2021

codecov bot commented Nov 8, 2021 •

edited

Loading

galipremsagar commented Nov 9, 2021

vuule commented Nov 9, 2021

davidwendt Nov 9, 2021

vuule Nov 9, 2021

davidwendt Nov 9, 2021 •

edited

Loading

galipremsagar commented Nov 10, 2021 •

edited

Loading

charlesbluca left a comment

ttnghia commented Nov 11, 2021

galipremsagar commented Nov 11, 2021

vuule commented Nov 11, 2021

Fix debrotli issue on CUDA 11.5 #9632

Fix debrotli issue on CUDA 11.5 #9632

Conversation

vuule commented Nov 8, 2021 • edited Loading

galipremsagar commented Nov 8, 2021

vuule commented Nov 8, 2021

codecov bot commented Nov 8, 2021 • edited Loading

Codecov Report

galipremsagar commented Nov 9, 2021

vuule commented Nov 9, 2021

davidwendt Nov 9, 2021

Choose a reason for hiding this comment

vuule Nov 9, 2021

Choose a reason for hiding this comment

davidwendt Nov 9, 2021 • edited Loading

Choose a reason for hiding this comment

galipremsagar commented Nov 10, 2021 • edited Loading

charlesbluca left a comment

Choose a reason for hiding this comment

ttnghia commented Nov 11, 2021

galipremsagar commented Nov 11, 2021

vuule commented Nov 11, 2021

vuule commented Nov 8, 2021 •

edited

Loading

codecov bot commented Nov 8, 2021 •

edited

Loading

davidwendt Nov 9, 2021 •

edited

Loading

galipremsagar commented Nov 10, 2021 •

edited

Loading