Add DELTA_BYTE_ARRAY encoder for Parquet (#15239)

Re-submission of #14938. Final (delta) piece of #13501.

Adds the ability to encode Parquet pages as DELTA_BYTE_ARRAY. Python testing wlll be added as a follow-on when per-column encoding selection is added to the python API (ref this [comment](#15081 (comment))).

Authors:
  - Ed Seidl (https://github.com/etseidl)
  - Vukasin Milovanovic (https://github.com/vuule)
  - Yunsong Wang (https://github.com/PointKernel)

Approvers:
  - Vukasin Milovanovic (https://github.com/vuule)
  - Yunsong Wang (https://github.com/PointKernel)

URL: #15239

cpp/src/io/parquet/page_delta_decode.cu

@@ -462,6 +462,14 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size)
     return;
   }
 
+  if (s->col.logical_type.has_value() && s->col.logical_type->type == LogicalType::DECIMAL) {
+    // we cannot read decimal encoded with DELTA_BYTE_ARRAY yet
+    if (t == 0) {
+      set_error(static_cast<kernel_error::value_type>(decode_error::INVALID_DATA_TYPE), error_code);
+    }
+    return;
+  }
+
   bool const has_repetition = s->col.max_level[level_type::REPETITION] > 0;
 
   // choose a character parallel string copy when the average string is longer than a warp
@@ -620,6 +628,14 @@ CUDF_KERNEL void __launch_bounds__(decode_block_size)
     return;
   }
 
+  if (s->col.logical_type.has_value() && s->col.logical_type->type == LogicalType::DECIMAL) {
+    // we cannot read decimal encoded with DELTA_LENGTH_BYTE_ARRAY yet
+    if (t == 0) {
+      set_error(static_cast<kernel_error::value_type>(decode_error::INVALID_DATA_TYPE), error_code);
+    }
+    return;
+  }
+
   bool const has_repetition = s->col.max_level[level_type::REPETITION] > 0;
 
   // copying logic from gpuDecodePageData.