From 0723f3f34a56b62a10e1715b95e807cd88a9a28d Mon Sep 17 00:00:00 2001
From: Lawrence Mitchell <wence@gmx.li>
Date: Mon, 13 Mar 2023 17:29:32 +0000
Subject: [PATCH] Ensure all of device bitmask is initialized in from_arrow
 (#12668)

libcudf bitmasks are allocated to a multiple of 64 bytes, in contrast the arrow spec only requires of a column with a null mask that "the validity bitmap must be large enough to have at least 1 bit for each array slot". When the number of rows is not a multiple of 64, the trailing part of the device allocation (which doesn't contribute to actually masking anything) is left uninitialized. While probably benign, this produces errors when running with compute-sanitizer in initcheck mode (since those data are touched and _are_ uninitialized).

To fix this, memset the trailing allocation to zero.

Closes #8873.

Authors:
  - Lawrence Mitchell (https://github.com/wence-)
  - Mark Harris (https://github.com/harrism)

Approvers:
  - David Wendt (https://github.com/davidwendt)
  - Bradley Dice (https://github.com/bdice)

URL: https://github.com/rapidsai/cudf/pull/12668
---
 cpp/src/interop/from_arrow.cu | 18 ++++++++++++------
 1 file changed, 12 insertions(+), 6 deletions(-)

diff --git a/cpp/src/interop/from_arrow.cu b/cpp/src/interop/from_arrow.cu
index 2d4501ec9f7..710f056dc9d 100644
--- a/cpp/src/interop/from_arrow.cu
+++ b/cpp/src/interop/from_arrow.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -104,16 +104,22 @@ struct dispatch_to_cudf_column {
     if (array.null_bitmap_data() == nullptr) {
       return std::make_unique<rmm::device_buffer>(0, stream, mr);
     }
-    auto mask = std::make_unique<rmm::device_buffer>(
-      bitmask_allocation_size_bytes(static_cast<size_type>(array.null_bitmap()->size() * CHAR_BIT)),
-      stream,
-      mr);
+    auto const null_bitmap_size = array.null_bitmap()->size();
+    auto const allocation_size =
+      bitmask_allocation_size_bytes(static_cast<size_type>(null_bitmap_size * CHAR_BIT));
+    auto mask        = std::make_unique<rmm::device_buffer>(allocation_size, stream, mr);
     auto mask_buffer = array.null_bitmap();
     CUDF_CUDA_TRY(cudaMemcpyAsync(mask->data(),
                                   reinterpret_cast<const uint8_t*>(mask_buffer->address()),
-                                  array.null_bitmap()->size(),
+                                  null_bitmap_size,
                                   cudaMemcpyDefault,
                                   stream.value()));
+    // Zero-initialize trailing padding bytes
+    auto const num_trailing_bytes = allocation_size - null_bitmap_size;
+    if (num_trailing_bytes > 0) {
+      auto trailing_bytes = static_cast<uint8_t*>(mask->data()) + null_bitmap_size;
+      CUDF_CUDA_TRY(cudaMemsetAsync(trailing_bytes, 0, num_trailing_bytes, stream.value()));
+    }
     return mask;
   }