diff --git a/.pre-commit-config.yaml b/.pre-commit-config.yaml index 238e5b44030..7e44091774f 100644 --- a/.pre-commit-config.yaml +++ b/.pre-commit-config.yaml @@ -63,7 +63,7 @@ repos: # Explicitly specify the pyproject.toml at the repo root, not per-project. args: ["--config=pyproject.toml"] - repo: https://github.com/pre-commit/mirrors-clang-format - rev: v16.0.1 + rev: v16.0.6 hooks: - id: clang-format types_or: [c, c++, cuda] diff --git a/cpp/benchmarks/iterator/iterator.cu b/cpp/benchmarks/iterator/iterator.cu index 7acf24c30a5..dcd13cf62c4 100644 --- a/cpp/benchmarks/iterator/iterator.cu +++ b/cpp/benchmarks/iterator/iterator.cu @@ -145,7 +145,7 @@ void BM_iterator(benchmark::State& state) cuda_event_timer raii(state, true); // flush_l2_cache = true, stream = 0 if (cub_or_thrust) { if (raw_or_iterator) { - raw_stream_bench_cub(hasnull_F, dev_result); // driven by raw pointer + raw_stream_bench_cub(hasnull_F, dev_result); // driven by raw pointer } else { iterator_bench_cub(hasnull_F, dev_result); // driven by riterator without nulls } diff --git a/cpp/benchmarks/stream_compaction/apply_boolean_mask.cpp b/cpp/benchmarks/stream_compaction/apply_boolean_mask.cpp index a6feaf04842..f78aa9fa654 100644 --- a/cpp/benchmarks/stream_compaction/apply_boolean_mask.cpp +++ b/cpp/benchmarks/stream_compaction/apply_boolean_mask.cpp @@ -59,8 +59,8 @@ void calculate_bandwidth(benchmark::State& state, cudf::size_type num_columns) int64_t const column_bytes_in = column_bytes_out; // we only read unmasked inputs int64_t const bytes_read = - (column_bytes_in + validity_bytes_in) * num_columns + // reading columns - mask_size; // reading boolean mask + (column_bytes_in + validity_bytes_in) * num_columns + // reading columns + mask_size; // reading boolean mask int64_t const bytes_written = (column_bytes_out + validity_bytes_out) * num_columns; // writing columns diff --git a/cpp/benchmarks/string/char_types.cpp b/cpp/benchmarks/string/char_types.cpp index 8e9e595fcef..59e6245fd41 100644 --- a/cpp/benchmarks/string/char_types.cpp +++ b/cpp/benchmarks/string/char_types.cpp @@ -43,7 +43,7 @@ static void bench_char_types(nvbench::state& state) state.set_cuda_stream(nvbench::make_cuda_stream_view(cudf::get_default_stream().value())); // gather some throughput statistics as well auto chars_size = input.chars_size(); - state.add_global_memory_reads(chars_size); // all bytes are read; + state.add_global_memory_reads(chars_size); // all bytes are read; if (api_type == "all") { state.add_global_memory_writes(num_rows); // output is a bool8 per row } else { diff --git a/cpp/benchmarks/string/extract.cpp b/cpp/benchmarks/string/extract.cpp index 9e67c5a5b52..135dadabbe4 100644 --- a/cpp/benchmarks/string/extract.cpp +++ b/cpp/benchmarks/string/extract.cpp @@ -43,7 +43,7 @@ static void bench_extract(nvbench::state& state) std::uniform_int_distribution words_dist(0, 999); std::vector samples(100); // 100 unique rows of data to reuse std::generate(samples.begin(), samples.end(), [&]() { - std::string row; // build a row of random tokens + std::string row; // build a row of random tokens while (static_cast(row.size()) < row_width) { row += std::to_string(words_dist(generator)) + " "; } diff --git a/cpp/include/cudf/column/column_device_view.cuh b/cpp/include/cudf/column/column_device_view.cuh index 05ef21bd750..35851a99822 100644 --- a/cpp/include/cudf/column/column_device_view.cuh +++ b/cpp/include/cudf/column/column_device_view.cuh @@ -1393,7 +1393,7 @@ struct pair_accessor { */ template struct pair_rep_accessor { - column_device_view const col; ///< column view of column in device + column_device_view const col; ///< column view of column in device using rep_type = device_storage_type_t; ///< representation type diff --git a/cpp/include/cudf/detail/copy_if.cuh b/cpp/include/cudf/detail/copy_if.cuh index 1dd91dcd865..ebe7e052b6d 100644 --- a/cpp/include/cudf/detail/copy_if.cuh +++ b/cpp/include/cudf/detail/copy_if.cuh @@ -133,7 +133,7 @@ __launch_bounds__(block_size) __global__ if (has_validity) { temp_valids[threadIdx.x] = false; // init shared memory if (threadIdx.x < cudf::detail::warp_size) temp_valids[block_size + threadIdx.x] = false; - __syncthreads(); // wait for init + __syncthreads(); // wait for init } if (mask_true) { diff --git a/cpp/include/cudf/detail/indexalator.cuh b/cpp/include/cudf/detail/indexalator.cuh index 0ab9da0dbd0..4731c4919e3 100644 --- a/cpp/include/cudf/detail/indexalator.cuh +++ b/cpp/include/cudf/detail/indexalator.cuh @@ -248,7 +248,7 @@ struct input_indexalator : base_indexalator { friend struct indexalator_factory; friend struct base_indexalator; // for CRTP - using reference = size_type const; // this keeps STL and thrust happy + using reference = size_type const; // this keeps STL and thrust happy input_indexalator() = default; input_indexalator(input_indexalator const&) = default; @@ -332,7 +332,7 @@ struct output_indexalator : base_indexalator { friend struct indexalator_factory; friend struct base_indexalator; // for CRTP - using reference = output_indexalator const&; // required for output iterators + using reference = output_indexalator const&; // required for output iterators output_indexalator() = default; output_indexalator(output_indexalator const&) = default; diff --git a/cpp/include/cudf/detail/join.hpp b/cpp/include/cudf/detail/join.hpp index 6fcf10aef57..b69632c83ca 100644 --- a/cpp/include/cudf/detail/join.hpp +++ b/cpp/include/cudf/detail/join.hpp @@ -78,8 +78,8 @@ struct hash_join { cudf::null_equality const _nulls_equal; ///< whether to consider nulls as equal cudf::table_view _build; ///< input table to build the hash map std::shared_ptr - _preprocessed_build; ///< input table preprocssed for row operators - map_type _hash_table; ///< hash table built on `_build` + _preprocessed_build; ///< input table preprocssed for row operators + map_type _hash_table; ///< hash table built on `_build` public: /** diff --git a/cpp/include/cudf/fixed_point/fixed_point.hpp b/cpp/include/cudf/fixed_point/fixed_point.hpp index 7c59c2f9194..13d8716c1df 100644 --- a/cpp/include/cudf/fixed_point/fixed_point.hpp +++ b/cpp/include/cudf/fixed_point/fixed_point.hpp @@ -829,5 +829,5 @@ using decimal32 = fixed_point; ///< 32-bit decima using decimal64 = fixed_point; ///< 64-bit decimal fixed point using decimal128 = fixed_point<__int128_t, Radix::BASE_10>; ///< 128-bit decimal fixed point -/** @} */ // end of group +/** @} */ // end of group } // namespace numeric diff --git a/cpp/include/cudf/groupby.hpp b/cpp/include/cudf/groupby.hpp index 6e575685daa..1c31e8777a8 100644 --- a/cpp/include/cudf/groupby.hpp +++ b/cpp/include/cudf/groupby.hpp @@ -386,8 +386,8 @@ class groupby { ///< indicates null order ///< of each column std::unique_ptr - _helper; ///< Helper object - ///< used by sort based implementation + _helper; ///< Helper object + ///< used by sort based implementation /** * @brief Get the sort helper object diff --git a/cpp/include/cudf/io/csv.hpp b/cpp/include/cudf/io/csv.hpp index c84ca7e6c73..b49a13a8ea9 100644 --- a/cpp/include/cudf/io/csv.hpp +++ b/cpp/include/cudf/io/csv.hpp @@ -213,7 +213,7 @@ class csv_reader_options { auto const max_row_bytes = 16 * 1024; // 16KB auto const column_bytes = 64; - auto const base_padding = 1024; // 1KB + auto const base_padding = 1024; // 1KB if (num_columns == 0) { // Use flat size if the number of columns is not known diff --git a/cpp/include/cudf/io/json.hpp b/cpp/include/cudf/io/json.hpp index 15dc2a614ad..d408d249a7f 100644 --- a/cpp/include/cudf/io/json.hpp +++ b/cpp/include/cudf/io/json.hpp @@ -207,7 +207,7 @@ class json_reader_options { auto const max_row_bytes = 16 * 1024; // 16KB auto const column_bytes = 64; - auto const base_padding = 1024; // 1KB + auto const base_padding = 1024; // 1KB if (num_columns == 0) { // Use flat size if the number of columns is not known diff --git a/cpp/include/cudf/io/orc_metadata.hpp b/cpp/include/cudf/io/orc_metadata.hpp index 623ee2e49fc..82d59803c25 100644 --- a/cpp/include/cudf/io/orc_metadata.hpp +++ b/cpp/include/cudf/io/orc_metadata.hpp @@ -111,10 +111,10 @@ struct string_statistics : minmax_statistics, sum_statistics count; ///< Count of `false` and `true` values + std::vector count; ///< count of `true` values }; /** @@ -141,8 +141,10 @@ using binary_statistics = sum_statistics; * the UNIX epoch. The `minimum_utc` and `maximum_utc` are the same values adjusted to UTC. */ struct timestamp_statistics : minmax_statistics { - std::optional minimum_utc; ///< minimum in milliseconds - std::optional maximum_utc; ///< maximum in milliseconds + std::optional minimum_utc; ///< minimum in milliseconds + std::optional maximum_utc; ///< maximum in milliseconds + std::optional minimum_nanos; ///< nanoseconds part of the minimum + std::optional maximum_nanos; ///< nanoseconds part of the maximum }; namespace orc { diff --git a/cpp/include/cudf/strings/capitalize.hpp b/cpp/include/cudf/strings/capitalize.hpp index 6d01ab047ba..57375e9ac6a 100644 --- a/cpp/include/cudf/strings/capitalize.hpp +++ b/cpp/include/cudf/strings/capitalize.hpp @@ -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. @@ -50,16 +50,18 @@ namespace strings { * * Any null string entries return corresponding null output column entries. * - * @throw cudf::logic_error if `delimiter.is_valid()` is `false`. + * @throw cudf::logic_error if `delimiter.is_valid()` is `false`. * - * @param input String column. - * @param delimiters Characters for identifying words to capitalize. + * @param input String column + * @param delimiters Characters for identifying words to capitalize + * @param stream CUDA stream used for device memory operations and kernel launches * @param mr Device memory resource used to allocate the returned column's device memory - * @return Column of strings capitalized from the input column. + * @return Column of strings capitalized from the input column */ std::unique_ptr capitalize( strings_column_view const& input, - string_scalar const& delimiters = string_scalar(""), + string_scalar const& delimiters = string_scalar("", true, cudf::get_default_stream()), + rmm::cuda_stream_view stream = cudf::get_default_stream(), rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** @@ -83,14 +85,16 @@ std::unique_ptr capitalize( * * Any null string entries return corresponding null output column entries. * - * @param input String column. - * @param sequence_type The character type that is used when identifying words. + * @param input String column + * @param sequence_type The character type that is used when identifying words + * @param stream CUDA stream used for device memory operations and kernel launches * @param mr Device memory resource used to allocate the returned column's device memory - * @return Column of titled strings. + * @return Column of titled strings */ std::unique_ptr title( strings_column_view const& input, string_character_types sequence_type = string_character_types::ALPHA, + rmm::cuda_stream_view stream = cudf::get_default_stream(), rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** @@ -112,12 +116,14 @@ std::unique_ptr title( * * Any null string entries result in corresponding null output column entries. * - * @param input String column. + * @param input String column + * @param stream CUDA stream used for device memory operations and kernel launches * @param mr Device memory resource used to allocate the returned column's device memory - * @return Column of type BOOL8. + * @return Column of type BOOL8 */ std::unique_ptr is_title( strings_column_view const& input, + rmm::cuda_stream_view stream = cudf::get_default_stream(), rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** @} */ // end of doxygen group diff --git a/cpp/include/cudf/strings/case.hpp b/cpp/include/cudf/strings/case.hpp index 06ba4f8d882..94191686a92 100644 --- a/cpp/include/cudf/strings/case.hpp +++ b/cpp/include/cudf/strings/case.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2022, NVIDIA CORPORATION. + * Copyright (c) 2019-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. @@ -38,11 +38,13 @@ namespace strings { * Any null entries create null entries in the output column. * * @param strings Strings instance for this operation. + * @param stream CUDA stream used for device memory operations and kernel launches * @param mr Device memory resource used to allocate the returned column's device memory. * @return New column of strings with characters converted. */ std::unique_ptr to_lower( strings_column_view const& strings, + rmm::cuda_stream_view stream = cudf::get_default_stream(), rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** @@ -55,11 +57,13 @@ std::unique_ptr to_lower( * Any null entries create null entries in the output column. * * @param strings Strings instance for this operation. + * @param stream CUDA stream used for device memory operations and kernel launches * @param mr Device memory resource used to allocate the returned column's device memory. * @return New column of strings with characters converted. */ std::unique_ptr to_upper( strings_column_view const& strings, + rmm::cuda_stream_view stream = cudf::get_default_stream(), rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** @@ -73,11 +77,13 @@ std::unique_ptr to_upper( * Any null entries create null entries in the output column. * * @param strings Strings instance for this operation. + * @param stream CUDA stream used for device memory operations and kernel launches * @param mr Device memory resource used to allocate the returned column's device memory. * @return New column of strings with characters converted. */ std::unique_ptr swapcase( strings_column_view const& strings, + rmm::cuda_stream_view stream = cudf::get_default_stream(), rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** @} */ // end of doxygen group diff --git a/cpp/include/cudf/strings/detail/convert/fixed_point_to_string.cuh b/cpp/include/cudf/strings/detail/convert/fixed_point_to_string.cuh new file mode 100644 index 00000000000..0ee26ec9ee2 --- /dev/null +++ b/cpp/include/cudf/strings/detail/convert/fixed_point_to_string.cuh @@ -0,0 +1,80 @@ +/* + * Copyright (c) 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. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ +#pragma once + +#include + +namespace cudf::strings::detail { + +/** + * @brief Returns the number of digits in the given fixed point number. + * + * @param value The value of the fixed point number + * @param scale The scale of the fixed point number + * @return int32_t The number of digits required to represent the fixed point number + */ +__device__ inline int32_t fixed_point_string_size(__int128_t const& value, int32_t scale) +{ + if (scale >= 0) return count_digits(value) + scale; + + auto const abs_value = numeric::detail::abs(value); + auto const exp_ten = numeric::detail::exp10<__int128_t>(-scale); + auto const fraction = count_digits(abs_value % exp_ten); + auto const num_zeros = std::max(0, (-scale - fraction)); + return static_cast(value < 0) + // sign if negative + count_digits(abs_value / exp_ten) + // integer + 1 + // decimal point + num_zeros + // zeros padding + fraction; // size of fraction +} + +/** + * @brief Converts the given fixed point number to a string. + * + * Caller is responsible for ensuring that the output buffer is large enough. The required output + * buffer size can be obtained by calling `fixed_point_string_size`. + * + * @param value The value of the fixed point number + * @param scale The scale of the fixed point number + * @param out_ptr The pointer to the output string + */ +__device__ inline void fixed_point_to_string(__int128_t const& value, int32_t scale, char* out_ptr) +{ + if (scale >= 0) { + out_ptr += integer_to_string(value, out_ptr); + thrust::generate_n(thrust::seq, out_ptr, scale, []() { return '0'; }); // add zeros + return; + } + + // scale < 0 + // write format: [-]integer.fraction + // where integer = abs(value) / (10^abs(scale)) + // fraction = abs(value) % (10^abs(scale)) + if (value < 0) *out_ptr++ = '-'; // add sign + auto const abs_value = numeric::detail::abs(value); + auto const exp_ten = numeric::detail::exp10<__int128_t>(-scale); + auto const num_zeros = std::max(0, (-scale - count_digits(abs_value % exp_ten))); + + out_ptr += integer_to_string(abs_value / exp_ten, out_ptr); // add the integer part + *out_ptr++ = '.'; // add decimal point + + thrust::generate_n(thrust::seq, out_ptr, num_zeros, []() { return '0'; }); // add zeros + out_ptr += num_zeros; + + integer_to_string(abs_value % exp_ten, out_ptr); // add the fraction part +} + +} // namespace cudf::strings::detail diff --git a/cpp/include/cudf/strings/detail/utf8.hpp b/cpp/include/cudf/strings/detail/utf8.hpp index df8e2885782..e04572535de 100644 --- a/cpp/include/cudf/strings/detail/utf8.hpp +++ b/cpp/include/cudf/strings/detail/utf8.hpp @@ -155,18 +155,18 @@ constexpr inline size_type from_char_utf8(char_utf8 character, char* str) constexpr uint32_t utf8_to_codepoint(cudf::char_utf8 utf8_char) { uint32_t unchr = 0; - if (utf8_char < 0x0000'0080) // single-byte pass thru + if (utf8_char < 0x0000'0080) // single-byte pass thru unchr = utf8_char; - else if (utf8_char < 0x0000'E000) // two bytes + else if (utf8_char < 0x0000'E000) // two bytes { - unchr = (utf8_char & 0x1F00) >> 2; // shift and - unchr |= (utf8_char & 0x003F); // unmask - } else if (utf8_char < 0x00F0'0000) // three bytes + unchr = (utf8_char & 0x1F00) >> 2; // shift and + unchr |= (utf8_char & 0x003F); // unmask + } else if (utf8_char < 0x00F0'0000) // three bytes { - unchr = (utf8_char & 0x0F'0000) >> 4; // get upper 4 bits - unchr |= (utf8_char & 0x00'3F00) >> 2; // shift and - unchr |= (utf8_char & 0x00'003F); // unmask - } else if (utf8_char <= 0xF800'0000u) // four bytes + unchr = (utf8_char & 0x0F'0000) >> 4; // get upper 4 bits + unchr |= (utf8_char & 0x00'3F00) >> 2; // shift and + unchr |= (utf8_char & 0x00'003F); // unmask + } else if (utf8_char <= 0xF800'0000u) // four bytes { unchr = (utf8_char & 0x0300'0000) >> 6; // upper 3 bits unchr |= (utf8_char & 0x003F'0000) >> 4; // next 6 bits @@ -185,20 +185,20 @@ constexpr uint32_t utf8_to_codepoint(cudf::char_utf8 utf8_char) constexpr cudf::char_utf8 codepoint_to_utf8(uint32_t unchr) { cudf::char_utf8 utf8 = 0; - if (unchr < 0x0000'0080) // single byte utf8 + if (unchr < 0x0000'0080) // single byte utf8 utf8 = unchr; - else if (unchr < 0x0000'0800) // double byte utf8 + else if (unchr < 0x0000'0800) // double byte utf8 { - utf8 = (unchr << 2) & 0x1F00; // shift bits for - utf8 |= (unchr & 0x3F); // utf8 encoding + utf8 = (unchr << 2) & 0x1F00; // shift bits for + utf8 |= (unchr & 0x3F); // utf8 encoding utf8 |= 0x0000'C080; - } else if (unchr < 0x0001'0000) // triple byte utf8 + } else if (unchr < 0x0001'0000) // triple byte utf8 { - utf8 = (unchr << 4) & 0x0F'0000; // upper 4 bits - utf8 |= (unchr << 2) & 0x00'3F00; // next 6 bits - utf8 |= (unchr & 0x3F); // last 6 bits + utf8 = (unchr << 4) & 0x0F'0000; // upper 4 bits + utf8 |= (unchr << 2) & 0x00'3F00; // next 6 bits + utf8 |= (unchr & 0x3F); // last 6 bits utf8 |= 0x00E0'8080; - } else if (unchr < 0x0011'0000) // quadruple byte utf8 + } else if (unchr < 0x0011'0000) // quadruple byte utf8 { utf8 = (unchr << 6) & 0x0700'0000; // upper 3 bits utf8 |= (unchr << 4) & 0x003F'0000; // next 6 bits diff --git a/cpp/include/cudf/table/row_operators.cuh b/cpp/include/cudf/table/row_operators.cuh index 599a85c8a54..4806f96c934 100644 --- a/cpp/include/cudf/table/row_operators.cuh +++ b/cpp/include/cudf/table/row_operators.cuh @@ -105,9 +105,9 @@ inline __device__ auto null_compare(bool lhs_is_null, bool rhs_is_null, null_ord { if (lhs_is_null and rhs_is_null) { // null (dictionary_wrapper const& lhs, using dictionary32 = dictionary_wrapper; ///< 32-bit integer indexed dictionary wrapper -/** @} */ // end of group +/** @} */ // end of group } // namespace cudf diff --git a/cpp/include/cudf_test/base_fixture.hpp b/cpp/include/cudf_test/base_fixture.hpp index b622d7c6b78..06aabbe4e9c 100644 --- a/cpp/include/cudf_test/base_fixture.hpp +++ b/cpp/include/cudf_test/base_fixture.hpp @@ -331,9 +331,9 @@ inline auto parse_cudf_test_opts(int argc, char** argv) cxxopts::Options options(argv[0], " - cuDF tests command line options"); char const* env_rmm_mode = std::getenv("GTEST_CUDF_RMM_MODE"); // Overridden by CLI options char const* env_stream_mode = - std::getenv("GTEST_CUDF_STREAM_MODE"); // Overridden by CLI options + std::getenv("GTEST_CUDF_STREAM_MODE"); // Overridden by CLI options char const* env_stream_error_mode = - std::getenv("GTEST_CUDF_STREAM_ERROR_MODE"); // Overridden by CLI options + std::getenv("GTEST_CUDF_STREAM_ERROR_MODE"); // Overridden by CLI options auto default_rmm_mode = env_rmm_mode ? env_rmm_mode : "pool"; auto default_stream_mode = env_stream_mode ? env_stream_mode : "default"; auto default_stream_error_mode = env_stream_error_mode ? env_stream_error_mode : "error"; diff --git a/cpp/include/nvtext/subword_tokenize.hpp b/cpp/include/nvtext/subword_tokenize.hpp index ac75f5e9147..72a899d70b4 100644 --- a/cpp/include/nvtext/subword_tokenize.hpp +++ b/cpp/include/nvtext/subword_tokenize.hpp @@ -44,7 +44,7 @@ struct hashed_vocabulary { std::unique_ptr bin_offsets; ///< uint16 column, containing the start index of each ///< bin in the flattened hash table std::unique_ptr - cp_metadata; ///< uint32 column, The code point metadata table to use for normalization + cp_metadata; ///< uint32 column, The code point metadata table to use for normalization std::unique_ptr aux_cp_table; ///< uint64 column, The auxiliary code point table to use for normalization }; diff --git a/cpp/scripts/run-clang-tidy.py b/cpp/scripts/run-clang-tidy.py index a617a4c0df7..e5e57dbf562 100644 --- a/cpp/scripts/run-clang-tidy.py +++ b/cpp/scripts/run-clang-tidy.py @@ -22,7 +22,7 @@ import shutil -EXPECTED_VERSION = "16.0.1" +EXPECTED_VERSION = "16.0.6" VERSION_REGEX = re.compile(r" LLVM version ([0-9.]+)") GPU_ARCH_REGEX = re.compile(r"sm_(\d+)") SPACES = re.compile(r"\s+") diff --git a/cpp/src/copying/contiguous_split.cu b/cpp/src/copying/contiguous_split.cu index e1a55ec5419..5ea56a05dcb 100644 --- a/cpp/src/copying/contiguous_split.cu +++ b/cpp/src/copying/contiguous_split.cu @@ -114,8 +114,8 @@ struct dst_buf_info { int bit_shift; // # of bits to shift right by (for validity buffers) size_type valid_count; // validity count for this block of work - int src_buf_index; // source buffer index - int dst_buf_index; // destination buffer index + int src_buf_index; // source buffer index + int dst_buf_index; // destination buffer index }; /** @@ -1384,7 +1384,7 @@ struct chunk_iteration_state { std::size_t starting_batch; ///< Starting batch index for the current iteration std::vector const h_num_buffs_per_iteration; ///< The count of batches per iteration std::vector const - h_size_of_buffs_per_iteration; ///< The size in bytes per iteration + h_size_of_buffs_per_iteration; ///< The size in bytes per iteration }; std::unique_ptr chunk_iteration_state::create( @@ -1989,7 +1989,7 @@ struct contiguous_split_state { // This can be 1 if `contiguous_split` is just packing and not splitting std::size_t const num_partitions; ///< The number of partitions to produce - size_type const num_src_bufs; ///< Number of source buffers including children + size_type const num_src_bufs; ///< Number of source buffers including children std::size_t const num_bufs; ///< Number of source buffers including children * number of splits diff --git a/cpp/src/groupby/sort/functors.hpp b/cpp/src/groupby/sort/functors.hpp index c378ac99727..be36956b929 100644 --- a/cpp/src/groupby/sort/functors.hpp +++ b/cpp/src/groupby/sort/functors.hpp @@ -94,12 +94,12 @@ struct store_result_functor { }; protected: - sort::sort_groupby_helper& helper; ///< Sort helper - cudf::detail::result_cache& cache; ///< cache of results to store into - column_view const& values; ///< Column of values to group and aggregate + sort::sort_groupby_helper& helper; ///< Sort helper + cudf::detail::result_cache& cache; ///< cache of results to store into + column_view const& values; ///< Column of values to group and aggregate - rmm::cuda_stream_view stream; ///< CUDA stream on which to execute kernels - rmm::mr::device_memory_resource* mr; ///< Memory resource to allocate space for results + rmm::cuda_stream_view stream; ///< CUDA stream on which to execute kernels + rmm::mr::device_memory_resource* mr; ///< Memory resource to allocate space for results sorted keys_are_sorted; ///< Whether the keys are sorted std::unique_ptr sorted_values; ///< Memoised grouped and sorted values diff --git a/cpp/src/io/avro/avro_gpu.cu b/cpp/src/io/avro/avro_gpu.cu index 2c634d9b590..365f6d6875c 100644 --- a/cpp/src/io/avro/avro_gpu.cu +++ b/cpp/src/io/avro/avro_gpu.cu @@ -303,7 +303,7 @@ avro_decode_row(schemadesc_s const* schema, // If within an array, check if we reached the last item if (array_repeat_count != 0 && array_children <= 0 && cur < end) { if (!--array_repeat_count) { - i = array_start; // Restart at the array parent + i = array_start; // Restart at the array parent } else { i = array_start + 1; // Restart after the array parent array_children = schema[array_start].count; diff --git a/cpp/src/io/comp/cpu_unbz2.cpp b/cpp/src/io/comp/cpu_unbz2.cpp index 7159ff30d7c..a116335b254 100644 --- a/cpp/src/io/comp/cpu_unbz2.cpp +++ b/cpp/src/io/comp/cpu_unbz2.cpp @@ -216,7 +216,7 @@ int32_t bz2_decompress_block(unbz_state_s* s) s->currBlockNo++; - skipbits(s, 32); // block CRC + skipbits(s, 32); // block CRC if (getbits(s, 1)) return BZ_DATA_ERROR; // blockRandomized not supported (old bzip versions) diff --git a/cpp/src/io/comp/debrotli.cu b/cpp/src/io/comp/debrotli.cu index 542ca031b7c..8bafd054bdb 100644 --- a/cpp/src/io/comp/debrotli.cu +++ b/cpp/src/io/comp/debrotli.cu @@ -121,7 +121,7 @@ __inline__ __device__ int brotli_context(int p1, int p2, int lut) struct huff_scratch_s { uint16_t code_length_histo[16]; uint8_t code_length_code_lengths[brotli_code_length_codes]; - int8_t offset[6]; // offsets in sorted table for each length + int8_t offset[6]; // offsets in sorted table for each length uint16_t lenvlctab[32]; uint16_t sorted[brotli_code_length_codes]; // symbols sorted by code length int16_t next_symbol[32]; @@ -1298,7 +1298,7 @@ static __device__ void InverseMoveToFrontTransform(debrotli_state_s* s, uint8_t* // Reinitialize elements that could have been changed. uint32_t i = 1; uint32_t upper_bound = s->mtf_upper_bound; - uint32_t* mtf = &s->mtf[1]; // Make mtf[-1] addressable. + uint32_t* mtf = &s->mtf[1]; // Make mtf[-1] addressable. auto* mtf_u8 = reinterpret_cast(mtf); uint32_t pattern = 0x0302'0100; // Little-endian diff --git a/cpp/src/io/comp/gpuinflate.cu b/cpp/src/io/comp/gpuinflate.cu index 42c4fbe7bea..8993815e560 100644 --- a/cpp/src/io/comp/gpuinflate.cu +++ b/cpp/src/io/comp/gpuinflate.cu @@ -124,11 +124,11 @@ struct inflate_state_s { uint8_t* outbase; ///< start of output buffer uint8_t* outend; ///< end of output buffer // Input state - uint8_t const* cur; ///< input buffer - uint8_t const* end; ///< end of input buffer + uint8_t const* cur; ///< input buffer + uint8_t const* end; ///< end of input buffer - uint2 bitbuf; ///< bit buffer (64-bit) - uint32_t bitpos; ///< position in bit buffer + uint2 bitbuf; ///< bit buffer (64-bit) + uint32_t bitpos; ///< position in bit buffer int32_t err; ///< Error status int btype; ///< current block type @@ -295,7 +295,7 @@ __device__ int construct( return 0; // complete, but decode() will fail // check for an over-subscribed or incomplete set of lengths - left = 1; // one possible code of zero length + left = 1; // one possible code of zero length for (len = 1; len <= max_bits; len++) { left <<= 1; // one more bit, double codes left left -= counts[len]; // deduct count from possible codes @@ -349,8 +349,8 @@ __device__ int init_dynamic(inflate_state_s* s) index = 0; while (index < nlen + ndist) { int symbol = decode(s, s->lencnt, s->lensym); - if (symbol < 0) return symbol; // invalid symbol - if (symbol < 16) // length in 0..15 + if (symbol < 0) return symbol; // invalid symbol + if (symbol < 16) // length in 0..15 lengths[index++] = symbol; else { // repeat instruction int len = 0; // last length to repeat, assume repeating zeros @@ -358,9 +358,9 @@ __device__ int init_dynamic(inflate_state_s* s) if (index == 0) return -5; // no last length! len = lengths[index - 1]; // last length symbol = 3 + getbits(s, 2); - } else if (symbol == 17) // repeat zero 3..10 times + } else if (symbol == 17) // repeat zero 3..10 times symbol = 3 + getbits(s, 3); - else // == 18, repeat zero 11..138 times + else // == 18, repeat zero 11..138 times symbol = 11 + getbits(s, 7); if (index + symbol > nlen + ndist) return -6; // too many lengths! while (symbol--) // repeat last or zero symbol times diff --git a/cpp/src/io/comp/uncomp.cpp b/cpp/src/io/comp/uncomp.cpp index 017fd8abb47..0d2d21333bb 100644 --- a/cpp/src/io/comp/uncomp.cpp +++ b/cpp/src/io/comp/uncomp.cpp @@ -28,7 +28,7 @@ #include // memset -#include // uncompress +#include // uncompress using cudf::host_span; @@ -47,7 +47,7 @@ struct gz_file_header_s { uint8_t os; // OS id }; -struct zip_eocd_s // end of central directory +struct zip_eocd_s // end of central directory { uint32_t sig; // 0x0605'4b50 uint16_t disk_id; // number of this disk @@ -59,7 +59,7 @@ struct zip_eocd_s // end of central directory // number uint16_t comment_len; // comment length (excluded from struct) }; -struct zip64_eocdl // end of central dir locator +struct zip64_eocdl // end of central dir locator { uint32_t sig; // 0x0706'4b50 uint32_t disk_start; // number of the disk with the start of the zip64 end of central directory @@ -67,7 +67,7 @@ struct zip64_eocdl // end of central dir locator uint32_t num_disks; // total number of disks }; -struct zip_cdfh_s // central directory file header +struct zip_cdfh_s // central directory file header { uint32_t sig; // 0x0201'4b50 uint16_t ver; // version made by @@ -111,7 +111,7 @@ struct bz2_file_header_s { struct gz_archive_s { gz_file_header_s const* fhdr; - uint16_t hcrc16; // header crc16 if present + uint16_t hcrc16; // header crc16 if present uint16_t xlen; uint8_t const* fxtra; // xlen bytes (optional) uint8_t const* fname; // zero-terminated original filename if present diff --git a/cpp/src/io/comp/unsnap.cu b/cpp/src/io/comp/unsnap.cu index a7a1cfd3f9e..c699502317f 100644 --- a/cpp/src/io/comp/unsnap.cu +++ b/cpp/src/io/comp/unsnap.cu @@ -45,7 +45,7 @@ void __device__ busy_wait(size_t cycles) struct unsnap_batch_s { int32_t len; // 1..64 = Number of bytes uint32_t - offset; // copy distance if greater than zero or negative of literal offset in byte stream + offset; // copy distance if greater than zero or negative of literal offset in byte stream }; /** diff --git a/cpp/src/io/json/json_column.cu b/cpp/src/io/json/json_column.cu index bdad16bd9f1..cabf904f020 100644 --- a/cpp/src/io/json/json_column.cu +++ b/cpp/src/io/json/json_column.cu @@ -169,7 +169,7 @@ reduce_to_column_tree(tree_meta_t& tree, }); // 4. unique_copy parent_node_ids, ranges - rmm::device_uvector column_levels(0, stream); // not required + rmm::device_uvector column_levels(0, stream); // not required rmm::device_uvector parent_col_ids(num_columns, stream); rmm::device_uvector col_range_begin(num_columns, stream); // Field names rmm::device_uvector col_range_end(num_columns, stream); diff --git a/cpp/src/io/json/nested_json_gpu.cu b/cpp/src/io/json/nested_json_gpu.cu index b691eaa8caf..0b49f97597d 100644 --- a/cpp/src/io/json/nested_json_gpu.cu +++ b/cpp/src/io/json/nested_json_gpu.cu @@ -762,18 +762,18 @@ auto get_translation_table(bool include_line_delimiter) nl_tokens({}), // LINE_BREAK {ValueBegin}}}; // OTHER pda_tlt[static_cast(pda_state_t::PD_BOA)] = { - { /*ROOT*/ - {ErrorBegin}, // OPENING_BRACE - {ErrorBegin}, // OPENING_BRACKET - {ErrorBegin}, // CLOSING_BRACE - {ErrorBegin}, // CLOSING_BRACKET - {ErrorBegin}, // QUOTE - {ErrorBegin}, // ESCAPE - {ErrorBegin}, // COMMA - {ErrorBegin}, // COLON - {ErrorBegin}, // WHITE_SPACE - nl_tokens({ErrorBegin}), // LINE_BREAK - {ErrorBegin}, // OTHER + { /*ROOT*/ + {ErrorBegin}, // OPENING_BRACE + {ErrorBegin}, // OPENING_BRACKET + {ErrorBegin}, // CLOSING_BRACE + {ErrorBegin}, // CLOSING_BRACKET + {ErrorBegin}, // QUOTE + {ErrorBegin}, // ESCAPE + {ErrorBegin}, // COMMA + {ErrorBegin}, // COLON + {ErrorBegin}, // WHITE_SPACE + nl_tokens({ErrorBegin}), // LINE_BREAK + {ErrorBegin}, // OTHER /*LIST*/ {StructBegin}, // OPENING_BRACE {ListBegin}, // OPENING_BRACKET @@ -799,18 +799,18 @@ auto get_translation_table(bool include_line_delimiter) nl_tokens({}), // LINE_BREAK {ErrorBegin}}}; // OTHER pda_tlt[static_cast(pda_state_t::PD_LON)] = { - { /*ROOT*/ - {ErrorBegin}, // OPENING_BRACE - {ErrorBegin}, // OPENING_BRACKET - {ErrorBegin}, // CLOSING_BRACE - {ErrorBegin}, // CLOSING_BRACKET - {ErrorBegin}, // QUOTE - {ErrorBegin}, // ESCAPE - {ErrorBegin}, // COMMA - {ErrorBegin}, // COLON - {ValueEnd}, // WHITE_SPACE - nl_tokens({ValueEnd}), // LINE_BREAK - {}, // OTHER + { /*ROOT*/ + {ErrorBegin}, // OPENING_BRACE + {ErrorBegin}, // OPENING_BRACKET + {ErrorBegin}, // CLOSING_BRACE + {ErrorBegin}, // CLOSING_BRACKET + {ErrorBegin}, // QUOTE + {ErrorBegin}, // ESCAPE + {ErrorBegin}, // COMMA + {ErrorBegin}, // COLON + {ValueEnd}, // WHITE_SPACE + nl_tokens({ValueEnd}), // LINE_BREAK + {}, // OTHER /*LIST*/ {ErrorBegin}, // OPENING_BRACE {ErrorBegin}, // OPENING_BRACKET @@ -824,17 +824,17 @@ auto get_translation_table(bool include_line_delimiter) nl_tokens({ValueEnd}), // LINE_BREAK {}, // OTHER /*STRUCT*/ - {ErrorBegin}, // OPENING_BRACE - {ErrorBegin}, // OPENING_BRACKET - {ValueEnd, StructMemberEnd, StructEnd}, // CLOSING_BRACE - {ErrorBegin}, // CLOSING_BRACKET - {ErrorBegin}, // QUOTE - {ErrorBegin}, // ESCAPE - {ValueEnd, StructMemberEnd}, // COMMA - {ErrorBegin}, // COLON - {ValueEnd}, // WHITE_SPACE - nl_tokens({ValueEnd}), // LINE_BREAK - {}}}; // OTHER + {ErrorBegin}, // OPENING_BRACE + {ErrorBegin}, // OPENING_BRACKET + {ValueEnd, StructMemberEnd, StructEnd}, // CLOSING_BRACE + {ErrorBegin}, // CLOSING_BRACKET + {ErrorBegin}, // QUOTE + {ErrorBegin}, // ESCAPE + {ValueEnd, StructMemberEnd}, // COMMA + {ErrorBegin}, // COLON + {ValueEnd}, // WHITE_SPACE + nl_tokens({ValueEnd}), // LINE_BREAK + {}}}; // OTHER pda_tlt[static_cast(pda_state_t::PD_STR)] = {{ /*ROOT*/ {}, // OPENING_BRACE @@ -974,17 +974,17 @@ auto get_translation_table(bool include_line_delimiter) nl_tokens({ErrorBegin}), // LINE_BREAK {ErrorBegin}, // OTHER /*STRUCT*/ - {ErrorBegin}, // OPENING_BRACE - {ErrorBegin}, // OPENING_BRACKET - {StructEnd}, // CLOSING_BRACE - {ErrorBegin}, // CLOSING_BRACKET - {StructMemberBegin, FieldNameBegin}, // QUOTE - {ErrorBegin}, // ESCAPE - {ErrorBegin}, // COMMA - {ErrorBegin}, // COLON - {}, // WHITE_SPACE - nl_tokens({}), // LINE_BREAK - {ErrorBegin}}}; // OTHER + {ErrorBegin}, // OPENING_BRACE + {ErrorBegin}, // OPENING_BRACKET + {StructEnd}, // CLOSING_BRACE + {ErrorBegin}, // CLOSING_BRACKET + {StructMemberBegin, FieldNameBegin}, // QUOTE + {ErrorBegin}, // ESCAPE + {ErrorBegin}, // COMMA + {ErrorBegin}, // COLON + {}, // WHITE_SPACE + nl_tokens({}), // LINE_BREAK + {ErrorBegin}}}; // OTHER pda_tlt[static_cast(pda_state_t::PD_FLN)] = {{ /*ROOT*/ {ErrorBegin}, // OPENING_BRACE @@ -1011,17 +1011,17 @@ auto get_translation_table(bool include_line_delimiter) nl_tokens({ErrorBegin}), // LINE_BREAK {ErrorBegin}, // OTHER /*STRUCT*/ - {}, // OPENING_BRACE - {}, // OPENING_BRACKET - {}, // CLOSING_BRACE - {}, // CLOSING_BRACKET - {FieldNameEnd}, // QUOTE - {}, // ESCAPE - {}, // COMMA - {}, // COLON - {}, // WHITE_SPACE - nl_tokens({}), // LINE_BREAK - {}}}; // OTHER + {}, // OPENING_BRACE + {}, // OPENING_BRACKET + {}, // CLOSING_BRACE + {}, // CLOSING_BRACKET + {FieldNameEnd}, // QUOTE + {}, // ESCAPE + {}, // COMMA + {}, // COLON + {}, // WHITE_SPACE + nl_tokens({}), // LINE_BREAK + {}}}; // OTHER pda_tlt[static_cast(pda_state_t::PD_FNE)] = {{ /*ROOT*/ {ErrorBegin}, // OPENING_BRACE @@ -1048,17 +1048,17 @@ auto get_translation_table(bool include_line_delimiter) nl_tokens({ErrorBegin}), // LINE_BREAK {ErrorBegin}, // OTHER /*STRUCT*/ - {}, // OPENING_BRACE - {}, // OPENING_BRACKET - {}, // CLOSING_BRACE - {}, // CLOSING_BRACKET - {}, // QUOTE - {}, // ESCAPE - {}, // COMMA - {}, // COLON - {}, // WHITE_SPACE - nl_tokens({}), // LINE_BREAK - {}}}; // OTHER + {}, // OPENING_BRACE + {}, // OPENING_BRACKET + {}, // CLOSING_BRACE + {}, // CLOSING_BRACKET + {}, // QUOTE + {}, // ESCAPE + {}, // COMMA + {}, // COLON + {}, // WHITE_SPACE + nl_tokens({}), // LINE_BREAK + {}}}; // OTHER pda_tlt[static_cast(pda_state_t::PD_PFN)] = {{ /*ROOT*/ {ErrorBegin}, // OPENING_BRACE @@ -1097,18 +1097,18 @@ auto get_translation_table(bool include_line_delimiter) nl_tokens({}), // LINE_BREAK {ErrorBegin}}}; // OTHER - pda_tlt[static_cast(pda_state_t::PD_ERR)] = {{ /*ROOT*/ - {}, // OPENING_BRACE - {}, // OPENING_BRACKET - {}, // CLOSING_BRACE - {}, // CLOSING_BRACKET - {}, // QUOTE - {}, // ESCAPE - {}, // COMMA - {}, // COLON - {}, // WHITE_SPACE - nl_tokens({}), // LINE_BREAK - {}, // OTHER + pda_tlt[static_cast(pda_state_t::PD_ERR)] = {{ /*ROOT*/ + {}, // OPENING_BRACE + {}, // OPENING_BRACKET + {}, // CLOSING_BRACE + {}, // CLOSING_BRACKET + {}, // QUOTE + {}, // ESCAPE + {}, // COMMA + {}, // COLON + {}, // WHITE_SPACE + nl_tokens({}), // LINE_BREAK + {}, // OTHER /*LIST*/ {}, // OPENING_BRACE {}, // OPENING_BRACKET diff --git a/cpp/src/io/orc/orc.cpp b/cpp/src/io/orc/orc.cpp index fc50b7118be..bc399b75ef9 100644 --- a/cpp/src/io/orc/orc.cpp +++ b/cpp/src/io/orc/orc.cpp @@ -178,7 +178,9 @@ void ProtobufReader::read(timestamp_statistics& s, size_t maxlen) auto op = std::tuple(field_reader(1, s.minimum), field_reader(2, s.maximum), field_reader(3, s.minimum_utc), - field_reader(4, s.maximum_utc)); + field_reader(4, s.maximum_utc), + field_reader(5, s.minimum_nanos), + field_reader(6, s.maximum_nanos)); function_builder(s, maxlen, op); } diff --git a/cpp/src/io/orc/orc_gpu.hpp b/cpp/src/io/orc/orc_gpu.hpp index 681cc0fb9d2..9b8df50a22a 100644 --- a/cpp/src/io/orc/orc_gpu.hpp +++ b/cpp/src/io/orc/orc_gpu.hpp @@ -157,7 +157,7 @@ struct EncChunk { uint8_t dtype_len; // data type length int32_t scale; // scale for decimals or timestamps - uint32_t* dict_index; // dictionary index from row index + uint32_t* dict_index; // dictionary index from row index uint32_t* decimal_offsets; orc_column_device_view const* column; }; diff --git a/cpp/src/io/orc/stats_enc.cu b/cpp/src/io/orc/stats_enc.cu index 069841980c1..69d7ec95acd 100644 --- a/cpp/src/io/orc/stats_enc.cu +++ b/cpp/src/io/orc/stats_enc.cu @@ -16,15 +16,16 @@ #include "orc_gpu.hpp" -#include #include +#include +#include + #include -namespace cudf { -namespace io { -namespace orc { -namespace gpu { +namespace cudf::io::orc::gpu { + +using strings::detail::fixed_point_string_size; constexpr unsigned int init_threads_per_group = 32; constexpr unsigned int init_groups_per_block = 4; @@ -58,13 +59,14 @@ __global__ void __launch_bounds__(init_threads_per_block) constexpr unsigned int buffersize_reduction_dim = 32; constexpr unsigned int block_size = buffersize_reduction_dim * buffersize_reduction_dim; constexpr unsigned int pb_fld_hdrlen = 1; -constexpr unsigned int pb_fld_hdrlen16 = 2; // > 127-byte length -constexpr unsigned int pb_fld_hdrlen32 = 5; // > 16KB length +constexpr unsigned int pb_fld_hdrlen32 = 5; +constexpr unsigned int pb_fldlen_int32 = 5; constexpr unsigned int pb_fldlen_int64 = 10; constexpr unsigned int pb_fldlen_float64 = 8; -constexpr unsigned int pb_fldlen_decimal = 40; // Assume decimal2string fits in 40 characters constexpr unsigned int pb_fldlen_bucket1 = 1 + pb_fldlen_int64; -constexpr unsigned int pb_fldlen_common = 2 * pb_fld_hdrlen + pb_fldlen_int64; +// statistics field number + number of values + has null +constexpr unsigned int pb_fldlen_common = + pb_fld_hdrlen + (pb_fld_hdrlen + pb_fldlen_int64) + 2 * pb_fld_hdrlen; template __global__ void __launch_bounds__(block_size, 1) @@ -87,21 +89,32 @@ __global__ void __launch_bounds__(block_size, 1) case dtype_int8: case dtype_int16: case dtype_int32: - case dtype_date32: case dtype_int64: - case dtype_timestamp64: stats_len = pb_fldlen_common + pb_fld_hdrlen + 3 * (pb_fld_hdrlen + pb_fldlen_int64); break; + case dtype_date32: + stats_len = pb_fldlen_common + pb_fld_hdrlen + 2 * (pb_fld_hdrlen + pb_fldlen_int64); + break; + case dtype_timestamp64: + stats_len = pb_fldlen_common + pb_fld_hdrlen + 4 * (pb_fld_hdrlen + pb_fldlen_int64) + + 2 * (pb_fld_hdrlen + pb_fldlen_int32); + break; case dtype_float32: case dtype_float64: stats_len = pb_fldlen_common + pb_fld_hdrlen + 3 * (pb_fld_hdrlen + pb_fldlen_float64); break; case dtype_decimal64: - case dtype_decimal128: - stats_len = pb_fldlen_common + pb_fld_hdrlen16 + 3 * (pb_fld_hdrlen + pb_fldlen_decimal); - break; + case dtype_decimal128: { + auto const scale = groups[idx].col_dtype.scale(); + auto const min_size = fixed_point_string_size(chunks[idx].min_value.d128_val, scale); + auto const max_size = fixed_point_string_size(chunks[idx].max_value.d128_val, scale); + auto const sum_size = fixed_point_string_size(chunks[idx].sum.d128_val, scale); + // common + total field length + encoded string lengths + strings + stats_len = pb_fldlen_common + pb_fld_hdrlen32 + 3 * (pb_fld_hdrlen + pb_fld_hdrlen32) + + min_size + max_size + sum_size; + } break; case dtype_string: - stats_len = pb_fldlen_common + pb_fld_hdrlen32 + 3 * (pb_fld_hdrlen + pb_fldlen_int64) + + stats_len = pb_fldlen_common + pb_fld_hdrlen32 + 3 * (pb_fld_hdrlen + pb_fld_hdrlen32) + chunks[idx].min_value.str_val.length + chunks[idx].max_value.str_val.length; break; case dtype_none: stats_len = pb_fldlen_common; @@ -126,9 +139,6 @@ struct stats_state_s { statistics_chunk chunk; statistics_merge_group group; statistics_dtype stats_dtype; //!< Statistics data type for this column - // ORC stats - uint64_t numberOfValues; - uint8_t hasNull; }; /* @@ -178,6 +188,15 @@ __device__ inline uint8_t* pb_put_binary(uint8_t* p, uint32_t id, void const* by return p + len; } +__device__ inline uint8_t* pb_put_decimal( + uint8_t* p, uint32_t id, __int128_t value, int32_t scale, int32_t len) +{ + p[0] = id * 8 + ProtofType::FIXEDLEN; + p = pb_encode_uint(p + 1, len); + strings::detail::fixed_point_to_string(value, scale, reinterpret_cast(p)); + return p + len; +} + // Protobuf field encoding for 64-bit raw encoding (double) __device__ inline uint8_t* pb_put_fixed64(uint8_t* p, uint32_t id, void const* raw64) { @@ -186,6 +205,15 @@ __device__ inline uint8_t* pb_put_fixed64(uint8_t* p, uint32_t id, void const* r return p + 9; } +// Splits a nanosecond timestamp into milliseconds and nanoseconds +__device__ std::pair split_nanosecond_timestamp(int64_t nano_count) +{ + auto const ns = cuda::std::chrono::nanoseconds(nano_count); + auto const ms_floor = cuda::std::chrono::floor(ns); + auto const ns_remainder = ns - ms_floor; + return {ms_floor.count(), ns_remainder.count()}; +} + /** * @brief Encode statistics in ORC protobuf format * @@ -228,12 +256,14 @@ __global__ void __launch_bounds__(encode_threads_per_block) // Encode and update actual bfr size if (idx < statistics_count && t == 0) { - s->chunk = chunks[idx]; - s->group = groups[idx]; - s->stats_dtype = s->group.stats_dtype; - s->base = blob_bfr + s->group.start_chunk; - s->end = blob_bfr + s->group.start_chunk + s->group.num_chunks; - uint8_t* cur = pb_put_uint(s->base, 1, s->chunk.non_nulls); + s->chunk = chunks[idx]; + s->group = groups[idx]; + s->stats_dtype = s->group.stats_dtype; + s->base = blob_bfr + s->group.start_chunk; + s->end = blob_bfr + s->group.start_chunk + s->group.num_chunks; + uint8_t* cur = pb_put_uint(s->base, 1, s->chunk.non_nulls); + cur = pb_put_uint(cur, 10, s->chunk.null_count != 0); // hasNull (bool) + uint8_t* fld_start = cur; switch (s->stats_dtype) { case dtype_int8: @@ -265,11 +295,14 @@ __global__ void __launch_bounds__(encode_threads_per_block) // optional double maximum = 2; // optional double sum = 3; // } - if (s->chunk.has_minmax) { + if (s->chunk.has_minmax || s->chunk.has_sum) { *cur = 3 * 8 + ProtofType::FIXEDLEN; cur += 2; - cur = pb_put_fixed64(cur, 1, &s->chunk.min_value.fp_val); - cur = pb_put_fixed64(cur, 2, &s->chunk.max_value.fp_val); + if (s->chunk.has_minmax) { + cur = pb_put_fixed64(cur, 1, &s->chunk.min_value.fp_val); + cur = pb_put_fixed64(cur, 2, &s->chunk.max_value.fp_val); + } + if (s->chunk.has_sum) { cur = pb_put_fixed64(cur, 3, &s->chunk.sum.fp_val); } fld_start[1] = cur - (fld_start + 2); } break; @@ -280,18 +313,25 @@ __global__ void __launch_bounds__(encode_threads_per_block) // optional string maximum = 2; // optional sint64 sum = 3; // sum will store the total length of all strings // } - if (s->chunk.has_minmax && s->chunk.has_sum) { - uint32_t sz = (pb_put_int(cur, 3, s->chunk.sum.i_val) - cur) + - (pb_put_uint(cur, 1, s->chunk.min_value.str_val.length) - cur) + - (pb_put_uint(cur, 2, s->chunk.max_value.str_val.length) - cur) + - s->chunk.min_value.str_val.length + s->chunk.max_value.str_val.length; + if (s->chunk.has_minmax || s->chunk.has_sum) { + uint32_t sz = 0; + if (s->chunk.has_minmax) { + sz += (pb_put_uint(cur, 1, s->chunk.min_value.str_val.length) - cur) + + (pb_put_uint(cur, 2, s->chunk.max_value.str_val.length) - cur) + + s->chunk.min_value.str_val.length + s->chunk.max_value.str_val.length; + } + if (s->chunk.has_sum) { sz += pb_put_int(cur, 3, s->chunk.sum.i_val) - cur; } + cur[0] = 4 * 8 + ProtofType::FIXEDLEN; cur = pb_encode_uint(cur + 1, sz); - cur = pb_put_binary( - cur, 1, s->chunk.min_value.str_val.ptr, s->chunk.min_value.str_val.length); - cur = pb_put_binary( - cur, 2, s->chunk.max_value.str_val.ptr, s->chunk.max_value.str_val.length); - cur = pb_put_int(cur, 3, s->chunk.sum.i_val); + + if (s->chunk.has_minmax) { + cur = pb_put_binary( + cur, 1, s->chunk.min_value.str_val.ptr, s->chunk.min_value.str_val.length); + cur = pb_put_binary( + cur, 2, s->chunk.max_value.str_val.ptr, s->chunk.max_value.str_val.length); + } + if (s->chunk.has_sum) { cur = pb_put_int(cur, 3, s->chunk.sum.i_val); } } break; case dtype_bool: @@ -299,8 +339,9 @@ __global__ void __launch_bounds__(encode_threads_per_block) // message BucketStatistics { // repeated uint64 count = 1 [packed=true]; // } - if (s->chunk.has_sum) { // Sum is equal to the number of 'true' values - cur[0] = 5 * 8 + ProtofType::FIXEDLEN; + if (s->chunk.has_sum) { + cur[0] = 5 * 8 + ProtofType::FIXEDLEN; + // count is equal to the number of 'true' values, despite what specs say cur = pb_put_packed_uint(cur + 2, 1, s->chunk.sum.u_val); fld_start[1] = cur - (fld_start + 2); } @@ -313,8 +354,33 @@ __global__ void __launch_bounds__(encode_threads_per_block) // optional string maximum = 2; // optional string sum = 3; // } - if (s->chunk.has_minmax) { - // TODO: Decimal support (decimal min/max stored as strings) + if (s->chunk.has_minmax or s->chunk.has_sum) { + auto const scale = s->group.col_dtype.scale(); + + uint32_t sz = 0; + auto const min_size = + s->chunk.has_minmax ? fixed_point_string_size(s->chunk.min_value.d128_val, scale) : 0; + auto const max_size = + s->chunk.has_minmax ? fixed_point_string_size(s->chunk.max_value.d128_val, scale) : 0; + if (s->chunk.has_minmax) { + // encoded string lengths, plus the strings + sz += (pb_put_uint(cur, 1, min_size) - cur) + min_size + + (pb_put_uint(cur, 1, max_size) - cur) + max_size; + } + auto const sum_size = + s->chunk.has_sum ? fixed_point_string_size(s->chunk.sum.d128_val, scale) : 0; + if (s->chunk.has_sum) { sz += (pb_put_uint(cur, 1, sum_size) - cur) + sum_size; } + + cur[0] = 6 * 8 + ProtofType::FIXEDLEN; + cur = pb_encode_uint(cur + 1, sz); + + if (s->chunk.has_minmax) { + cur = pb_put_decimal(cur, 1, s->chunk.min_value.d128_val, scale, min_size); // minimum + cur = pb_put_decimal(cur, 2, s->chunk.max_value.d128_val, scale, max_size); // maximum + } + if (s->chunk.has_sum) { + cur = pb_put_decimal(cur, 3, s->chunk.sum.d128_val, scale, sum_size); // sum + } } break; case dtype_date32: @@ -338,12 +404,24 @@ __global__ void __launch_bounds__(encode_threads_per_block) // optional sint64 maximum = 2; // optional sint64 minimumUtc = 3; // min,max values saved as milliseconds since UNIX epoch // optional sint64 maximumUtc = 4; + // optional int32 minimumNanos = 5; // lower 6 TS digits for min/max to achieve nanosecond + // precision optional int32 maximumNanos = 6; // } if (s->chunk.has_minmax) { cur[0] = 9 * 8 + ProtofType::FIXEDLEN; cur += 2; - cur = pb_put_int(cur, 3, s->chunk.min_value.i_val); // minimumUtc - cur = pb_put_int(cur, 4, s->chunk.max_value.i_val); // maximumUtc + auto const [min_ms, min_ns_remainder] = + split_nanosecond_timestamp(s->chunk.min_value.i_val); + auto const [max_ms, max_ns_remainder] = + split_nanosecond_timestamp(s->chunk.max_value.i_val); + + // minimum/maximum are the same as minimumUtc/maximumUtc as we always write files in UTC + cur = pb_put_int(cur, 1, min_ms); // minimum + cur = pb_put_int(cur, 2, max_ms); // maximum + cur = pb_put_int(cur, 3, min_ms); // minimumUtc + cur = pb_put_int(cur, 4, max_ms); // maximumUtc + cur = pb_put_int(cur, 5, min_ns_remainder); // minimumNanos + cur = pb_put_int(cur, 6, max_ns_remainder); // maximumNanos fld_start[1] = cur - (fld_start + 2); } break; @@ -403,7 +481,4 @@ void orc_encode_statistics(uint8_t* blob_bfr, blob_bfr, groups, chunks, statistics_count); } -} // namespace gpu -} // namespace orc -} // namespace io -} // namespace cudf +} // namespace cudf::io::orc::gpu diff --git a/cpp/src/io/orc/stripe_data.cu b/cpp/src/io/orc/stripe_data.cu index b66ca827119..3edcd3d83b2 100644 --- a/cpp/src/io/orc/stripe_data.cu +++ b/cpp/src/io/orc/stripe_data.cu @@ -367,14 +367,14 @@ inline __device__ uint32_t varint_length(volatile orc_bytestream_s* bs, int pos) if (zbit) { return 5 + (zbit >> 3); // up to 9x7 bits } else if ((sizeof(T) <= 8) || (bytestream_readbyte(bs, pos + 9) <= 0x7f)) { - return 10; // up to 70 bits + return 10; // up to 70 bits } else { uint64_t next64 = bytestream_readu64(bs, pos + 10); zbit = __ffsll((~next64) & 0x8080'8080'8080'8080ull); if (zbit) { return 10 + (zbit >> 3); // Up to 18x7 bits (126) } else { - return 19; // Up to 19x7 bits (133) + return 19; // Up to 19x7 bits (133) } } } diff --git a/cpp/src/io/parquet/compact_protocol_reader.cpp b/cpp/src/io/parquet/compact_protocol_reader.cpp index 92fcd151925..ae11af92f78 100644 --- a/cpp/src/io/parquet/compact_protocol_reader.cpp +++ b/cpp/src/io/parquet/compact_protocol_reader.cpp @@ -168,7 +168,7 @@ bool CompactProtocolReader::read(LogicalType* l) ParquetFieldUnion(2, l->isset.MAP, l->MAP), ParquetFieldUnion(3, l->isset.LIST, l->LIST), ParquetFieldUnion(4, l->isset.ENUM, l->ENUM), - ParquetFieldUnion(5, l->isset.DECIMAL, l->DECIMAL), // read the struct + ParquetFieldUnion(5, l->isset.DECIMAL, l->DECIMAL), // read the struct ParquetFieldUnion(6, l->isset.DATE, l->DATE), ParquetFieldUnion(7, l->isset.TIME, l->TIME), // read the struct ParquetFieldUnion(8, l->isset.TIMESTAMP, l->TIMESTAMP), // read the struct diff --git a/cpp/src/io/parquet/compact_protocol_writer.cpp b/cpp/src/io/parquet/compact_protocol_writer.cpp index b2a89129645..b2c0c97c52d 100644 --- a/cpp/src/io/parquet/compact_protocol_writer.cpp +++ b/cpp/src/io/parquet/compact_protocol_writer.cpp @@ -315,7 +315,7 @@ inline void CompactProtocolFieldWriter::field_struct(int field, T const& val) if constexpr (not std::is_empty_v) { writer.write(val); // write the struct if it's not empty } else { - put_byte(0); // otherwise, add a stop field + put_byte(0); // otherwise, add a stop field } current_field_value = field; } diff --git a/cpp/src/io/parquet/delta_binary.cuh b/cpp/src/io/parquet/delta_binary.cuh index 4fc8b9cfb8e..2382e4aafdf 100644 --- a/cpp/src/io/parquet/delta_binary.cuh +++ b/cpp/src/io/parquet/delta_binary.cuh @@ -90,16 +90,16 @@ inline __device__ zigzag128_t get_zz128(uint8_t const*& cur, uint8_t const* end) } struct delta_binary_decoder { - uint8_t const* block_start; // start of data, but updated as data is read - uint8_t const* block_end; // end of data - uleb128_t block_size; // usually 128, must be multiple of 128 - uleb128_t mini_block_count; // usually 4, chosen such that block_size/mini_block_count is a - // multiple of 32 - uleb128_t value_count; // total values encoded in the block - zigzag128_t last_value; // last value decoded, initialized to first_value from header - - uint32_t values_per_mb; // block_size / mini_block_count, must be multiple of 32 - uint32_t current_value_idx; // current value index, initialized to 0 at start of block + uint8_t const* block_start; // start of data, but updated as data is read + uint8_t const* block_end; // end of data + uleb128_t block_size; // usually 128, must be multiple of 128 + uleb128_t mini_block_count; // usually 4, chosen such that block_size/mini_block_count is a + // multiple of 32 + uleb128_t value_count; // total values encoded in the block + zigzag128_t last_value; // last value decoded, initialized to first_value from header + + uint32_t values_per_mb; // block_size / mini_block_count, must be multiple of 32 + uint32_t current_value_idx; // current value index, initialized to 0 at start of block zigzag128_t cur_min_delta; // min delta for the block uint32_t cur_mb; // index of the current mini-block within the block diff --git a/cpp/src/io/parquet/page_delta_decode.cu b/cpp/src/io/parquet/page_delta_decode.cu index e79a479388f..35f33a761be 100644 --- a/cpp/src/io/parquet/page_delta_decode.cu +++ b/cpp/src/io/parquet/page_delta_decode.cu @@ -85,7 +85,7 @@ __global__ void __launch_bounds__(96) gpuDecodeDeltaBinary( if (t < 2 * warp_size) { // warp0..1 target_pos = min(src_pos + 2 * batch_size, s->nz_count + batch_size); - } else { // warp2 + } else { // warp2 target_pos = min(s->nz_count, src_pos + batch_size); } __syncthreads(); diff --git a/cpp/src/io/parquet/page_enc.cu b/cpp/src/io/parquet/page_enc.cu index 0af561be8da..fe0dbb85124 100644 --- a/cpp/src/io/parquet/page_enc.cu +++ b/cpp/src/io/parquet/page_enc.cu @@ -1858,8 +1858,8 @@ __device__ std::pair get_extremum(statistics_val const* s } case dtype_int64: case dtype_timestamp64: - case dtype_float64: - case dtype_decimal64: return {stats_val, sizeof(int64_t)}; + case dtype_float64: return {stats_val, sizeof(int64_t)}; + case dtype_decimal64: case dtype_decimal128: byte_reverse128(stats_val->d128_val, scratch); return {scratch, sizeof(__int128_t)}; diff --git a/cpp/src/io/parquet/parquet.hpp b/cpp/src/io/parquet/parquet.hpp index a729f28d672..f7318bb9935 100644 --- a/cpp/src/io/parquet/parquet.hpp +++ b/cpp/src/io/parquet/parquet.hpp @@ -365,8 +365,8 @@ struct ColumnIndex { std::vector> min_values; // lower bound for values in each page std::vector> max_values; // upper bound for values in each page BoundaryOrder boundary_order = - BoundaryOrder::UNORDERED; // Indicates if min and max values are ordered - std::vector null_counts; // Optional count of null values per page + BoundaryOrder::UNORDERED; // Indicates if min and max values are ordered + std::vector null_counts; // Optional count of null values per page }; // bit space we are reserving in column_buffer::user_data diff --git a/cpp/src/io/parquet/parquet_gpu.hpp b/cpp/src/io/parquet/parquet_gpu.hpp index e82b6abc13d..a3cc37dee4f 100644 --- a/cpp/src/io/parquet/parquet_gpu.hpp +++ b/cpp/src/io/parquet/parquet_gpu.hpp @@ -299,7 +299,7 @@ struct ColumnChunkDesc { int8_t converted_type; // converted type enum LogicalType logical_type; // logical type int8_t decimal_precision; // Decimal precision - int32_t ts_clock_rate; // output timestamp clock frequency (0=default, 1000=ms, 1000000000=ns) + int32_t ts_clock_rate; // output timestamp clock frequency (0=default, 1000=ms, 1000000000=ns) int32_t src_col_index; // my input column index int32_t src_col_schema; // my schema index in the file @@ -396,16 +396,16 @@ constexpr uint32_t encoding_to_mask(Encoding encoding) struct EncColumnChunk { parquet_column_device_view const* col_desc; //!< Column description size_type col_desc_id; - PageFragment* fragments; //!< First fragment in chunk - uint8_t* uncompressed_bfr; //!< Uncompressed page data - uint8_t* compressed_bfr; //!< Compressed page data - statistics_chunk const* stats; //!< Fragment statistics - uint32_t bfr_size; //!< Uncompressed buffer size - uint32_t compressed_size; //!< Compressed buffer size - uint32_t max_page_data_size; //!< Max data size (excluding header) of any page in this chunk - uint32_t page_headers_size; //!< Sum of size of all page headers - size_type start_row; //!< First row of chunk - uint32_t num_rows; //!< Number of rows in chunk + PageFragment* fragments; //!< First fragment in chunk + uint8_t* uncompressed_bfr; //!< Uncompressed page data + uint8_t* compressed_bfr; //!< Compressed page data + statistics_chunk const* stats; //!< Fragment statistics + uint32_t bfr_size; //!< Uncompressed buffer size + uint32_t compressed_size; //!< Compressed buffer size + uint32_t max_page_data_size; //!< Max data size (excluding header) of any page in this chunk + uint32_t page_headers_size; //!< Sum of size of all page headers + size_type start_row; //!< First row of chunk + uint32_t num_rows; //!< Number of rows in chunk size_type num_values; //!< Number of values in chunk. Different from num_rows for nested types uint32_t first_fragment; //!< First fragment of chunk EncPage* pages; //!< Ptr to pages that belong to this chunk diff --git a/cpp/src/io/parquet/reader_impl_preprocess.cu b/cpp/src/io/parquet/reader_impl_preprocess.cu index bde73c3dd96..a2db0de26bb 100644 --- a/cpp/src/io/parquet/reader_impl_preprocess.cu +++ b/cpp/src/io/parquet/reader_impl_preprocess.cu @@ -1673,7 +1673,7 @@ void reader::impl::preprocess_pages(size_t skip_rows, // - we will be doing a chunked read gpu::ComputePageSizes(pages, chunks, - 0, // 0-max size_t. process all possible rows + 0, // 0-max size_t. process all possible rows std::numeric_limits::max(), true, // compute num_rows chunk_read_limit > 0, // compute string sizes diff --git a/cpp/src/io/statistics/statistics_type_identification.cuh b/cpp/src/io/statistics/statistics_type_identification.cuh index 32931d7d34d..ea8c71f0dcb 100644 --- a/cpp/src/io/statistics/statistics_type_identification.cuh +++ b/cpp/src/io/statistics/statistics_type_identification.cuh @@ -49,15 +49,15 @@ enum class is_int96_timestamp { YES, NO }; template struct conversion_map; -// Every timestamp or duration type is converted to milliseconds in ORC statistics +// Every timestamp or duration type is converted to nanoseconds in ORC statistics template struct conversion_map { - using types = std::tuple, - std::pair, - std::pair, - std::pair, - std::pair, - std::pair>; + using types = std::tuple, + std::pair, + std::pair, + std::pair, + std::pair, + std::pair>; }; // In Parquet timestamps and durations with second resolution are converted to @@ -125,7 +125,7 @@ class extrema_type { using non_arithmetic_extrema_type = typename std::conditional_t< cudf::is_fixed_point() or cudf::is_duration() or cudf::is_timestamp(), - typename std::conditional_t, __int128_t, int64_t>, + typename std::conditional_t(), __int128_t, int64_t>, typename std::conditional_t< std::is_same_v, string_view, @@ -134,8 +134,7 @@ class extrema_type { // unsigned int/bool -> uint64_t // signed int -> int64_t // float/double -> double - // decimal32/64 -> int64_t - // decimal128 -> __int128_t + // decimal32/64/128 -> __int128_t // duration_[T] -> int64_t // string_view -> string_view // byte_array_view -> byte_array_view diff --git a/cpp/src/io/statistics/typed_statistics_chunk.cuh b/cpp/src/io/statistics/typed_statistics_chunk.cuh index d007209a12a..e6ec1471cb7 100644 --- a/cpp/src/io/statistics/typed_statistics_chunk.cuh +++ b/cpp/src/io/statistics/typed_statistics_chunk.cuh @@ -244,9 +244,9 @@ get_untyped_chunk(typed_statistics_chunk const& chunk) stat.null_count = chunk.null_count; stat.has_minmax = chunk.has_minmax; stat.has_sum = [&]() { - if (!chunk.has_minmax) return false; // invalidate the sum if overflow or underflow is possible if constexpr (std::is_floating_point_v or std::is_integral_v) { + if (!chunk.has_minmax) { return true; } return std::numeric_limits::max() / chunk.non_nulls >= static_cast(chunk.maximum_value) and std::numeric_limits::lowest() / chunk.non_nulls <= diff --git a/cpp/src/join/join.cu b/cpp/src/join/join.cu index 8210f3114d6..ae025b1a213 100644 --- a/cpp/src/join/join.cu +++ b/cpp/src/join/join.cu @@ -73,7 +73,7 @@ left_join(table_view const& left_input, // Make sure any dictionary columns have matched key sets. // This will return any new dictionary columns created as well as updated table_views. auto matched = cudf::dictionary::detail::match_dictionaries( - {left_input, right_input}, // these should match + {left_input, right_input}, // these should match stream, rmm::mr::get_current_device_resource()); // temporary objects returned // now rebuild the table views with the updated ones @@ -98,7 +98,7 @@ full_join(table_view const& left_input, // Make sure any dictionary columns have matched key sets. // This will return any new dictionary columns created as well as updated table_views. auto matched = cudf::dictionary::detail::match_dictionaries( - {left_input, right_input}, // these should match + {left_input, right_input}, // these should match stream, rmm::mr::get_current_device_resource()); // temporary objects returned // now rebuild the table views with the updated ones diff --git a/cpp/src/quantiles/tdigest/tdigest_aggregation.cu b/cpp/src/quantiles/tdigest/tdigest_aggregation.cu index 2ce55e10fb1..9e8b75ae3b6 100644 --- a/cpp/src/quantiles/tdigest/tdigest_aggregation.cu +++ b/cpp/src/quantiles/tdigest/tdigest_aggregation.cu @@ -459,7 +459,7 @@ __global__ void generate_cluster_limits_kernel(int delta, int adjusted_w_index = nearest_w_index; if ((last_inserted_index < 0) || // if we haven't inserted anything yet (nearest_w_index == - last_inserted_index)) { // if we land in the same bucket as the previous cap + last_inserted_index)) { // if we land in the same bucket as the previous cap // force the value into this bucket adjusted_w_index = (last_inserted_index == group_size - 1) diff --git a/cpp/src/rolling/detail/rolling_collect_list.cuh b/cpp/src/rolling/detail/rolling_collect_list.cuh index 9f74a961e12..39d15ed716f 100644 --- a/cpp/src/rolling/detail/rolling_collect_list.cuh +++ b/cpp/src/rolling/detail/rolling_collect_list.cuh @@ -116,7 +116,7 @@ std::unique_ptr create_collect_gather_map(column_view const& child_offse thrust::make_counting_iterator(per_row_mapping.size()), gather_map->mutable_view().template begin(), [d_offsets = - child_offsets.template begin(), // E.g. [0, 2, 5, 8, 11, 13] + child_offsets.template begin(), // E.g. [0, 2, 5, 8, 11, 13] d_groups = per_row_mapping.template begin(), // E.g. [0,0, 1,1,1, 2,2,2, 3,3,3, 4,4] d_prev = preceding_iter] __device__(auto i) { diff --git a/cpp/src/strings/capitalize.cu b/cpp/src/strings/capitalize.cu index 4e248922702..c555031b588 100644 --- a/cpp/src/strings/capitalize.cu +++ b/cpp/src/strings/capitalize.cu @@ -287,25 +287,28 @@ std::unique_ptr is_title(strings_column_view const& input, std::unique_ptr capitalize(strings_column_view const& input, string_scalar const& delimiter, + rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); - return detail::capitalize(input, delimiter, cudf::get_default_stream(), mr); + return detail::capitalize(input, delimiter, stream, mr); } std::unique_ptr title(strings_column_view const& input, string_character_types sequence_type, + rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); - return detail::title(input, sequence_type, cudf::get_default_stream(), mr); + return detail::title(input, sequence_type, stream, mr); } std::unique_ptr is_title(strings_column_view const& input, + rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); - return detail::is_title(input, cudf::get_default_stream(), mr); + return detail::is_title(input, stream, mr); } } // namespace strings diff --git a/cpp/src/strings/case.cu b/cpp/src/strings/case.cu index c5fe7a19f53..8f4c2ee574a 100644 --- a/cpp/src/strings/case.cu +++ b/cpp/src/strings/case.cu @@ -310,24 +310,27 @@ std::unique_ptr swapcase(strings_column_view const& strings, // APIs std::unique_ptr to_lower(strings_column_view const& strings, + rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); - return detail::to_lower(strings, cudf::get_default_stream(), mr); + return detail::to_lower(strings, stream, mr); } std::unique_ptr to_upper(strings_column_view const& strings, + rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); - return detail::to_upper(strings, cudf::get_default_stream(), mr); + return detail::to_upper(strings, stream, mr); } std::unique_ptr swapcase(strings_column_view const& strings, + rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); - return detail::swapcase(strings, cudf::get_default_stream(), mr); + return detail::swapcase(strings, stream, mr); } } // namespace strings diff --git a/cpp/src/strings/char_types/char_types.cu b/cpp/src/strings/char_types/char_types.cu index b87fb80fcc2..0c0ad0ad29e 100644 --- a/cpp/src/strings/char_types/char_types.cu +++ b/cpp/src/strings/char_types/char_types.cu @@ -139,9 +139,9 @@ struct filter_chars_fn { { auto const code_point = detail::utf8_to_codepoint(ch); auto const flag = code_point <= 0x00'FFFF ? d_flags[code_point] : 0; - if (flag == 0) // all types pass unless specifically identified + if (flag == 0) // all types pass unless specifically identified return (types_to_remove == ALL_TYPES); - if (types_to_keep == ALL_TYPES) // filter case + if (types_to_keep == ALL_TYPES) // filter case return (types_to_remove & flag) != 0; return (types_to_keep & flag) == 0; // keep case } diff --git a/cpp/src/strings/convert/convert_datetime.cu b/cpp/src/strings/convert/convert_datetime.cu index cca06ca0739..8a953d778ed 100644 --- a/cpp/src/strings/convert/convert_datetime.cu +++ b/cpp/src/strings/convert/convert_datetime.cu @@ -317,8 +317,8 @@ struct parse_datetime { bytes_read -= left; break; } - case 'u': [[fallthrough]]; // day of week: Mon(1)-Sat(6),Sun(7) - case 'w': { // day of week; Sun(0),Mon(1)-Sat(6) + case 'u': [[fallthrough]]; // day of week: Mon(1)-Sat(6),Sun(7) + case 'w': { // day of week; Sun(0),Mon(1)-Sat(6) auto const [weekday, left] = parse_int(ptr, item.length); timeparts.weekday = // 0 is mapped to 7 for chrono library static_cast((item.value == 'w' && weekday == 0) ? 7 : weekday); @@ -1000,7 +1000,7 @@ struct datetime_formatter_fn { case 'S': // second copy_value = timeparts.second; break; - case 'f': // sub-second + case 'f': // sub-second { char subsecond_digits[] = "000000000"; // 9 max digits int const digits = [] { diff --git a/cpp/src/strings/convert/convert_durations.cu b/cpp/src/strings/convert/convert_durations.cu index 863f76b9b98..6ab70825a6b 100644 --- a/cpp/src/strings/convert/convert_durations.cu +++ b/cpp/src/strings/convert/convert_durations.cu @@ -576,7 +576,7 @@ struct parse_duration { item_length++; // : timeparts->second = parse_second(ptr + item_length, item_length); break; - case 'r': // hh:MM:SS AM/PM + case 'r': // hh:MM:SS AM/PM timeparts->hour = parse_hour(ptr, item_length); item_length++; // : timeparts->minute = parse_minute(ptr + item_length, item_length); diff --git a/cpp/src/strings/convert/convert_fixed_point.cu b/cpp/src/strings/convert/convert_fixed_point.cu index a3336258d3e..51aab9faeba 100644 --- a/cpp/src/strings/convert/convert_fixed_point.cu +++ b/cpp/src/strings/convert/convert_fixed_point.cu @@ -21,7 +21,7 @@ #include #include #include -#include +#include #include #include #include @@ -200,62 +200,19 @@ struct from_fixed_point_fn { size_type* d_offsets{}; char* d_chars{}; - /** - * @brief Calculates the size of the string required to convert the element, in base-10 format. - * - * Output format is [-]integer.fraction - */ - __device__ int32_t compute_output_size(DecimalType value) - { - auto const scale = d_decimals.type().scale(); - - if (scale >= 0) return count_digits(value) + scale; - - auto const abs_value = numeric::detail::abs(value); - auto const exp_ten = numeric::detail::exp10(-scale); - auto const fraction = count_digits(abs_value % exp_ten); - auto const num_zeros = std::max(0, (-scale - fraction)); - return static_cast(value < 0) + // sign if negative - count_digits(abs_value / exp_ten) + // integer - 1 + // decimal point - num_zeros + // zeros padding - fraction; // size of fraction - } - /** * @brief Converts a decimal element into a string. * * The value is converted into base-10 digits [0-9] * plus the decimal point and a negative sign prefix. */ - __device__ void decimal_to_string(size_type idx) + __device__ void fixed_point_element_to_string(size_type idx) { auto const value = d_decimals.element(idx); auto const scale = d_decimals.type().scale(); char* d_buffer = d_chars + d_offsets[idx]; - if (scale >= 0) { - d_buffer += integer_to_string(value, d_buffer); - thrust::generate_n(thrust::seq, d_buffer, scale, []() { return '0'; }); // add zeros - return; - } - - // scale < 0 - // write format: [-]integer.fraction - // where integer = abs(value) / (10^abs(scale)) - // fraction = abs(value) % (10^abs(scale)) - if (value < 0) *d_buffer++ = '-'; // add sign - auto const abs_value = numeric::detail::abs(value); - auto const exp_ten = numeric::detail::exp10(-scale); - auto const num_zeros = std::max(0, (-scale - count_digits(abs_value % exp_ten))); - - d_buffer += integer_to_string(abs_value / exp_ten, d_buffer); // add the integer part - *d_buffer++ = '.'; // add decimal point - - thrust::generate_n(thrust::seq, d_buffer, num_zeros, []() { return '0'; }); // add zeros - d_buffer += num_zeros; - - integer_to_string(abs_value % exp_ten, d_buffer); // add the fraction part + fixed_point_to_string(value, scale, d_buffer); } __device__ void operator()(size_type idx) @@ -265,9 +222,10 @@ struct from_fixed_point_fn { return; } if (d_chars != nullptr) { - decimal_to_string(idx); + fixed_point_element_to_string(idx); } else { - d_offsets[idx] = compute_output_size(d_decimals.element(idx)); + d_offsets[idx] = + fixed_point_string_size(d_decimals.element(idx), d_decimals.type().scale()); } } }; diff --git a/cpp/src/strings/convert/convert_floats.cu b/cpp/src/strings/convert/convert_floats.cu index ab1e6870937..32167589ab4 100644 --- a/cpp/src/strings/convert/convert_floats.cu +++ b/cpp/src/strings/convert/convert_floats.cu @@ -284,7 +284,7 @@ struct ftos_converter { while (pb != buffer) // reverses the digits *ptr++ = *--pb; // e.g. 54321 -> 12345 } else - *ptr++ = '0'; // always include at least .0 + *ptr++ = '0'; // always include at least .0 // exponent if (exp10) { *ptr++ = 'e'; @@ -310,7 +310,7 @@ struct ftos_converter { { if (std::isnan(value)) return 3; // NaN bool bneg = false; - if (signbit(value)) { // handles -0.0 too + if (signbit(value)) { // handles -0.0 too value = -value; bneg = true; } @@ -337,7 +337,7 @@ struct ftos_converter { ++count; // always include .0 // exponent if (exp10) { - count += 2; // 'e±' + count += 2; // 'e±' if (exp10 < 0) exp10 = -exp10; count += (int)(exp10 < 10); // padding while (exp10 > 0) { diff --git a/cpp/src/strings/convert/convert_integers.cu b/cpp/src/strings/convert/convert_integers.cu index 260c3393f3c..5597d2831c0 100644 --- a/cpp/src/strings/convert/convert_integers.cu +++ b/cpp/src/strings/convert/convert_integers.cu @@ -76,7 +76,7 @@ struct string_to_integer_check_fn { auto const digit = static_cast(chr - '0'); auto const bound_check = (bound_val - sign * digit) / IntegerType{10} * sign; if (value > bound_check) return false; - value = value* IntegerType{10} + digit; + value = value * IntegerType{10} + digit; } return true; diff --git a/cpp/src/strings/convert/convert_ipv4.cu b/cpp/src/strings/convert/convert_ipv4.cu index 4606aba6d17..adb72cb0263 100644 --- a/cpp/src/strings/convert/convert_ipv4.cu +++ b/cpp/src/strings/convert/convert_ipv4.cu @@ -197,7 +197,7 @@ std::unique_ptr is_ipv4(strings_column_view const& strings, if (d_str.empty()) return false; constexpr int max_ip = 255; // values must be in [0,255] int ip_vals[4] = {-1, -1, -1, -1}; - int ipv_idx = 0; // index into ip_vals + int ipv_idx = 0; // index into ip_vals for (auto const ch : d_str) { if ((ch >= '0') && (ch <= '9')) { auto const ip_val = ip_vals[ipv_idx]; diff --git a/cpp/src/strings/convert/convert_urls.cu b/cpp/src/strings/convert/convert_urls.cu index 71b6c09310e..9efa148cfd2 100644 --- a/cpp/src/strings/convert/convert_urls.cu +++ b/cpp/src/strings/convert/convert_urls.cu @@ -107,9 +107,9 @@ struct url_encoder_fn { out_ptr = copy_and_increment(out_ptr, hex, 2); // add them to the output } } - } else // these are to be utf-8 url-encoded + } else // these are to be utf-8 url-encoded { - uint8_t char_bytes[4]; // holds utf-8 bytes for one character + uint8_t char_bytes[4]; // holds utf-8 bytes for one character size_type char_width = from_char_utf8(ch, reinterpret_cast(char_bytes)); nbytes += char_width * 3; // '%' plus 2 hex chars per byte (example: é is %C3%A9) // process each byte in this current character diff --git a/cpp/src/strings/json/json_path.cu b/cpp/src/strings/json/json_path.cu index 2d2691e0518..c56752f5429 100644 --- a/cpp/src/strings/json/json_path.cu +++ b/cpp/src/strings/json/json_path.cu @@ -984,7 +984,7 @@ std::unique_ptr get_json_object(cudf::strings_column_view const& c col.size(), rmm::device_buffer{0, stream, mr}, // no data cudf::detail::create_null_mask(col.size(), mask_state::ALL_NULL, stream, mr), - col.size()); // null count + col.size()); // null count } constexpr int block_size = 512; diff --git a/cpp/src/strings/regex/regcomp.cpp b/cpp/src/strings/regex/regcomp.cpp index 5fd098a872e..b7a7f19369d 100644 --- a/cpp/src/strings/regex/regcomp.cpp +++ b/cpp/src/strings/regex/regcomp.cpp @@ -184,9 +184,9 @@ class regex_parser { int32_t _id_cclass_d{-1}; // digits [0-9] int32_t _id_cclass_D{-1}; // not digits - char32_t _chr{}; // last lex'd char - int32_t _cclass_id{}; // last lex'd class - int16_t _min_count{}; // data for counted operators + char32_t _chr{}; // last lex'd char + int32_t _cclass_id{}; // last lex'd class + int16_t _min_count{}; // data for counted operators int16_t _max_count{}; std::vector _items; @@ -361,9 +361,9 @@ class regex_parser { auto [q, n_chr] = next_char(); if (n_chr == 0) { return 0; } // malformed: '[x-' - if (!q && n_chr == ']') { // handles: '[x-]' + if (!q && n_chr == ']') { // handles: '[x-]' literals.push_back(chr); - literals.push_back(chr); // add '-' as literal + literals.push_back(chr); // add '-' as literal break; } // normal case: '[a-z]' @@ -749,7 +749,7 @@ class regex_parser { // infinite repeats if (n > 0) { // append '+' after last repetition out.push_back(regex_parser::Item{item.type == COUNTED ? PLUS : PLUS_LAZY, 0}); - } else { // copy it once then append '*' + } else { // copy it once then append '*' out.insert(out.end(), begin, end); out.push_back(regex_parser::Item{item.type == COUNTED ? STAR : STAR_LAZY, 0}); } @@ -1095,7 +1095,7 @@ void reprog::build_start_ids() ids.pop(); reinst const& inst = _insts[id]; if (inst.type == OR) { - if (inst.u2.left_id != id) // prevents infinite while-loop here + if (inst.u2.left_id != id) // prevents infinite while-loop here ids.push(inst.u2.left_id); if (inst.u1.right_id != id) // prevents infinite while-loop here ids.push(inst.u1.right_id); diff --git a/cpp/src/strings/regex/regcomp.h b/cpp/src/strings/regex/regcomp.h index aa2cb363b80..ab912ace0df 100644 --- a/cpp/src/strings/regex/regcomp.h +++ b/cpp/src/strings/regex/regcomp.h @@ -77,16 +77,16 @@ constexpr int32_t NCCLASS_D{1 << 5}; // not CCLASS_D or '\n' * @brief Structure of an encoded regex instruction */ struct reinst { - int32_t type; /* operator type or instruction type */ + int32_t type; /* operator type or instruction type */ union { int32_t cls_id; /* class pointer */ char32_t c; /* character */ int32_t subid; /* sub-expression id for RBRA and LBRA */ int32_t right_id; /* right child of OR */ } u1; - union { /* regexec relies on these two being in the same union */ - int32_t left_id; /* left child of OR */ - int32_t next_id; /* next instruction for CAT & LBRA */ + union { /* regexec relies on these two being in the same union */ + int32_t left_id; /* left child of OR */ + int32_t next_id; /* next instruction for CAT & LBRA */ } u2; int32_t reserved4; }; diff --git a/cpp/src/strings/regex/regex.cuh b/cpp/src/strings/regex/regex.cuh index 19d82380350..c1abbd78b43 100644 --- a/cpp/src/strings/regex/regex.cuh +++ b/cpp/src/strings/regex/regex.cuh @@ -253,21 +253,21 @@ class reprog_device { reprog_device(reprog const&); - int32_t _startinst_id; // first instruction id - int32_t _num_capturing_groups; // instruction groups - int32_t _insts_count; // number of instructions - int32_t _starts_count; // number of start-insts ids - int32_t _classes_count; // number of classes - int32_t _max_insts; // for partitioning working memory + int32_t _startinst_id; // first instruction id + int32_t _num_capturing_groups; // instruction groups + int32_t _insts_count; // number of instructions + int32_t _starts_count; // number of start-insts ids + int32_t _classes_count; // number of classes + int32_t _max_insts; // for partitioning working memory uint8_t const* _codepoint_flags{}; // table of character types reinst const* _insts{}; // array of regex instructions int32_t const* _startinst_ids{}; // array of start instruction ids reclass_device const* _classes{}; // array of regex classes - std::size_t _prog_size{}; // total size of this instance - void* _buffer{}; // working memory buffer - int32_t _thread_count{}; // threads available in working memory + std::size_t _prog_size{}; // total size of this instance + void* _buffer{}; // working memory buffer + int32_t _thread_count{}; // threads available in working memory }; /** diff --git a/cpp/src/strings/regex/regex.inl b/cpp/src/strings/regex/regex.inl index c5205ae7789..ce12dc17aa4 100644 --- a/cpp/src/strings/regex/regex.inl +++ b/cpp/src/strings/regex/regex.inl @@ -146,17 +146,17 @@ __device__ __forceinline__ bool reclass_device::is_match(char32_t const ch, uint32_t codept = utf8_to_codepoint(ch); if (codept > 0x00'FFFF) return false; int8_t fl = codepoint_flags[codept]; - if ((builtins & CCLASS_W) && ((ch == '_') || IS_ALPHANUM(fl))) // \w + if ((builtins & CCLASS_W) && ((ch == '_') || IS_ALPHANUM(fl))) // \w return true; - if ((builtins & CCLASS_S) && IS_SPACE(fl)) // \s + if ((builtins & CCLASS_S) && IS_SPACE(fl)) // \s return true; - if ((builtins & CCLASS_D) && IS_DIGIT(fl)) // \d + if ((builtins & CCLASS_D) && IS_DIGIT(fl)) // \d return true; if ((builtins & NCCLASS_W) && ((ch != '\n') && (ch != '_') && !IS_ALPHANUM(fl))) // \W return true; - if ((builtins & NCCLASS_S) && !IS_SPACE(fl)) // \S + if ((builtins & NCCLASS_S) && !IS_SPACE(fl)) // \S return true; - if ((builtins & NCCLASS_D) && ((ch != '\n') && !IS_DIGIT(fl))) // \D + if ((builtins & NCCLASS_D) && ((ch != '\n') && !IS_DIGIT(fl))) // \D return true; // return false; diff --git a/cpp/src/strings/replace/replace_re.cu b/cpp/src/strings/replace/replace_re.cu index 460074a5296..81ddb937be5 100644 --- a/cpp/src/strings/replace/replace_re.cu +++ b/cpp/src/strings/replace/replace_re.cu @@ -68,7 +68,7 @@ struct replace_regex_fn { if (!match) { break; } // no more matches auto const [start_pos, end_pos] = match_positions_to_bytes(*match, d_str, last_pos); - nbytes += d_repl.size_bytes() - (end_pos - start_pos); // add new size + nbytes += d_repl.size_bytes() - (end_pos - start_pos); // add new size if (out_ptr) { // replace: // i:bbbbsssseeee diff --git a/cpp/src/strings/split/partition.cu b/cpp/src/strings/split/partition.cu index 099f5978992..0c7d119ea38 100644 --- a/cpp/src/strings/split/partition.cu +++ b/cpp/src/strings/split/partition.cu @@ -170,7 +170,7 @@ struct rpartition_fn : public partition_fn { --itr; pos = check_delimiter(idx, d_str, itr); } - if (pos < 0) // delimiter not found + if (pos < 0) // delimiter not found { d_indices_left[idx] = string_index_pair{"", 0}; // two empty d_indices_delim[idx] = string_index_pair{"", 0}; // strings diff --git a/cpp/src/strings/split/split.cuh b/cpp/src/strings/split/split.cuh index e76d8ac1c60..dc0b04af388 100644 --- a/cpp/src/strings/split/split.cuh +++ b/cpp/src/strings/split/split.cuh @@ -190,7 +190,7 @@ struct split_tokenizer_fn : base_split_tokenizer { device_span d_delimiters, device_span d_tokens) const { - auto const base_ptr = get_base_ptr(); // d_positions values based on this + auto const base_ptr = get_base_ptr(); // d_positions values based on this auto str_ptr = d_str.data(); auto const str_end = str_ptr + d_str.size_bytes(); // end of the string auto const token_count = static_cast(d_tokens.size()); diff --git a/cpp/src/strings/split/split_re.cu b/cpp/src/strings/split/split_re.cu index 9aeb6b69bdc..3be5937297f 100644 --- a/cpp/src/strings/split/split_re.cu +++ b/cpp/src/strings/split/split_re.cu @@ -91,7 +91,7 @@ struct token_reader_fn { } else { if (direction == split_direction::FORWARD) { break; } // we are done for (auto l = 0; l < token_idx - 1; ++l) { - d_result[l] = d_result[l + 1]; // shift left + d_result[l] = d_result[l + 1]; // shift left } d_result[token_idx - 1] = token; } diff --git a/cpp/src/strings/utilities.cu b/cpp/src/strings/utilities.cu index 57a868485df..c8c68d19ce6 100644 --- a/cpp/src/strings/utilities.cu +++ b/cpp/src/strings/utilities.cu @@ -86,9 +86,9 @@ thread_safe_per_context_cache d_special_case_mappings; } // namespace - /** - * @copydoc cudf::strings::detail::get_character_flags_table - */ +/** + * @copydoc cudf::strings::detail::get_character_flags_table + */ character_flags_table_type const* get_character_flags_table() { return d_character_codepoint_flags.find_or_initialize([&](void) { diff --git a/cpp/src/text/normalize.cu b/cpp/src/text/normalize.cu index 78dfb6bf1a6..1b07b0785f5 100644 --- a/cpp/src/text/normalize.cu +++ b/cpp/src/text/normalize.cu @@ -70,7 +70,7 @@ struct normalize_spaces_fn { cudf::string_view const single_space(" ", 1); auto const d_str = d_strings.element(idx); char* buffer = d_chars ? d_chars + d_offsets[idx] : nullptr; - char* optr = buffer; // running output pointer + char* optr = buffer; // running output pointer cudf::size_type nbytes = 0; // holds the number of bytes per output string @@ -146,7 +146,7 @@ struct codepoint_to_utf8_fn { char* out_ptr = d_chars + d_offsets[idx]; for (uint32_t jdx = 0; jdx < count; ++jdx) { uint32_t code_point = *str_cps++; - if (code_point < UTF8_1BYTE) // ASCII range + if (code_point < UTF8_1BYTE) // ASCII range *out_ptr++ = static_cast(code_point); else if (code_point < UTF8_2BYTE) { // create two-byte UTF-8 // b00001xxx:byyyyyyyy => b110xxxyy:b10yyyyyy diff --git a/cpp/src/text/replace.cu b/cpp/src/text/replace.cu index d122f048a4e..34916e121dc 100644 --- a/cpp/src/text/replace.cu +++ b/cpp/src/text/replace.cu @@ -114,7 +114,7 @@ using strings_iterator = cudf::column_device_view::const_iterator= end) { break; } // done checking for pairs // skip to the next adjacent pair diff --git a/cpp/src/text/subword/load_merges_file.cu b/cpp/src/text/subword/load_merges_file.cu index 1f1b90b3f49..db6ad2e2dd2 100644 --- a/cpp/src/text/subword/load_merges_file.cu +++ b/cpp/src/text/subword/load_merges_file.cu @@ -93,7 +93,7 @@ std::unique_ptr initialize_merge_pairs_map( auto merge_pairs_map = std::make_unique( static_cast(input.size() * 2), // capacity is 2x; cuco::empty_key{-1}, - cuco::empty_value{-1}, // empty value is not used + cuco::empty_value{-1}, // empty value is not used bpe_equal{input}, probe_scheme{bpe_hasher{input}}, hash_table_allocator_type{default_allocator{}, stream}, diff --git a/cpp/src/text/utilities/tokenize_ops.cuh b/cpp/src/text/utilities/tokenize_ops.cuh index fbd2d1efcff..a84e94a6924 100644 --- a/cpp/src/text/utilities/tokenize_ops.cuh +++ b/cpp/src/text/utilities/tokenize_ops.cuh @@ -230,7 +230,7 @@ struct multi_delimiter_strings_tokenizer { }); if (itr_find != delimiters_end) { // found delimiter auto token_size = static_cast((curr_ptr - data_ptr) - last_pos); - if (token_size > 0) // we only care about non-zero sized tokens + if (token_size > 0) // we only care about non-zero sized tokens { if (d_str_tokens) d_str_tokens[token_idx] = string_index_pair{data_ptr + last_pos, token_size}; diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index a69dc9bf2f8..4923ef5c903 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -627,6 +627,7 @@ ConfigureTest(STREAM_CONCATENATE_TEST streams/concatenate_test.cpp STREAM_MODE t ConfigureTest(STREAM_FILLING_TEST streams/filling_test.cpp STREAM_MODE testing) ConfigureTest(STREAM_REPLACE_TEST streams/replace_test.cpp STREAM_MODE testing) ConfigureTest(STREAM_SEARCH_TEST streams/search_test.cpp STREAM_MODE testing) +ConfigureTest(STREAM_STRINGS_TEST streams/strings/case_test.cpp STREAM_MODE testing) # ################################################################################################## # Install tests #################################################################################### diff --git a/cpp/tests/groupby/merge_lists_tests.cpp b/cpp/tests/groupby/merge_lists_tests.cpp index 991473c5023..f2909f870aa 100644 --- a/cpp/tests/groupby/merge_lists_tests.cpp +++ b/cpp/tests/groupby/merge_lists_tests.cpp @@ -374,7 +374,7 @@ TEST_F(GroupbyMergeListsTest, StringsColumnInput) "" /*NULL*/, "" /*NULL*/, "German Shepherd", - "" /*NULL*/ + "" /*NULL*/ }, nulls_at({3, 4, 5, 7})}, // key = "dog" lists_col{{"Whale", "" /*NULL*/, "Polar Bear"}, null_at(1)}, // key = "unknown" diff --git a/cpp/tests/groupby/merge_sets_tests.cpp b/cpp/tests/groupby/merge_sets_tests.cpp index 67ff61563bb..5fc7e68b524 100644 --- a/cpp/tests/groupby/merge_sets_tests.cpp +++ b/cpp/tests/groupby/merge_sets_tests.cpp @@ -333,7 +333,7 @@ TEST_F(GroupbyMergeSetsTest, StringsColumnInput) lists_col{{"" /*NULL*/, "" /*NULL*/, "" /*NULL*/}, all_nulls()} // key = "dog" }; auto const lists3 = lists_col{ - lists_col{"Fuji", "Red Delicious"}, // key = "apple" + lists_col{"Fuji", "Red Delicious"}, // key = "apple" lists_col{{"" /*NULL*/, "Corgi", "German Shepherd", "" /*NULL*/, "Golden Retriever"}, nulls_at({0, 3})}, // key = "dog" lists_col{{"Seeedless", "Mini"}, no_nulls()} // key = "water melon" @@ -343,14 +343,14 @@ TEST_F(GroupbyMergeSetsTest, StringsColumnInput) merge_sets(vcol_views{keys1, keys2, keys3}, vcol_views{lists1, lists2, lists3}); auto const expected_keys = strings_col{"apple", "banana", "dog", "unknown", "water melon"}; auto const expected_lists = lists_col{ - lists_col{"Fuji", "Honey Bee", "Red Delicious"}, // key = "apple" - lists_col{"Green", "Yellow"}, // key = "banana" + lists_col{"Fuji", "Honey Bee", "Red Delicious"}, // key = "apple" + lists_col{"Green", "Yellow"}, // key = "banana" lists_col{{ "Corgi", "German Shepherd", "Golden Retriever", "Poodle", "" /*NULL*/ }, - null_at(4)}, // key = "dog" - lists_col{{"Polar Bear", "Whale", "" /*NULL*/}, null_at(2)}, // key = "unknown" - lists_col{{"Mini", "Seeedless"}, no_nulls()} // key = "water melon" + null_at(4)}, // key = "dog" + lists_col{{"Polar Bear", "Whale", "" /*NULL*/}, null_at(2)}, // key = "unknown" + lists_col{{"Mini", "Seeedless"}, no_nulls()} // key = "water melon" }; CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_keys, *out_keys, verbosity); diff --git a/cpp/tests/io/orc_test.cpp b/cpp/tests/io/orc_test.cpp index cff7b1cf081..890ef914713 100644 --- a/cpp/tests/io/orc_test.cpp +++ b/cpp/tests/io/orc_test.cpp @@ -976,6 +976,10 @@ TEST_F(OrcReaderTest, CombinedSkipRowTest) TEST_F(OrcStatisticsTest, Basic) { auto sequence = cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i; }); + auto ts_sequence = + cudf::detail::make_counting_transform_iterator(0, [](auto i) { return (i - 4) * 1000002; }); + auto dec_sequence = + cudf::detail::make_counting_transform_iterator(0, [&](auto i) { return i * 1001; }); auto validity = cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i % 2; }); std::vector strings{ @@ -986,11 +990,17 @@ TEST_F(OrcStatisticsTest, Basic) sequence, sequence + num_rows, validity); column_wrapper col2( sequence, sequence + num_rows, validity); - column_wrapper col3{strings.begin(), strings.end()}; - column_wrapper col4(sequence, sequence + num_rows); - column_wrapper col5( - sequence, sequence + num_rows, validity); - table_view expected({col1, col2, col3, col4, col5}); + str_col col3{strings.begin(), strings.end()}; + column_wrapper col4( + ts_sequence, ts_sequence + num_rows, validity); + column_wrapper col5( + ts_sequence, ts_sequence + num_rows, validity); + bool_col col6({true, true, true, true, true, false, false, false, false}, validity); + + cudf::test::fixed_point_column_wrapper col7( + dec_sequence, dec_sequence + num_rows, numeric::scale_type{-1}); + + table_view expected({col1, col2, col3, col4, col5, col6, col7}); auto filepath = temp_env->get_temp_filepath("OrcStatsMerge.orc"); @@ -1000,16 +1010,21 @@ TEST_F(OrcStatisticsTest, Basic) auto const stats = cudf::io::read_parsed_orc_statistics(cudf::io::source_info{filepath}); - auto const expected_column_names = - std::vector{"", "_col0", "_col1", "_col2", "_col3", "_col4"}; + auto expected_column_names = std::vector{""}; + std::generate_n( + std::back_inserter(expected_column_names), + expected.num_columns(), + [starting_index = 0]() mutable { return "_col" + std::to_string(starting_index++); }); EXPECT_EQ(stats.column_names, expected_column_names); auto validate_statistics = [&](std::vector const& stats) { + ASSERT_EQ(stats.size(), expected.num_columns() + 1); auto& s0 = stats[0]; EXPECT_EQ(*s0.number_of_values, 9ul); auto& s1 = stats[1]; EXPECT_EQ(*s1.number_of_values, 4ul); + EXPECT_TRUE(*s1.has_null); auto& ts1 = std::get(s1.type_specific_stats); EXPECT_EQ(*ts1.minimum, 1); EXPECT_EQ(*ts1.maximum, 7); @@ -1017,30 +1032,55 @@ TEST_F(OrcStatisticsTest, Basic) auto& s2 = stats[2]; EXPECT_EQ(*s2.number_of_values, 4ul); + EXPECT_TRUE(*s2.has_null); auto& ts2 = std::get(s2.type_specific_stats); EXPECT_EQ(*ts2.minimum, 1.); EXPECT_EQ(*ts2.maximum, 7.); - // No sum ATM, filed #7087 - ASSERT_FALSE(ts2.sum); + EXPECT_EQ(*ts2.sum, 16.); auto& s3 = stats[3]; EXPECT_EQ(*s3.number_of_values, 9ul); + EXPECT_FALSE(*s3.has_null); auto& ts3 = std::get(s3.type_specific_stats); EXPECT_EQ(*ts3.minimum, "Friday"); EXPECT_EQ(*ts3.maximum, "Wednesday"); EXPECT_EQ(*ts3.sum, 58ul); auto& s4 = stats[4]; - EXPECT_EQ(*s4.number_of_values, 9ul); - EXPECT_EQ(std::get(s4.type_specific_stats).count[0], 8ul); + EXPECT_EQ(*s4.number_of_values, 4ul); + EXPECT_TRUE(*s4.has_null); + auto& ts4 = std::get(s4.type_specific_stats); + EXPECT_EQ(*ts4.minimum, -4); + EXPECT_EQ(*ts4.maximum, 3); + EXPECT_EQ(*ts4.minimum_utc, -4); + EXPECT_EQ(*ts4.maximum_utc, 3); + EXPECT_EQ(*ts4.minimum_nanos, 999994); + EXPECT_EQ(*ts4.maximum_nanos, 6); auto& s5 = stats[5]; EXPECT_EQ(*s5.number_of_values, 4ul); + EXPECT_TRUE(*s5.has_null); auto& ts5 = std::get(s5.type_specific_stats); - EXPECT_EQ(*ts5.minimum_utc, 1000); - EXPECT_EQ(*ts5.maximum_utc, 7000); - ASSERT_FALSE(ts5.minimum); - ASSERT_FALSE(ts5.maximum); + EXPECT_EQ(*ts5.minimum, -3001); + EXPECT_EQ(*ts5.maximum, 3000); + EXPECT_EQ(*ts5.minimum_utc, -3001); + EXPECT_EQ(*ts5.maximum_utc, 3000); + EXPECT_EQ(*ts5.minimum_nanos, 994000); + EXPECT_EQ(*ts5.maximum_nanos, 6000); + + auto& s6 = stats[6]; + EXPECT_EQ(*s6.number_of_values, 4ul); + EXPECT_TRUE(*s6.has_null); + auto& ts6 = std::get(s6.type_specific_stats); + EXPECT_EQ(ts6.count[0], 2); + + auto& s7 = stats[7]; + EXPECT_EQ(*s7.number_of_values, 9ul); + EXPECT_FALSE(*s7.has_null); + auto& ts7 = std::get(s7.type_specific_stats); + EXPECT_EQ(*ts7.minimum, "0.0"); + EXPECT_EQ(*ts7.maximum, "800.8"); + EXPECT_EQ(*ts7.sum, "3603.6"); }; validate_statistics(stats.file_stats); @@ -1259,9 +1299,8 @@ TEST_F(OrcStatisticsTest, Overflow) TEST_F(OrcStatisticsTest, HasNull) { - // cudf's ORC writer doesn't yet support the ability to encode the hasNull value in statistics so - // we're embedding a file created using pyorc - // + // This test can now be implemented with libcudf; keeping the pyorc version to keep the test + // inputs diversified // Method to create file: // >>> import pyorc // >>> output = open("./temp.orc", "wb") @@ -1861,4 +1900,38 @@ TEST_F(OrcWriterTest, EmptyChildStringColumn) CUDF_TEST_EXPECT_TABLES_EQUAL(expected, result.tbl->view()); } +template +void check_all_null_stats(cudf::io::column_statistics const& stats) +{ + EXPECT_EQ(stats.number_of_values, 0); + EXPECT_TRUE(stats.has_null); + + auto const ts = std::get(stats.type_specific_stats); + EXPECT_FALSE(ts.minimum.has_value()); + EXPECT_FALSE(ts.maximum.has_value()); + EXPECT_TRUE(ts.sum.has_value()); + EXPECT_EQ(*ts.sum, 0); +} + +TEST_F(OrcStatisticsTest, AllNulls) +{ + float64_col double_col({0., 0., 0.}, cudf::test::iterators::all_nulls()); + int32_col int_col({0, 0, 0}, cudf::test::iterators::all_nulls()); + str_col string_col({"", "", ""}, cudf::test::iterators::all_nulls()); + + cudf::table_view expected({int_col, double_col, string_col}); + + std::vector out_buffer; + cudf::io::orc_writer_options out_opts = + cudf::io::orc_writer_options::builder(cudf::io::sink_info{&out_buffer}, expected); + cudf::io::write_orc(out_opts); + + auto const stats = cudf::io::read_parsed_orc_statistics( + cudf::io::source_info{out_buffer.data(), out_buffer.size()}); + + check_all_null_stats(stats.file_stats[1]); + check_all_null_stats(stats.file_stats[2]); + check_all_null_stats(stats.file_stats[3]); +} + CUDF_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/io/parquet_test.cpp b/cpp/tests/io/parquet_test.cpp index 64aca091686..81e0e12eeb9 100644 --- a/cpp/tests/io/parquet_test.cpp +++ b/cpp/tests/io/parquet_test.cpp @@ -2166,7 +2166,7 @@ TEST_F(ParquetChunkedWriterTest, ForcedNullabilityList) cudf::io::table_input_metadata metadata(table1); metadata.column_metadata[0].set_nullability(true); // List is nullable at first (root) level metadata.column_metadata[0].child(1).set_nullability( - false); // non-nullable at second (leaf) level + false); // non-nullable at second (leaf) level metadata.column_metadata[1].set_nullability(true); auto filepath = temp_env->get_temp_filepath("ChunkedListNullable.parquet"); @@ -5880,7 +5880,7 @@ TEST_F(ParquetMetadataReaderTest, TestNested) EXPECT_EQ(out_map_col.type_kind(), cudf::io::parquet::TypeKind::UNDEFINED_TYPE); // map ASSERT_EQ(out_map_col.num_children(), 1); - EXPECT_EQ(out_map_col.child(0).name(), "key_value"); // key_value (named in parquet writer) + EXPECT_EQ(out_map_col.child(0).name(), "key_value"); // key_value (named in parquet writer) ASSERT_EQ(out_map_col.child(0).num_children(), 2); EXPECT_EQ(out_map_col.child(0).child(0).name(), "key"); // key (named in parquet writer) EXPECT_EQ(out_map_col.child(0).child(1).name(), "value"); // value (named in parquet writer) @@ -5897,7 +5897,7 @@ TEST_F(ParquetMetadataReaderTest, TestNested) ASSERT_EQ(out_list_col.child(0).num_children(), 1); auto const& out_list_struct_col = out_list_col.child(0).child(0); - EXPECT_EQ(out_list_struct_col.name(), "element"); // elements (named in parquet writer) + EXPECT_EQ(out_list_struct_col.name(), "element"); // elements (named in parquet writer) EXPECT_EQ(out_list_struct_col.type_kind(), cudf::io::parquet::TypeKind::UNDEFINED_TYPE); // struct ASSERT_EQ(out_list_struct_col.num_children(), 2); diff --git a/cpp/tests/lists/reverse_tests.cpp b/cpp/tests/lists/reverse_tests.cpp index a899d387c3e..00dc13c5812 100644 --- a/cpp/tests/lists/reverse_tests.cpp +++ b/cpp/tests/lists/reverse_tests.cpp @@ -370,8 +370,8 @@ TYPED_TEST(ListsReverseTypedTest, InputListsOfStructsWithNulls) "Kiwi", "Cherry", "Banana", - "", /*NULL*/ - "", /*NULL*/ + "", /*NULL*/ + "", /*NULL*/ "Apple", "", /*NULL*/ "Banana", // end list1 @@ -436,8 +436,8 @@ TYPED_TEST(ListsReverseTypedTest, InputListsOfStructsWithNulls) "Kiwi", "Cherry", "Banana", - "", /*NULL*/ - "", /*NULL*/ + "", /*NULL*/ + "", /*NULL*/ "Apple", "", /*NULL*/ "Banana", // end list1 diff --git a/cpp/tests/lists/set_operations/difference_distinct_tests.cpp b/cpp/tests/lists/set_operations/difference_distinct_tests.cpp index bf7ebc902ba..84c51f256b7 100644 --- a/cpp/tests/lists/set_operations/difference_distinct_tests.cpp +++ b/cpp/tests/lists/set_operations/difference_distinct_tests.cpp @@ -571,7 +571,7 @@ TEST_F(SetDifferenceTest, InputListsOfNestedStructsHaveNull) "" /*NULL*/, "" /*NULL*/, "" /*NULL*/, "Apple", "Banana", "Cherry", "Kiwi", // end list1 "" /*NULL*/, "Bear", "Cat", "Dog", "Duck", - "Panda", // end list2 + "Panda", // end list2 "ÁÁÁ", "ÉÉÉÉÉ", "ÁBC", "ÁÁÁ", "ÍÍÍÍÍ", "" /*NULL*/, "XYZ", "ÁBC" // end list3 diff --git a/cpp/tests/lists/set_operations/intersect_distinct_tests.cpp b/cpp/tests/lists/set_operations/intersect_distinct_tests.cpp index dbccf06036b..11f98af3520 100644 --- a/cpp/tests/lists/set_operations/intersect_distinct_tests.cpp +++ b/cpp/tests/lists/set_operations/intersect_distinct_tests.cpp @@ -514,7 +514,7 @@ TEST_F(SetIntersectTest, InputListsOfNestedStructsHaveNull) null, // end list1 null, // end list2 null, - null // end list3 + null // end list3 }, all_nulls()}; auto grandchild2 = strings_col{{ @@ -522,7 +522,7 @@ TEST_F(SetIntersectTest, InputListsOfNestedStructsHaveNull) "Apple", // end list1 "" /*NULL*/, // end list2 "ÁÁÁ", - "ÉÉÉÉÉ" // end list3 + "ÉÉÉÉÉ" // end list3 }, nulls_at({0, 2})}; auto child1 = structs_col{{grandchild1, grandchild2}, null_at(0)}; diff --git a/cpp/tests/lists/set_operations/union_distinct_tests.cpp b/cpp/tests/lists/set_operations/union_distinct_tests.cpp index 5cc0897351d..e33ea31541b 100644 --- a/cpp/tests/lists/set_operations/union_distinct_tests.cpp +++ b/cpp/tests/lists/set_operations/union_distinct_tests.cpp @@ -560,7 +560,7 @@ TEST_F(SetUnionTest, InputListsOfNestedStructsHaveNull) auto grandchild2 = strings_col{{ "" /*NULL*/, "Apple", "Banana", "Cherry", "Kiwi", "Banana", "Cherry", - "Kiwi", // end list1 + "Kiwi", // end list1 "" /*NULL*/, "Bear", "Cat", "Dog", "Duck", "Panda", "Bear", "Cat", "Dog", "Duck", "Panda", // end list2 @@ -597,7 +597,7 @@ TEST_F(SetUnionTest, InputListsOfNestedStructsHaveNull) { "" /*NULL*/, "" /*NULL*/, "" /*NULL*/, "" /*NULL*/, "" /*NULL*/, "" /*NULL*/, "Apple", "Apple", "Banana", "Cherry", "Kiwi", "Banana", "Cherry", - "Kiwi", // end list1 + "Kiwi", // end list1 "" /*NULL*/, "" /*NULL*/, "Bear", "Cat", "Dog", "Duck", "Panda", "Bear", "Cat", "Dog", "Duck", "Panda", // end list2 "ÁÁÁ", "ÁÁÁ", "ÉÉÉÉÉ", "ÉÉÉÉÉ", "ÁBC", "ÁÁÁ", "ÍÍÍÍÍ", diff --git a/cpp/tests/lists/stream_compaction/distinct_tests.cpp b/cpp/tests/lists/stream_compaction/distinct_tests.cpp index 57d1714c255..fbc637f9315 100644 --- a/cpp/tests/lists/stream_compaction/distinct_tests.cpp +++ b/cpp/tests/lists/stream_compaction/distinct_tests.cpp @@ -529,7 +529,7 @@ TEST_F(ListDistinctTest, InputListsOfStructsHaveNull) 2, 3, 3, - 3}, // end list3 + 3}, // end list3 nulls_at({1, 6, 12, 13})}; auto child2 = strings_col{{ // begin list1 "XXX", /*NULL*/ @@ -551,7 +551,7 @@ TEST_F(ListDistinctTest, InputListsOfStructsHaveNull) "ÁBC", "ÁÁÁ", "ÍÍÍÍÍ", - "", /*NULL*/ + "", /*NULL*/ "XYZ", "ÁBC"}, // end list3 nulls_at({6, 17})}; @@ -670,7 +670,7 @@ TEST_F(ListDistinctTest, InputListsOfNestedStructsHaveNull) "ÁBC", "ÁÁÁ", "ÍÍÍÍÍ", - "", /*NULL*/ + "", /*NULL*/ "XYZ", "ÁBC" // end list3 }, @@ -729,8 +729,8 @@ TEST_F(ListDistinctTest, InputListsOfStructsOfLists) floats_lists{3, 4, 5}, // end list2 // begin list3 floats_lists{}, - floats_lists{}, // end list3 - // begin list4 + floats_lists{}, // end list3 + // begin list4 floats_lists{6, 7}, floats_lists{6, 7}, floats_lists{6, 7}}; diff --git a/cpp/tests/reshape/interleave_columns_tests.cpp b/cpp/tests/reshape/interleave_columns_tests.cpp index eba6c961bbb..e8ea9d619c5 100644 --- a/cpp/tests/reshape/interleave_columns_tests.cpp +++ b/cpp/tests/reshape/interleave_columns_tests.cpp @@ -806,7 +806,7 @@ TYPED_TEST(ListsColumnsInterleaveTypedTest, SlicedInputListsOfListsWithNulls) ListsCol{ListsCol{{null, 11}, null_at(0)}, ListsCol{{22, null, null}, nulls_at({1, 2})}}, // don't care ListsCol{ListsCol{{null, 11}, null_at(0)}, - ListsCol{{22, null, null}, nulls_at({1, 2})}} // don't care + ListsCol{{22, null, null}, nulls_at({1, 2})}} // don't care }; auto const col1 = cudf::slice(col1_original, {3, 6})[0]; diff --git a/cpp/tests/rolling/range_rolling_window_test.cpp b/cpp/tests/rolling/range_rolling_window_test.cpp index 585383f28f8..eed9db1fe04 100644 --- a/cpp/tests/rolling/range_rolling_window_test.cpp +++ b/cpp/tests/rolling/range_rolling_window_test.cpp @@ -91,7 +91,7 @@ struct window_exec { ScalarT preceding; // Preceding window scalar. ScalarT following; // Following window scalar. cudf::size_type min_periods = 1; -}; // struct window_exec; +}; // struct window_exec; struct RangeRollingTest : public cudf::test::BaseFixture {}; diff --git a/cpp/tests/sort/segmented_sort_tests.cpp b/cpp/tests/sort/segmented_sort_tests.cpp index b3f98eb54b9..da9666cbc74 100644 --- a/cpp/tests/sort/segmented_sort_tests.cpp +++ b/cpp/tests/sort/segmented_sort_tests.cpp @@ -270,7 +270,7 @@ TEST_F(SegmentedSortInt, Sliced) column_wrapper expected2{{0, 1, 3, 2, 4, 5, 6}}; column_wrapper expected3{{0, 1, 2, 3, 4, 5, 6}}; // clang-format on - auto slice = cudf::slice(col1, {4, 11})[0]; // 7 elements + auto slice = cudf::slice(col1, {4, 11})[0]; // 7 elements cudf::table_view input{{slice}}; auto seg_slice = cudf::slice(segments2, {2, 4})[0]; // 2 elements diff --git a/cpp/tests/streams/strings/case_test.cpp b/cpp/tests/streams/strings/case_test.cpp new file mode 100644 index 00000000000..df3eabd773a --- /dev/null +++ b/cpp/tests/streams/strings/case_test.cpp @@ -0,0 +1,55 @@ +/* + * Copyright (c) 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. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include +#include + +#include +#include +#include + +class StringsCaseTest : public cudf::test::BaseFixture {}; + +TEST_F(StringsCaseTest, LowerUpper) +{ + auto const input = + cudf::test::strings_column_wrapper({"", + "The quick brown fox", + "jumps over the lazy dog.", + "all work and no play makes Jack a dull boy", + R"(!"#$%&'()*+,-./0123456789:;<=>?@[\]^_`{|}~)"}); + auto view = cudf::strings_column_view(input); + + cudf::strings::to_lower(view, cudf::test::get_default_stream()); + cudf::strings::to_upper(view, cudf::test::get_default_stream()); + cudf::strings::swapcase(view, cudf::test::get_default_stream()); +} + +TEST_F(StringsCaseTest, Capitalize) +{ + auto const input = + cudf::test::strings_column_wrapper({"", + "The Quick Brown Fox", + "jumps over the lazy dog", + "all work and no play makes Jack a dull boy"}); + auto view = cudf::strings_column_view(input); + + auto const delimiter = cudf::string_scalar(" ", true, cudf::test::get_default_stream()); + cudf::strings::capitalize(view, delimiter, cudf::test::get_default_stream()); + cudf::strings::is_title(view, cudf::test::get_default_stream()); + cudf::strings::title( + view, cudf::strings::string_character_types::ALPHA, cudf::test::get_default_stream()); +} diff --git a/cpp/tests/strings/chars_types_tests.cpp b/cpp/tests/strings/chars_types_tests.cpp index a16da41af7a..c595977c269 100644 --- a/cpp/tests/strings/chars_types_tests.cpp +++ b/cpp/tests/strings/chars_types_tests.cpp @@ -50,17 +50,17 @@ TEST_P(CharsTypes, AllTypes) "\t\r\n\f "}; bool expecteds[] = {false, false, false, false, false, false, false, false, - false, false, false, false, false, true, false, false, // decimal + false, false, false, false, false, true, false, false, // decimal false, false, false, false, false, false, false, false, - false, true, false, true, false, true, false, false, // numeric + false, true, false, true, false, true, false, false, // numeric false, false, false, false, false, false, false, false, - false, false, false, true, false, true, false, false, // digit + false, false, false, true, false, true, false, false, // digit true, true, false, true, false, false, false, false, - false, false, false, false, false, false, true, false, // alpha + false, false, false, false, false, false, true, false, // alpha false, false, false, false, false, false, false, false, - false, false, false, false, false, false, false, true, // space + false, false, false, false, false, false, false, true, // space false, false, false, true, false, false, false, false, - false, false, false, false, false, false, false, false, // upper + false, false, false, false, false, false, false, false, // upper false, true, false, false, false, false, false, false, false, false, false, false, false, false, true, false}; // lower diff --git a/cpp/tests/strings/durations_tests.cpp b/cpp/tests/strings/durations_tests.cpp index 0c7a1ad8042..1902f907f43 100644 --- a/cpp/tests/strings/durations_tests.cpp +++ b/cpp/tests/strings/durations_tests.cpp @@ -398,7 +398,7 @@ TEST_F(StringsDurationsTest, ParseSingle) "-59", "999", "-999", - "", // error + "", // error "01", ""}; // error auto size = cudf::column_view(string_src).size(); @@ -449,7 +449,7 @@ TEST_F(StringsDurationsTest, ParseMultiple) "-59:00:00", "999:00:00", "-999:00:00", - "", // error + "", // error "01:01:01", ""}; // error auto size = cudf::column_view(string_src).size(); @@ -503,7 +503,7 @@ TEST_F(StringsDurationsTest, ParseSubsecond) "-59:00:00", "999:00:00", "-999:00:00", - "", // error + "", // error "01:01:01", ""}; // error auto size = cudf::column_view(string_src).size(); @@ -660,7 +660,7 @@ TEST_F(StringsDurationsTest, ParseCompoundSpecifier) "09:00 AM", // error "", // error "01:01:01", - ""}; // error + ""}; // error cudf::test::fixed_width_column_wrapper expected_s3( {0, diff --git a/cpp/tests/utilities/column_utilities.cu b/cpp/tests/utilities/column_utilities.cu index bae402155e9..620e0bfe8de 100644 --- a/cpp/tests/utilities/column_utilities.cu +++ b/cpp/tests/utilities/column_utilities.cu @@ -440,7 +440,7 @@ class corresponding_rows_not_equivalent { // Must handle inf and nan separately if (std::isinf(x) || std::isinf(y)) { - return x != y; // comparison of (inf==inf) returns true + return x != y; // comparison of (inf==inf) returns true } else if (std::isnan(x) || std::isnan(y)) { return std::isnan(x) != std::isnan(y); // comparison of (nan==nan) returns false } else { diff --git a/python/cudf/cudf/tests/test_orc.py b/python/cudf/cudf/tests/test_orc.py index aafc8831bf4..07aa5430f4f 100644 --- a/python/cudf/cudf/tests/test_orc.py +++ b/python/cudf/cudf/tests/test_orc.py @@ -633,16 +633,19 @@ def test_orc_write_statistics(tmpdir, datadir, nrows, stats_freq): for col in gdf: if "minimum" in file_stats[0][col]: stats_min = file_stats[0][col]["minimum"] - actual_min = gdf[col].min() - assert normalized_equals(actual_min, stats_min) + if stats_min is not None: + actual_min = gdf[col].min() + assert normalized_equals(actual_min, stats_min) if "maximum" in file_stats[0][col]: stats_max = file_stats[0][col]["maximum"] - actual_max = gdf[col].max() - assert normalized_equals(actual_max, stats_max) + if stats_max is not None: + actual_max = gdf[col].max() + assert normalized_equals(actual_max, stats_max) if "number_of_values" in file_stats[0][col]: stats_num_vals = file_stats[0][col]["number_of_values"] - actual_num_vals = gdf[col].count() - assert stats_num_vals == actual_num_vals + if stats_num_vals is not None: + actual_num_vals = gdf[col].count() + assert stats_num_vals == actual_num_vals # compare stripe statistics with actual min/max for stripe_idx in range(0, orc_file.nstripes): @@ -651,21 +654,24 @@ def test_orc_write_statistics(tmpdir, datadir, nrows, stats_freq): stripe_df = cudf.DataFrame(stripe.to_pandas()) for col in stripe_df: if "minimum" in stripes_stats[stripe_idx][col]: - actual_min = stripe_df[col].min() stats_min = stripes_stats[stripe_idx][col]["minimum"] - assert normalized_equals(actual_min, stats_min) + if stats_min is not None: + actual_min = stripe_df[col].min() + assert normalized_equals(actual_min, stats_min) if "maximum" in stripes_stats[stripe_idx][col]: - actual_max = stripe_df[col].max() stats_max = stripes_stats[stripe_idx][col]["maximum"] - assert normalized_equals(actual_max, stats_max) + if stats_max is not None: + actual_max = stripe_df[col].max() + assert normalized_equals(actual_max, stats_max) if "number_of_values" in stripes_stats[stripe_idx][col]: stats_num_vals = stripes_stats[stripe_idx][col][ "number_of_values" ] - actual_num_vals = stripe_df[col].count() - assert stats_num_vals == actual_num_vals + if stats_num_vals is not None: + actual_num_vals = stripe_df[col].count() + assert stats_num_vals == actual_num_vals @pytest.mark.parametrize("stats_freq", ["STRIPE", "ROWGROUP"]) @@ -733,16 +739,19 @@ def test_orc_chunked_write_statistics(tmpdir, datadir, nrows, stats_freq): for col in expect: if "minimum" in file_stats[0][col]: stats_min = file_stats[0][col]["minimum"] - actual_min = expect[col].min() - assert normalized_equals(actual_min, stats_min) + if stats_min is not None: + actual_min = expect[col].min() + assert normalized_equals(actual_min, stats_min) if "maximum" in file_stats[0][col]: stats_max = file_stats[0][col]["maximum"] - actual_max = expect[col].max() - assert normalized_equals(actual_max, stats_max) + if stats_max is not None: + actual_max = expect[col].max() + assert normalized_equals(actual_max, stats_max) if "number_of_values" in file_stats[0][col]: stats_num_vals = file_stats[0][col]["number_of_values"] - actual_num_vals = expect[col].count() - assert stats_num_vals == actual_num_vals + if stats_num_vals is not None: + actual_num_vals = expect[col].count() + assert stats_num_vals == actual_num_vals # compare stripe statistics with actual min/max for stripe_idx in range(0, orc_file.nstripes): @@ -751,21 +760,24 @@ def test_orc_chunked_write_statistics(tmpdir, datadir, nrows, stats_freq): stripe_df = cudf.DataFrame(stripe.to_pandas()) for col in stripe_df: if "minimum" in stripes_stats[stripe_idx][col]: - actual_min = stripe_df[col].min() stats_min = stripes_stats[stripe_idx][col]["minimum"] - assert normalized_equals(actual_min, stats_min) + if stats_min is not None: + actual_min = stripe_df[col].min() + assert normalized_equals(actual_min, stats_min) if "maximum" in stripes_stats[stripe_idx][col]: - actual_max = stripe_df[col].max() stats_max = stripes_stats[stripe_idx][col]["maximum"] - assert normalized_equals(actual_max, stats_max) + if stats_max is not None: + actual_max = stripe_df[col].max() + assert normalized_equals(actual_max, stats_max) if "number_of_values" in stripes_stats[stripe_idx][col]: stats_num_vals = stripes_stats[stripe_idx][col][ "number_of_values" ] - actual_num_vals = stripe_df[col].count() - assert stats_num_vals == actual_num_vals + if stats_num_vals is not None: + actual_num_vals = stripe_df[col].count() + assert stats_num_vals == actual_num_vals @pytest.mark.parametrize("nrows", [1, 100, 6000000])