Read FIXED_LEN_BYTE_ARRAY as binary in parquet reader

cpp/src/io/parquet/page_data.cu

@@ -41,7 +41,7 @@ inline __device__ void gpuOutputString(volatile page_state_s* s,
                                        void* dstv)
 {
   auto [ptr, len] = gpuGetStringData(s, sb, src_pos);
-  if (s->dtype_len == 4) {
+  if (s->dtype_len == 4 and (s->col.data_type & 7) == BYTE_ARRAY) {
     // Output hash. This hash value is used if the option to convert strings to
     // categoricals is enabled. The seed value is chosen arbitrarily.
     uint32_t constexpr hash_seed = 33;
@@ -773,8 +773,10 @@ __global__ void __launch_bounds__(decode_block_size) gpuDecodePageData(
   if (s->dict_base) {
     out_thread0 = (s->dict_bits > 0) ? 64 : 32;
   } else {
-    out_thread0 =
-      ((s->col.data_type & 7) == BOOLEAN || (s->col.data_type & 7) == BYTE_ARRAY) ? 64 : 32;
+    out_thread0 = ((s->col.data_type & 7) == BOOLEAN || (s->col.data_type & 7) == BYTE_ARRAY ||
+                   (s->col.data_type & 7) == FIXED_LEN_BYTE_ARRAY)
+                    ? 64
+                    : 32;
   }
 
   PageNestingDecodeInfo* nesting_info_base = s->nesting_info;
@@ -811,7 +813,8 @@ __global__ void __launch_bounds__(decode_block_size) gpuDecodePageData(
         src_target_pos = gpuDecodeDictionaryIndices<false>(s, sb, src_target_pos, t & 0x1f).first;
       } else if ((s->col.data_type & 7) == BOOLEAN) {
         src_target_pos = gpuDecodeRleBooleans(s, sb, src_target_pos, t & 0x1f);
-      } else if ((s->col.data_type & 7) == BYTE_ARRAY) {
+      } else if ((s->col.data_type & 7) == BYTE_ARRAY or
+                 (s->col.data_type & 7) == FIXED_LEN_BYTE_ARRAY) {
         gpuInitStringDescriptors<false>(s, sb, src_target_pos, t & 0x1f);
       }
       if (t == 32) { *(volatile int32_t*)&s->dict_pos = src_target_pos; }
@@ -881,6 +884,12 @@ __global__ void __launch_bounds__(decode_block_size) gpuDecodePageData(
               }
               break;
           }
+        } else if (dtype == FIXED_LEN_BYTE_ARRAY) {
+          gpuOutputString(s,
+                          sb,
+                          val_src_pos,
+                          nesting_info_base[leaf_level_index].data_out +
+                            static_cast<size_t>(dst_pos) * sizeof(string_index_pair));
         } else if (dtype == INT96) {
           gpuOutputInt96Timestamp(s, sb, val_src_pos, static_cast<int64_t*>(dst));
         } else if (dtype_len == 8) {

cpp/src/io/parquet/page_decode.cuh

@@ -418,12 +418,20 @@ __device__ size_type gpuInitStringDescriptors(page_state_s volatile* s,
 
     while (pos < target_pos) {
       int len;
-      if (k + 4 <= dict_size) {
-        len = (cur[k]) | (cur[k + 1] << 8) | (cur[k + 2] << 16) | (cur[k + 3] << 24);
-        k += 4;
-        if (k + len > dict_size) { len = 0; }
+      if ((s->col.data_type & 7) == FIXED_LEN_BYTE_ARRAY) {
+        if (k < dict_size) {
+          len = s->dtype_len_in;
+        } else {
+          len = 0;
+        }
       } else {
-        len = 0;
+        if (k + 4 <= dict_size) {
+          len = (cur[k]) | (cur[k + 1] << 8) | (cur[k + 2] << 16) | (cur[k + 3] << 24);
+          k += 4;
+          if (k + len > dict_size) { len = 0; }
+        } else {
+          len = 0;
+        }
       }
       if constexpr (!sizes_only) {
         sb->dict_idx[rolling_index(pos)] = k;
@@ -1198,7 +1206,10 @@ inline __device__ bool setupLocalPageInfo(page_state_s* const s,
               uint32_t len = idx < max_depth - 1 ? sizeof(cudf::size_type) : s->dtype_len;
               // if this is a string column, then dtype_len is a lie. data will be offsets rather
               // than (ptr,len) tuples.
-              if (data_type == BYTE_ARRAY && s->dtype_len != 4) { len = sizeof(cudf::size_type); }
+              if ((data_type == BYTE_ARRAY && s->dtype_len != 4) ||
+                  data_type == FIXED_LEN_BYTE_ARRAY) {
+                len = sizeof(cudf::size_type);
+              }
               nesting_info->data_out += (output_offset * len);
             }
             if (nesting_info->string_out != nullptr) {

cpp/src/io/parquet/page_string_decode.cu

@@ -487,6 +487,36 @@ __device__ size_t totalDictEntriesSize(uint8_t const* data,
   return sum_l;
 }
 
+/**
+ * @brief Compute string size information for fixed len byte array strings.
+ *
+ * @param input_size Size of each fixed length byte array
+ * @param data_size Length of data
+ * @param start_value Do not count values that occur before this index
+ * @param end_value Do not count values that occur after this index
+ */
+__device__ size_t totalFixedLenByteArraySize(int input_size,
+                                             int data_size,
+                                             int start_value,
+                                             int end_value)
+{
+  int const t      = threadIdx.x;
+  int pos          = 0;
+  size_t total_len = 0;
+
+  // This step is purely serial
+  if (!t) {
+    int k = 0;
+    while (pos < end_value && k < data_size) {
+      k += input_size;
+      if (pos >= start_value) { total_len += input_size; }  // TODO not sure this is necessary
+      pos++;
+    }
+  }
+
+  return total_len;
+}
+
 /**
  * @brief Compute string size information for plain encoded strings.
  *
@@ -625,8 +655,12 @@ __global__ void __launch_bounds__(preprocess_block_size) gpuComputePageStringSiz
       break;
     case Encoding::PLAIN:
       dict_size = static_cast<int32_t>(end - data);
-      str_bytes = is_bounds_pg ? totalPlainEntriesSize(data, dict_size, start_value, end_value)
-                               : dict_size - sizeof(int) * (pp->num_input_values - pp->num_nulls);
+      if ((s->col.data_type & 7) == FIXED_LEN_BYTE_ARRAY) {
+        str_bytes = totalFixedLenByteArraySize(s->dtype_len_in, dict_size, start_value, end_value);
+      } else {
+        str_bytes = is_bounds_pg ? totalPlainEntriesSize(data, dict_size, start_value, end_value)
+                                 : dict_size - sizeof(int) * (pp->num_input_values - pp->num_nulls);
+      }
       break;
   }
 

cpp/src/io/parquet/parquet_gpu.hpp

@@ -429,8 +429,9 @@ struct EncPage {
  */
 constexpr bool is_string_col(ColumnChunkDesc const& chunk)
 {
-  return (chunk.data_type & 7) == BYTE_ARRAY and (chunk.data_type >> 3) != 4 and
-         chunk.converted_type != DECIMAL;
+  return ((chunk.data_type & 7) == BYTE_ARRAY and (chunk.data_type >> 3) != 4 and
+          chunk.converted_type != DECIMAL) or
+         (chunk.data_type & 7) == FIXED_LEN_BYTE_ARRAY;
 }
 
 /**