From 07e6cc170b46d9e84b01d42df4788813435d9896 Mon Sep 17 00:00:00 2001 From: Karthikeyan Natarajan Date: Wed, 5 Oct 2022 09:46:10 +0530 Subject: [PATCH 1/4] fix decimal benchmark input data generation (scale also was in data) --- cpp/benchmarks/common/generate_input.cu | 49 ++++++++++++++----------- 1 file changed, 28 insertions(+), 21 deletions(-) diff --git a/cpp/benchmarks/common/generate_input.cu b/cpp/benchmarks/common/generate_input.cu index 890a78bb9bf..e2b271f67c0 100644 --- a/cpp/benchmarks/common/generate_input.cu +++ b/cpp/benchmarks/common/generate_input.cu @@ -247,39 +247,43 @@ struct random_value_fn()>> { */ template struct random_value_fn()>> { - using rep = typename T::rep; - rep const lower_bound; - rep const upper_bound; + using DeviceType = cudf::device_storage_type_t; + using rep = typename T::rep; + DeviceType const lower_bound; + DeviceType const upper_bound; distribution_fn dist; std::optional scale; - random_value_fn(distribution_params const& desc) + random_value_fn(distribution_params const& desc) : lower_bound{desc.lower_bound}, upper_bound{desc.upper_bound}, - dist{make_distribution(desc.id, desc.lower_bound, desc.upper_bound)} + dist{make_distribution(desc.id, desc.lower_bound, desc.upper_bound)} { } - rmm::device_uvector operator()(thrust::minstd_rand& engine, unsigned size) + [[nodiscard]] numeric::scale_type get_scale() const + { + return scale.value_or(numeric::scale_type{0}); + } + + rmm::device_uvector operator()(thrust::minstd_rand& engine, unsigned size) { if (not scale.has_value()) { - int const max_scale = std::numeric_limits::digits10; + int const max_scale = std::numeric_limits::digits10; std::uniform_int_distribution scale_dist{-max_scale, max_scale}; std::mt19937 engine_scale(engine()); scale = numeric::scale_type{scale_dist(engine_scale)}; } auto const ints = dist(engine, size); - rmm::device_uvector result(size, cudf::default_stream_value); + rmm::device_uvector result(size, cudf::default_stream_value); // Clamp the generated random value to the specified range - thrust::transform(thrust::device, - ints.begin(), - ints.end(), - result.begin(), - [scale = *(this->scale), - upper_bound = this->upper_bound, - lower_bound = this->lower_bound] __device__(auto int_value) { - return T{std::clamp(int_value, lower_bound, upper_bound), scale}; - }); + thrust::transform( + thrust::device, + ints.begin(), + ints.end(), + result.begin(), + [upper_bound = this->upper_bound, lower_bound = this->lower_bound] __device__( + auto int_value) { return std::clamp(int_value, lower_bound, upper_bound); }); return result; } }; @@ -393,6 +397,8 @@ std::unique_ptr create_random_column(data_profile const& profile, thrust::minstd_rand& engine, cudf::size_type num_rows) { + using DeviceType = cudf::device_storage_type_t; + numeric::scale_type scale{}; // Bernoulli distribution auto valid_dist = random_value_fn( distribution_params{1. - profile.get_null_probability().value_or(0)}); @@ -400,7 +406,7 @@ std::unique_ptr create_random_column(data_profile const& profile, // Distribution for picking elements from the array of samples auto const avg_run_len = profile.get_avg_run_length(); - rmm::device_uvector data(0, cudf::default_stream_value); + rmm::device_uvector data(0, cudf::default_stream_value); rmm::device_uvector null_mask(0, cudf::default_stream_value); if (profile.get_cardinality() == 0 and avg_run_len == 1) { @@ -412,11 +418,12 @@ std::unique_ptr create_random_column(data_profile const& profile, : profile_cardinality; }(); rmm::device_uvector samples_null_mask = valid_dist(engine, cardinality); - rmm::device_uvector samples = value_dist(engine, cardinality); + rmm::device_uvector samples = value_dist(engine, cardinality); + if constexpr (cudf::is_fixed_point()) { scale = value_dist.get_scale(); } // generate n samples and gather. auto const sample_indices = sample_indices_with_run_length(avg_run_len, cardinality, num_rows, engine); - data = rmm::device_uvector(num_rows, cudf::default_stream_value); + data = rmm::device_uvector(num_rows, cudf::default_stream_value); null_mask = rmm::device_uvector(num_rows, cudf::default_stream_value); thrust::gather( thrust::device, sample_indices.begin(), sample_indices.end(), samples.begin(), data.begin()); @@ -431,7 +438,7 @@ std::unique_ptr create_random_column(data_profile const& profile, cudf::detail::valid_if(null_mask.begin(), null_mask.end(), thrust::identity{}); return std::make_unique( - cudf::data_type{cudf::type_to_id()}, + cudf::data_type{cudf::type_to_id(), scale}, num_rows, data.release(), profile.get_null_probability().has_value() ? std::move(result_bitmask) : rmm::device_buffer{}); From 0f36e7e2f14bc626a9a1e54a041ee371407c1fb2 Mon Sep 17 00:00:00 2001 From: Karthikeyan Natarajan Date: Thu, 6 Oct 2022 11:24:09 +0530 Subject: [PATCH 2/4] remove T::rep usage completely --- cpp/benchmarks/common/generate_input.cu | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/cpp/benchmarks/common/generate_input.cu b/cpp/benchmarks/common/generate_input.cu index e2b271f67c0..ef65177c941 100644 --- a/cpp/benchmarks/common/generate_input.cu +++ b/cpp/benchmarks/common/generate_input.cu @@ -248,10 +248,9 @@ struct random_value_fn()>> { template struct random_value_fn()>> { using DeviceType = cudf::device_storage_type_t; - using rep = typename T::rep; DeviceType const lower_bound; DeviceType const upper_bound; - distribution_fn dist; + distribution_fn dist; std::optional scale; random_value_fn(distribution_params const& desc) From 91bbdfe2c1aa197831586e0ec7cb8165e473fc32 Mon Sep 17 00:00:00 2001 From: Karthikeyan Natarajan Date: Thu, 6 Oct 2022 23:21:47 +0530 Subject: [PATCH 3/4] IILE constexpr if for datatype scale arg for only decimal types. --- cpp/benchmarks/common/generate_input.cu | 28 +++++++++++++++---------- 1 file changed, 17 insertions(+), 11 deletions(-) diff --git a/cpp/benchmarks/common/generate_input.cu b/cpp/benchmarks/common/generate_input.cu index ef65177c941..39665594076 100644 --- a/cpp/benchmarks/common/generate_input.cu +++ b/cpp/benchmarks/common/generate_input.cu @@ -260,19 +260,19 @@ struct random_value_fn()>> { { } - [[nodiscard]] numeric::scale_type get_scale() const - { - return scale.value_or(numeric::scale_type{0}); - } - - rmm::device_uvector operator()(thrust::minstd_rand& engine, unsigned size) + [[nodiscard]] numeric::scale_type get_scale(thrust::minstd_rand& engine) { if (not scale.has_value()) { - int const max_scale = std::numeric_limits::digits10; + constexpr int max_scale = std::numeric_limits::digits10; std::uniform_int_distribution scale_dist{-max_scale, max_scale}; std::mt19937 engine_scale(engine()); scale = numeric::scale_type{scale_dist(engine_scale)}; } + return scale.value_or(numeric::scale_type{0}); + } + + rmm::device_uvector operator()(thrust::minstd_rand& engine, unsigned size) + { auto const ints = dist(engine, size); rmm::device_uvector result(size, cudf::default_stream_value); // Clamp the generated random value to the specified range @@ -396,13 +396,19 @@ std::unique_ptr create_random_column(data_profile const& profile, thrust::minstd_rand& engine, cudf::size_type num_rows) { - using DeviceType = cudf::device_storage_type_t; - numeric::scale_type scale{}; // Bernoulli distribution auto valid_dist = random_value_fn( distribution_params{1. - profile.get_null_probability().value_or(0)}); auto value_dist = random_value_fn{profile.get_distribution_params()}; + using DeviceType = cudf::device_storage_type_t; + cudf::data_type const dtype = [&]() { + if constexpr (cudf::is_fixed_point()) + return cudf::data_type{cudf::type_to_id(), value_dist.get_scale(engine)}; + else + return cudf::data_type{cudf::type_to_id()}; + }(); + // Distribution for picking elements from the array of samples auto const avg_run_len = profile.get_avg_run_length(); rmm::device_uvector data(0, cudf::default_stream_value); @@ -418,7 +424,7 @@ std::unique_ptr create_random_column(data_profile const& profile, }(); rmm::device_uvector samples_null_mask = valid_dist(engine, cardinality); rmm::device_uvector samples = value_dist(engine, cardinality); - if constexpr (cudf::is_fixed_point()) { scale = value_dist.get_scale(); } + // generate n samples and gather. auto const sample_indices = sample_indices_with_run_length(avg_run_len, cardinality, num_rows, engine); @@ -437,7 +443,7 @@ std::unique_ptr create_random_column(data_profile const& profile, cudf::detail::valid_if(null_mask.begin(), null_mask.end(), thrust::identity{}); return std::make_unique( - cudf::data_type{cudf::type_to_id(), scale}, + dtype, num_rows, data.release(), profile.get_null_probability().has_value() ? std::move(result_bitmask) : rmm::device_buffer{}); From 99ff5198a56e85cdc470db26971ec5ef1708d646 Mon Sep 17 00:00:00 2001 From: Karthikeyan Natarajan Date: Thu, 6 Oct 2022 23:32:37 +0530 Subject: [PATCH 4/4] remove unnecessary transform clamp --- cpp/benchmarks/common/generate_input.cu | 12 +----------- 1 file changed, 1 insertion(+), 11 deletions(-) diff --git a/cpp/benchmarks/common/generate_input.cu b/cpp/benchmarks/common/generate_input.cu index 39665594076..2bcdaa6760c 100644 --- a/cpp/benchmarks/common/generate_input.cu +++ b/cpp/benchmarks/common/generate_input.cu @@ -273,17 +273,7 @@ struct random_value_fn()>> { rmm::device_uvector operator()(thrust::minstd_rand& engine, unsigned size) { - auto const ints = dist(engine, size); - rmm::device_uvector result(size, cudf::default_stream_value); - // Clamp the generated random value to the specified range - thrust::transform( - thrust::device, - ints.begin(), - ints.end(), - result.begin(), - [upper_bound = this->upper_bound, lower_bound = this->lower_bound] __device__( - auto int_value) { return std::clamp(int_value, lower_bound, upper_bound); }); - return result; + return dist(engine, size); } };