From 2d3640c8619a614eff14831ac84f0c6d07cc8d04 Mon Sep 17 00:00:00 2001 From: Jihoon Son Date: Mon, 24 Jun 2024 15:18:22 -0700 Subject: [PATCH 1/2] All kernels should have hidden visibility Signed-off-by: Jihoon Son --- src/main/cpp/faultinj/faultinj.cu | 4 ++-- src/main/cpp/src/bloom_filter.cu | 2 +- src/main/cpp/src/cast_string.cu | 5 ++--- src/main/cpp/src/cast_string_to_float.cu | 2 +- src/main/cpp/src/parse_uri.cu | 4 ++-- 5 files changed, 8 insertions(+), 9 deletions(-) diff --git a/src/main/cpp/faultinj/faultinj.cu b/src/main/cpp/faultinj/faultinj.cu index 13065a81ed..fcb4b3a12d 100644 --- a/src/main/cpp/faultinj/faultinj.cu +++ b/src/main/cpp/faultinj/faultinj.cu @@ -136,12 +136,12 @@ CUptiResult cuptiInitialize(void) return status; } -__global__ void faultInjectorKernelAssert(void) +__global__ static void faultInjectorKernelAssert(void) { assert(0 && "faultInjectorKernelAssert triggered"); } -__global__ void faultInjectorKernelTrap(void) { asm("trap;"); } +__global__ static void faultInjectorKernelTrap(void) { asm("trap;"); } boost::optional lookupConfig( boost::optional domainConfigs, diff --git a/src/main/cpp/src/bloom_filter.cu b/src/main/cpp/src/bloom_filter.cu index 5dfdd582ef..98ed58a62b 100644 --- a/src/main/cpp/src/bloom_filter.cu +++ b/src/main/cpp/src/bloom_filter.cu @@ -60,7 +60,7 @@ __device__ inline std::pair gpu_get_hash_ma } template -__global__ void gpu_bloom_filter_put(cudf::bitmask_type* const bloom_filter, +CUDF_KERNEL void gpu_bloom_filter_put(cudf::bitmask_type* const bloom_filter, cudf::size_type bloom_filter_bits, cudf::column_device_view input, cudf::size_type num_hashes) diff --git a/src/main/cpp/src/cast_string.cu b/src/main/cpp/src/cast_string.cu index bfbbc3777d..230431aecc 100644 --- a/src/main/cpp/src/cast_string.cu +++ b/src/main/cpp/src/cast_string.cu @@ -156,7 +156,7 @@ process_value(bool first_value, T current_val, T const new_digit, bool adding) * @param ansi_mode true if ansi mode is required, which is more strict and throws */ template -void __global__ string_to_integer_kernel(T* out, +void CUDF_KERNEL string_to_integer_kernel(T* out, bitmask_type* validity, const char* const chars, size_type const* offsets, @@ -386,10 +386,9 @@ __device__ thrust::optional> validate_and_exponent * @param scale scale of desired decimals * @param precision precision of desired decimals * @param ansi_mode true if ansi mode is required, which is more strict and throws - * @return __global__ */ template -__global__ void string_to_decimal_kernel(T* out, +CUDF_KERNEL void string_to_decimal_kernel(T* out, bitmask_type* validity, const char* const chars, size_type const* offsets, diff --git a/src/main/cpp/src/cast_string_to_float.cu b/src/main/cpp/src/cast_string_to_float.cu index cd7de88110..5b04bfe10a 100644 --- a/src/main/cpp/src/cast_string_to_float.cu +++ b/src/main/cpp/src/cast_string_to_float.cu @@ -618,7 +618,7 @@ class string_to_float { }; template -__global__ void string_to_float_kernel(T* out, +CUDF_KERNEL void string_to_float_kernel(T* out, bitmask_type* validity, int32_t* ansi_except, size_type* valid_count, diff --git a/src/main/cpp/src/parse_uri.cu b/src/main/cpp/src/parse_uri.cu index 0e57366358..d7ab08b1bc 100644 --- a/src/main/cpp/src/parse_uri.cu +++ b/src/main/cpp/src/parse_uri.cu @@ -770,7 +770,7 @@ uri_parts __device__ validate_uri(const char* str, * @param out_offsets Offsets to the start of the chunks * @param out_validity Bitmask of validity data, updated in function */ -__global__ void parse_uri_char_counter(column_device_view const in_strings, +CUDF_KERNEL void parse_uri_char_counter(column_device_view const in_strings, URI_chunks chunk, char const* const base_ptr, size_type* const out_lengths, @@ -850,7 +850,7 @@ __global__ void parse_uri_char_counter(column_device_view const in_strings, * @param offsets Offset value of each string associated with `out_chars` * @param out_chars Character buffer for the output string column */ -__global__ void parse_uri(column_device_view const in_strings, +CUDF_KERNEL void parse_uri(column_device_view const in_strings, char const* const base_ptr, size_type const* const src_offsets, size_type const* const offsets, From 4b9493067618ffdd64166d7824ec188e14c14f15 Mon Sep 17 00:00:00 2001 From: Jihoon Son Date: Tue, 25 Jun 2024 10:06:18 -0700 Subject: [PATCH 2/2] style fix Signed-off-by: Jihoon Son --- src/main/cpp/src/bloom_filter.cu | 6 ++--- src/main/cpp/src/cast_string.cu | 30 ++++++++++++------------ src/main/cpp/src/cast_string_to_float.cu | 14 +++++------ src/main/cpp/src/parse_uri.cu | 20 ++++++++-------- 4 files changed, 35 insertions(+), 35 deletions(-) diff --git a/src/main/cpp/src/bloom_filter.cu b/src/main/cpp/src/bloom_filter.cu index 98ed58a62b..da4e3c5cb9 100644 --- a/src/main/cpp/src/bloom_filter.cu +++ b/src/main/cpp/src/bloom_filter.cu @@ -61,9 +61,9 @@ __device__ inline std::pair gpu_get_hash_ma template CUDF_KERNEL void gpu_bloom_filter_put(cudf::bitmask_type* const bloom_filter, - cudf::size_type bloom_filter_bits, - cudf::column_device_view input, - cudf::size_type num_hashes) + cudf::size_type bloom_filter_bits, + cudf::column_device_view input, + cudf::size_type num_hashes) { size_t const tid = threadIdx.x + blockIdx.x * blockDim.x; if (tid >= input.size()) { return; } diff --git a/src/main/cpp/src/cast_string.cu b/src/main/cpp/src/cast_string.cu index 230431aecc..4267daae37 100644 --- a/src/main/cpp/src/cast_string.cu +++ b/src/main/cpp/src/cast_string.cu @@ -157,13 +157,13 @@ process_value(bool first_value, T current_val, T const new_digit, bool adding) */ template void CUDF_KERNEL string_to_integer_kernel(T* out, - bitmask_type* validity, - const char* const chars, - size_type const* offsets, - bitmask_type const* incoming_null_mask, - size_type num_rows, - bool ansi_mode, - bool strip) + bitmask_type* validity, + const char* const chars, + size_type const* offsets, + bitmask_type const* incoming_null_mask, + size_type num_rows, + bool ansi_mode, + bool strip) { auto const group = cooperative_groups::this_thread_block(); auto const warp = cooperative_groups::tiled_partition(group); @@ -389,14 +389,14 @@ __device__ thrust::optional> validate_and_exponent */ template CUDF_KERNEL void string_to_decimal_kernel(T* out, - bitmask_type* validity, - const char* const chars, - size_type const* offsets, - bitmask_type const* incoming_null_mask, - size_type num_rows, - int32_t scale, - int32_t precision, - bool strip) + bitmask_type* validity, + const char* const chars, + size_type const* offsets, + bitmask_type const* incoming_null_mask, + size_type num_rows, + int32_t scale, + int32_t precision, + bool strip) { auto const group = cooperative_groups::this_thread_block(); auto const warp = cooperative_groups::tiled_partition(group); diff --git a/src/main/cpp/src/cast_string_to_float.cu b/src/main/cpp/src/cast_string_to_float.cu index 5b04bfe10a..c19a2a10fe 100644 --- a/src/main/cpp/src/cast_string_to_float.cu +++ b/src/main/cpp/src/cast_string_to_float.cu @@ -619,13 +619,13 @@ class string_to_float { template CUDF_KERNEL void string_to_float_kernel(T* out, - bitmask_type* validity, - int32_t* ansi_except, - size_type* valid_count, - const char* const chars, - size_type const* offsets, - bitmask_type const* incoming_null_mask, - size_type const num_rows) + bitmask_type* validity, + int32_t* ansi_except, + size_type* valid_count, + const char* const chars, + size_type const* offsets, + bitmask_type const* incoming_null_mask, + size_type const num_rows) { size_type const tid = threadIdx.x + (blockDim.x * blockIdx.x); size_type const row = tid / 32; diff --git a/src/main/cpp/src/parse_uri.cu b/src/main/cpp/src/parse_uri.cu index d7ab08b1bc..f0a78f4f52 100644 --- a/src/main/cpp/src/parse_uri.cu +++ b/src/main/cpp/src/parse_uri.cu @@ -771,12 +771,12 @@ uri_parts __device__ validate_uri(const char* str, * @param out_validity Bitmask of validity data, updated in function */ CUDF_KERNEL void parse_uri_char_counter(column_device_view const in_strings, - URI_chunks chunk, - char const* const base_ptr, - size_type* const out_lengths, - size_type* const out_offsets, - bitmask_type* out_validity, - thrust::optional query_match) + URI_chunks chunk, + char const* const base_ptr, + size_type* const out_lengths, + size_type* const out_offsets, + bitmask_type* out_validity, + thrust::optional query_match) { // thread per row auto const tid = cudf::detail::grid_1d::global_thread_id(); @@ -851,10 +851,10 @@ CUDF_KERNEL void parse_uri_char_counter(column_device_view const in_strings, * @param out_chars Character buffer for the output string column */ CUDF_KERNEL void parse_uri(column_device_view const in_strings, - char const* const base_ptr, - size_type const* const src_offsets, - size_type const* const offsets, - char* const out_chars) + char const* const base_ptr, + size_type const* const src_offsets, + size_type const* const offsets, + char* const out_chars) { auto const tid = cudf::detail::grid_1d::global_thread_id();