Skip to content

Commit 30446ea

Browse files
TanishGudiseclaude0ax1
authored
fix(cuda): bound dictionary gathers by the dictionary length (#10915)
## Summary - Tracking Issue: #10914 The standalone dictionary kernels in `dict.cu` gathered `values[codes[i]]` for every row. Nullable codes can hold any value at null positions (`DictArray::try_new` doesn't bound-check them), so these gathers read outside the dictionary. The fused planner already rejects nullable-code dictionaries for this reason and sends them to these kernels. Develop's own fuzzer, extended here, crashes on this within 71k runs. ## Changes - `dict.cu`: every kernel (primitive, decimal, bool and the 16-byte view gather) takes `values_len`. A code at or beyond it writes a zero value. In a valid array only null positions can hold such a code, so the zero is never observed. - `dict.rs` and `benches/dynamic_dispatch_cuda.rs`: pass `values_len` at every launch. - Tests (12), each compared with the expected result through `assert_arrays_eq`: - out-of-range codes under nulls for string, decimal, bool and primitive values, with codes up to `u32::MAX` and `u64::MAX`; - the dictionary the fuzzer crashed on (one u8 value, null code 3,960,089,894); - device-resident codes with nullable values and codes; - all-null codes; - slices across 1,023/1,024/1,025 and 2,047/2,048; - u16 codes with 65,536 values and u32 codes with 70,000; - empty, 12-, 13- and 4,097-byte strings; - non-UTF-8 binary values. - `fuzz/src/gpu/mod.rs`: `compress_gpu` also generates string and binary dictionaries, dictionaries with arbitrary u8/u16/u32 codes under nulls, and strings compressed with the CUDA preset. Evidence (A40, CUDA 13.0): | Check | develop's kernels | this PR | |---|---|---| | 5 bug tests: out-of-range codes under nulls (string, decimal, bool, primitive) and the fuzzer's input | 5 fail with `CUDA_ERROR_ILLEGAL_ADDRESS` | 5 pass | | 7 edge-case tests (all-null codes, block-boundary slices, full u16/u32 ranges, inline boundary, binary, device codes) | pass | 7 pass | | `compute-sanitizer --tool memcheck` on the 6 out-of-range tests | 1-2 errors each, `Invalid __global__ read` of 1, 8 or 16 bytes | 0 errors | | `compress_gpu` fuzz (CI flags, ASAN, empty corpus) | crash after 70,941 runs, `CUDA_ERROR_ILLEGAL_ADDRESS` | 100,000 runs in 283 s, 0 crashes | | `vortex-cuda` lib suite | — | 719 passed, 0 failed | Under nsys, the bounds check costs nothing measurable on numeric dictionaries: 1.574 vs 1.563 ms per iteration on CMSprovider and 0.825 vs 0.832 ms on taxi. On develop's kernels the `u32::MAX`/`u64::MAX` tests leave the CUDA context in a sticky `CUDA_ERROR_ILLEGAL_ADDRESS` state. Run in one process, they also fail the unrelated tests that follow, about 336 of them. A regression of this fix would show up the same way, as a wall of failures led by these tests. Everything in the table ran on this PR alone, except the develop fuzz crash: that ran develop's kernels under this PR's extended fuzz target. ## AI assistance Developed with Claude Code (Claude Opus 5.5), which wrote most of the code, tests and this description under my direction. I reviewed the diff and the measurements. 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Signed-off-by: Tanish Gudise <[email protected]> Co-authored-by: Claude Opus 5.5 <[email protected]> Co-authored-by: Alexander Droste <[email protected]>
1 parent 500f8b1 commit 30446ea

4 files changed

Lines changed: 334 additions & 8 deletions

File tree

‎fuzz/src/gpu/mod.rs‎

Lines changed: 84 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -13,9 +13,17 @@ use vortex_array::ArrayRef;
1313
use vortex_array::Canonical;
1414
use vortex_array::IntoArray;
1515
use vortex_array::VortexSessionExecute;
16+
use vortex_array::arrays::DictArray;
17+
use vortex_array::arrays::PrimitiveArray;
18+
use vortex_array::arrays::arbitrary::ArbitraryArray;
19+
use vortex_array::arrays::arbitrary::ArbitraryArrayConfig;
20+
use vortex_array::arrays::arbitrary::ArbitraryWith;
1621
use vortex_array::arrays::dict::ArbitraryDictArray;
22+
use vortex_array::dtype::DType;
1723
use vortex_array::dtype::Nullability;
1824
use vortex_array::dtype::PType;
25+
use vortex_array::validity::Validity;
26+
use vortex_buffer::Buffer;
1927
use vortex_error::vortex_err;
2028

2129
use crate::SESSION;
@@ -28,13 +36,16 @@ use crate::error::VortexFuzzResult;
2836
pub enum GpuEncodingKind {
2937
/// Dictionary encoding with GPU take support.
3038
Dict,
39+
/// Dictionary whose nullable codes hold arbitrary values, possibly outside the dictionary,
40+
/// under null positions.
41+
DictGarbageNullCodes,
3142
}
3243

3344
impl<'a> Arbitrary<'a> for GpuEncodingKind {
3445
fn arbitrary(u: &mut Unstructured<'a>) -> Result<Self> {
35-
// Currently only Dict is supported
36-
match u.int_in_range(0..=0)? {
46+
match u.int_in_range(0..=1)? {
3747
0 => Ok(GpuEncodingKind::Dict),
48+
1 => Ok(GpuEncodingKind::DictGarbageNullCodes),
3849
_ => unreachable!(),
3950
}
4051
}
@@ -58,14 +69,83 @@ impl<'a> Arbitrary<'a> for FuzzCompressGpu {
5869
.0
5970
.into_array()
6071
}
72+
GpuEncodingKind::DictGarbageNullCodes => arbitrary_dict_with_garbage_null_codes(u)?,
6173
};
6274

6375
Ok(FuzzCompressGpu { array })
6476
}
6577
}
6678

79+
/// Generate a random UTF-8 or binary DType.
80+
fn arbitrary_gpu_string_dtype(u: &mut Unstructured) -> Result<DType> {
81+
let nullability: Nullability = u.arbitrary()?;
82+
Ok(if u.arbitrary()? {
83+
DType::Utf8(nullability)
84+
} else {
85+
DType::Binary(nullability)
86+
})
87+
}
88+
89+
/// Generate a dictionary whose codes are in range at valid positions and arbitrary, including far
90+
/// outside the dictionary, at null positions.
91+
fn arbitrary_dict_with_garbage_null_codes(u: &mut Unstructured) -> Result<ArrayRef> {
92+
let dtype = if u.arbitrary()? {
93+
arbitrary_gpu_string_dtype(u)?
94+
} else {
95+
arbitrary_gpu_primitive_dtype(u)?
96+
};
97+
let values_len = u.int_in_range(1..=20)?;
98+
let values = ArbitraryArray::arbitrary_with_config(
99+
u,
100+
&ArbitraryArrayConfig {
101+
dtype: Some(dtype),
102+
len: values_len..=values_len,
103+
},
104+
)?
105+
.0;
106+
let len = u.int_in_range(0..=200)?;
107+
let valid: Vec<bool> = (0..len).map(|_| u.arbitrary()).collect::<Result<_>>()?;
108+
let values_len = u32::try_from(values_len).map_err(|_| arbitrary::Error::IncorrectFormat)?;
109+
let validity: Validity = valid.iter().copied().collect();
110+
let codes = match u.int_in_range(0..=2)? {
111+
0 => PrimitiveArray::new(
112+
Buffer::from(garbage_codes::<u8>(u, &valid, values_len)?),
113+
validity,
114+
),
115+
1 => PrimitiveArray::new(
116+
Buffer::from(garbage_codes::<u16>(u, &valid, values_len)?),
117+
validity,
118+
),
119+
_ => PrimitiveArray::new(
120+
Buffer::from(garbage_codes::<u32>(u, &valid, values_len)?),
121+
validity,
122+
),
123+
};
124+
let dict = DictArray::try_new(codes.into_array(), values)
125+
.map_err(|_| arbitrary::Error::IncorrectFormat)?;
126+
Ok(dict.into_array())
127+
}
128+
129+
/// Codes in `0..values_len` at valid positions and any value of `T` at null positions.
130+
fn garbage_codes<'a, T>(u: &mut Unstructured<'a>, valid: &[bool], values_len: u32) -> Result<Vec<T>>
131+
where
132+
T: Arbitrary<'a> + TryFrom<u32>,
133+
{
134+
valid
135+
.iter()
136+
.map(|&is_valid| {
137+
if is_valid {
138+
T::try_from(u.int_in_range(0..=values_len - 1)?)
139+
.map_err(|_| arbitrary::Error::IncorrectFormat)
140+
} else {
141+
u.arbitrary::<T>()
142+
}
143+
})
144+
.collect()
145+
}
146+
67147
/// Generate a random primitive DType suitable for GPU operations.
68-
fn arbitrary_gpu_primitive_dtype(u: &mut Unstructured) -> Result<vortex_array::dtype::DType> {
148+
fn arbitrary_gpu_primitive_dtype(u: &mut Unstructured) -> Result<DType> {
69149
let nullability: Nullability = u.arbitrary()?;
70150
let ptype = match u.int_in_range(0..=9)? {
71151
0 => PType::U8,
@@ -80,7 +160,7 @@ fn arbitrary_gpu_primitive_dtype(u: &mut Unstructured) -> Result<vortex_array::d
80160
9 => PType::F64,
81161
_ => unreachable!(),
82162
};
83-
Ok(vortex_array::dtype::DType::Primitive(ptype, nullability))
163+
Ok(DType::Primitive(ptype, nullability))
84164
}
85165

86166
/// Run the GPU decompression fuzzer.

‎vortex-cuda/benches/dynamic_dispatch_cuda.rs‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -625,9 +625,11 @@ mod standalone {
625625
{
626626
let mut launch = cuda_ctx.stream().launch_builder(&self.dict_fn);
627627
let codes_len = self.codes_len as u64;
628+
let values_len = self.values_len as u64;
628629
launch.arg(&self.codes_buf);
629630
launch.arg(&codes_len);
630631
launch.arg(&self.values_alp_buf);
632+
launch.arg(&values_len);
631633
launch.arg(&mut self.output_buf);
632634
unsafe {
633635
launch

‎vortex-cuda/kernels/src/dict.cu‎

Lines changed: 11 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -8,10 +8,14 @@
88
#include "config.cuh"
99
#include "types.cuh"
1010

11+
// Codes under null positions may hold any value, and so may corrupt input, so every gather checks
12+
// the code against the dictionary length. A code outside the dictionary writes a zero value instead
13+
// of reading past the end of `values`.
1114
template <typename ValueT, typename IndexT>
1215
__device__ void dict_kernel(const IndexT *const __restrict codes,
1316
uint64_t codes_len,
1417
const ValueT *const __restrict values,
18+
uint64_t values_len,
1519
ValueT *const __restrict output) {
1620
const uint32_t elements_per_block = blockDim.x * ELEMENTS_PER_THREAD;
1721
const uint64_t block_start = static_cast<uint64_t>(blockIdx.x) * elements_per_block;
@@ -20,7 +24,7 @@ __device__ void dict_kernel(const IndexT *const __restrict codes,
2024

2125
for (uint64_t idx = block_start + threadIdx.x; idx < block_end; idx += blockDim.x) {
2226
const IndexT code = codes[idx];
23-
output[idx] = values[code];
27+
output[idx] = static_cast<uint64_t>(code) < values_len ? values[code] : ValueT {};
2428
}
2529
}
2630

@@ -39,6 +43,7 @@ template <typename IndexT>
3943
__device__ void dict_bool_kernel(const IndexT *const __restrict codes,
4044
uint64_t codes_len,
4145
const uint8_t *const __restrict values,
46+
uint64_t values_len,
4247
uint64_t values_bit_offset,
4348
uint8_t *const __restrict output) {
4449
const uint64_t output_len = (codes_len + 7) / 8;
@@ -53,7 +58,7 @@ __device__ void dict_bool_kernel(const IndexT *const __restrict codes,
5358
#pragma unroll
5459
for (uint32_t bit = 0; bit < 8; ++bit) {
5560
const uint64_t row = row_start + bit;
56-
if (row < codes_len) {
61+
if (row < codes_len && static_cast<uint64_t>(codes[row]) < values_len) {
5762
const uint64_t value_idx = values_bit_offset + static_cast<uint64_t>(codes[row]);
5863
const uint8_t value = (values[value_idx / 8] >> (value_idx % 8)) & 1;
5964
packed |= static_cast<uint8_t>(value << bit);
@@ -69,8 +74,9 @@ __device__ void dict_bool_kernel(const IndexT *const __restrict codes,
6974
const IndexType *const __restrict codes, \
7075
uint64_t codes_len, \
7176
const ValueType *const __restrict values, \
77+
uint64_t values_len, \
7278
ValueType *const __restrict output) { \
73-
dict_kernel<ValueType, IndexType>(codes, codes_len, values, output); \
79+
dict_kernel<ValueType, IndexType>(codes, codes_len, values, values_len, output); \
7480
}
7581

7682
// Generate dict kernel for all index types (unsigned integers) for a given value type
@@ -84,9 +90,10 @@ __device__ void dict_bool_kernel(const IndexT *const __restrict codes,
8490
extern "C" __global__ void dict_bool_##index_suffix(const IndexType *const __restrict codes, \
8591
uint64_t codes_len, \
8692
const uint8_t *const __restrict values, \
93+
uint64_t values_len, \
8794
uint64_t values_bit_offset, \
8895
uint8_t *const __restrict output) { \
89-
dict_bool_kernel<IndexType>(codes, codes_len, values, values_bit_offset, output); \
96+
dict_bool_kernel<IndexType>(codes, codes_len, values, values_len, values_bit_offset, output); \
9097
}
9198

9299
// Generate fixed-width kernels for all native ptypes and decimal values.

0 commit comments

Comments
 (0)