[WIP] Fix orc reader RLEv2 reader

CHANGELOG.md

@@ -95,6 +95,7 @@
 - PR #5446 Fix compile error caused by out-of-date PR merge (4990)
 - PR #5459 Fix str.translate to convert table characters to UTF-8
 - PR #5465 Fix benchmark out of memory errors due to multiple initialization
+- PR #5473 Fix RLEv2 patched base in ORC reader
 - PR #5472 Fix str concat issue with indexed series
 
 

cpp/src/io/orc/stripe_data.cu

@@ -642,6 +642,20 @@ static const __device__ __constant__ uint8_t kRLEv2_W[32] = {
   1,  2,  3,  4,  5,  6,  7,  8,  9,  10, 11, 12, 13, 14, 15, 16,
   17, 18, 19, 20, 21, 22, 23, 24, 26, 28, 30, 32, 40, 48, 56, 64};
 
+/**
+ * @brief Maps the RLEv2 patch size (pw + pgw) to number of bits
+ *
+ * Patch size (in bits) is only allowed to be from the below set. If `pw + pgw == 34` then the size
+ * of the patch in the file is the smallest size in the set that can fit 34 bits i.e.
+ * `ClosestFixedBitsMap[34] == 40`
+ *
+ * @see https://github.com/apache/orc/commit/9faf7f5147a7bc69
+ */
+static const __device__ __constant__ uint8_t ClosestFixedBitsMap[65] = {
+  1,  1,  2,  3,  4,  5,  6,  7,  8,  9,  10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21,
+  22, 23, 24, 26, 26, 28, 28, 30, 30, 32, 32, 40, 40, 40, 40, 40, 40, 40, 40, 48, 48, 48,
+  48, 48, 48, 48, 48, 56, 56, 56, 56, 56, 56, 56, 56, 64, 64, 64, 64, 64, 64, 64, 64};
+
 /**
  * @brief ORC Integer RLEv2 decoding
  *
@@ -812,7 +826,7 @@ static __device__ uint32_t Integer_RLEv2(
           bytestream_readbe(bs, pos * 8 + i * w, w, v);
           vals[base + i] = v;
         } else if (mode == 2) {
-          uint32_t ofs   = bytestream_readbits64(bs, pos * 8 + i * w, w);
+          uint64_t ofs   = bytestream_readbits64(bs, pos * 8 + i * w, w);
           vals[base + i] = rle->baseval.u64[r] + ofs;
         } else {
           int64_t delta = rle->delta[r], ofs;
@@ -834,9 +848,10 @@ static __device__ uint32_t Integer_RLEv2(
       uint32_t pgw      = 1 + ((pw_byte3 >> 5) & 7);  // patch gap width, 1 to 8 bits
       uint32_t pll      = pw_byte3 & 0x1f;            // patch list length
       if (pll != 0) {
+        uint32_t pgw_pw_len = ClosestFixedBitsMap[pw + pgw];
         uint64_t patch_pos64 =
           (tr < pll)
-            ? bytestream_readbits64(bs, pos * 8 + ((n * w + 7) & ~7) + tr * (pgw + pw), pgw + pw)
+            ? bytestream_readbits64(bs, pos * 8 + ((n * w + 7) & ~7) + tr * (pgw + pw), pgw_pw_len)
             : 0;
         uint32_t patch_pos;
         T patch = 1;

python/cudf/cudf/tests/test_orc.py

@@ -67,6 +67,7 @@ def _make_path_or_buf(src):
                 "double1",
             ],
         ),
+        ("TestOrcFile.RLEv2.orc", ["x", "y"]),
         ("TestOrcFile.testSnappy.orc", None),
         ("TestOrcFile.demo-12-zlib.orc", ["_col2", "_col3", "_col4", "_col5"]),
     ],