Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Fixing potential overflow in array read in string to float cast #1591

Merged
merged 1 commit into from
Nov 28, 2023
Merged

Fixing potential overflow in array read in string to float cast #1591

Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
8 changes: 6 additions & 2 deletions src/main/cpp/src/cast_string_to_float.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -194,6 +194,8 @@ class string_to_float {
compute_validity(_valid, _except);
}

static constexpr int max_safe_digits = 19;

private:
// shuffle down to remove whitespace
__device__ void remove_leading_whitespace()
Expand Down Expand Up @@ -319,7 +321,6 @@ class string_to_float {

// have we seen a valid digit yet?
bool seen_valid_digit = false;
constexpr int max_safe_digits = 19;
do {
int num_chars = _blen - _bpos;

Expand Down Expand Up @@ -603,7 +604,9 @@ __global__ void string_to_float_kernel(T* out,
size_type const row = tid / 32;
if (row >= num_rows) { return; }

__shared__ uint64_t ipow[19];
// one more than max safe digits to ensure that we can reference
// max_safe_digits into the array.
__shared__ uint64_t ipow[string_to_float<T, block_size>::max_safe_digits + 1];
if (threadIdx.x == 0) {
ipow[0] = 1;
ipow[1] = 10;
Expand All @@ -624,6 +627,7 @@ __global__ void string_to_float_kernel(T* out,
ipow[16] = 10000000000000000;
ipow[17] = 100000000000000000;
ipow[18] = 1000000000000000000;
ipow[19] = 10000000000000000000;
}
__syncthreads();

Expand Down