From 4cbba8f6b222494ec62df254b42415b55bbccd0e Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 1 Feb 2023 16:40:31 +0000 Subject: [PATCH 1/2] Ensure all of device bitmask is initialized in from_arrow 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. --- cpp/src/interop/from_arrow.cu | 17 +++++++++++------ 1 file changed, 11 insertions(+), 6 deletions(-) diff --git a/cpp/src/interop/from_arrow.cu b/cpp/src/interop/from_arrow.cu index 2d4501ec9f7..3d2279c5afe 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,21 @@ struct dispatch_to_cudf_column { if (array.null_bitmap_data() == nullptr) { return std::make_unique(0, stream, mr); } - auto mask = std::make_unique( - bitmask_allocation_size_bytes(static_cast(array.null_bitmap()->size() * CHAR_BIT)), - stream, - mr); + auto null_bitmap_size = array.null_bitmap()->size(); + auto allocation_size = + bitmask_allocation_size_bytes(static_cast(null_bitmap_size * CHAR_BIT)); + auto mask = std::make_unique(allocation_size, stream, mr); auto mask_buffer = array.null_bitmap(); CUDF_CUDA_TRY(cudaMemcpyAsync(mask->data(), reinterpret_cast(mask_buffer->address()), - array.null_bitmap()->size(), + null_bitmap_size, cudaMemcpyDefault, stream.value())); + auto num_zeros = allocation_size - null_bitmap_size; + if (num_zeros > 0) { + auto zero_after = static_cast(mask->data()) + null_bitmap_size; + CUDF_CUDA_TRY(cudaMemsetAsync(zero_after, 0, num_zeros, stream.value())); + } return mask; } From 1dc9e643deacbfe4b030c357899266ca1c2e2f54 Mon Sep 17 00:00:00 2001 From: Lawrence Mitchell Date: Wed, 1 Feb 2023 17:28:01 +0000 Subject: [PATCH 2/2] Minor changes in review --- cpp/src/interop/from_arrow.cu | 13 +++++++------ 1 file changed, 7 insertions(+), 6 deletions(-) diff --git a/cpp/src/interop/from_arrow.cu b/cpp/src/interop/from_arrow.cu index 3d2279c5afe..710f056dc9d 100644 --- a/cpp/src/interop/from_arrow.cu +++ b/cpp/src/interop/from_arrow.cu @@ -104,8 +104,8 @@ struct dispatch_to_cudf_column { if (array.null_bitmap_data() == nullptr) { return std::make_unique(0, stream, mr); } - auto null_bitmap_size = array.null_bitmap()->size(); - auto allocation_size = + auto const null_bitmap_size = array.null_bitmap()->size(); + auto const allocation_size = bitmask_allocation_size_bytes(static_cast(null_bitmap_size * CHAR_BIT)); auto mask = std::make_unique(allocation_size, stream, mr); auto mask_buffer = array.null_bitmap(); @@ -114,10 +114,11 @@ struct dispatch_to_cudf_column { null_bitmap_size, cudaMemcpyDefault, stream.value())); - auto num_zeros = allocation_size - null_bitmap_size; - if (num_zeros > 0) { - auto zero_after = static_cast(mask->data()) + null_bitmap_size; - CUDF_CUDA_TRY(cudaMemsetAsync(zero_after, 0, num_zeros, 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(mask->data()) + null_bitmap_size; + CUDF_CUDA_TRY(cudaMemsetAsync(trailing_bytes, 0, num_trailing_bytes, stream.value())); } return mask; }