From 34d1212f5e5a3d60e66f35b3250b435f2c4a0e1d Mon Sep 17 00:00:00 2001 From: Navid Saremi Date: Mon, 23 Mar 2020 12:37:43 +0000 Subject: [PATCH 1/3] Adjusted get_random_data to include runs of same values --- test/rocprim/test_utils.hpp | 47 ++++++++++++++++++++++++++++++++++--- 1 file changed, 44 insertions(+), 3 deletions(-) diff --git a/test/rocprim/test_utils.hpp b/test/rocprim/test_utils.hpp index 853ff1b9f..9e0dc0d60 100644 --- a/test/rocprim/test_utils.hpp +++ b/test/rocprim/test_utils.hpp @@ -43,6 +43,9 @@ std::ostream& operator<<(std::ostream& stream, const rocprim::half& value) namespace test_utils { +static constexpr uint32_t random_data_generation_segments = 8; +static constexpr uint32_t random_data_generation_repeat_strides = 2; + // Support half operators on host side ROCPRIM_HOST inline @@ -209,7 +212,26 @@ inline auto get_random_data(size_t size, T min, T max, int seed_value) gen.seed(seed_value); std::uniform_int_distribution distribution(min, max); std::vector data(size); - std::generate(data.begin(), data.end(), [&]() { return distribution(gen); }); + uint32_t segment_size = size / random_data_generation_segments; + for(uint32_t segment_index = 0; segment_index < random_data_generation_segments; segment_index++) + { + if(segment_index % random_data_generation_repeat_strides == 0) + { + T repeated_value = distribution(gen); + std::fill( + data.begin() + segment_size * segment_index, + data.begin() + segment_size * (segment_index + 1), + repeated_value); + + } + else + { + std::generate( + data.begin() + segment_size * segment_index, + data.begin() + segment_size * (segment_index + 1), + [&]() { return distribution(gen); }); + } + } return data; } @@ -224,7 +246,26 @@ inline auto get_random_data(size_t size, T min, T max, int seed_value) using dis_type = typename std::conditional::value, float, T>::type; std::uniform_real_distribution distribution(min, max); std::vector data(size); - std::generate(data.begin(), data.end(), [&]() { return distribution(gen); }); + uint32_t segment_size = size / random_data_generation_segments; + for(uint32_t segment_index = 0; segment_index < random_data_generation_segments; segment_index++) + { + if(segment_index % random_data_generation_repeat_strides == 0) + { + T repeated_value = distribution(gen); + std::fill( + data.begin() + segment_size * segment_index, + data.begin() + segment_size * (segment_index + 1), + repeated_value); + + } + else + { + std::generate( + data.begin() + segment_size * segment_index, + data.begin() + segment_size * (segment_index + 1), + [&]() { return distribution(gen); }); + } + } return data; } @@ -252,7 +293,7 @@ template inline auto get_random_value(T min, T max, int seed_value) -> typename std::enable_if::value, T>::type { - return get_random_data(1, min, max, seed_value)[0]; + return get_random_data(random_data_generation_segments, min, max, seed_value)[0]; } // Can't use std::prefix_sum for inclusive/exclusive scan, because From 8da7479cca528b879fb3649d283726c9e36253a8 Mon Sep 17 00:00:00 2001 From: Navid Saremi Date: Wed, 25 Mar 2020 14:40:26 +0000 Subject: [PATCH 2/3] Adding different threshold for half types. Adjust custom_test_type to use the new distribution --- test/rocprim/test_block_scan.cpp | 4 +- test/rocprim/test_utils.hpp | 145 ++++++++++++++++++++++++------- 2 files changed, 114 insertions(+), 35 deletions(-) diff --git a/test/rocprim/test_block_scan.cpp b/test/rocprim/test_block_scan.cpp index 1b9369143..295855504 100644 --- a/test/rocprim/test_block_scan.cpp +++ b/test/rocprim/test_block_scan.cpp @@ -244,10 +244,10 @@ struct static_run_algo } // Verifying results - test_utils::assert_near(output, expected, 0.01); + test_utils::assert_near(output, expected, test_utils::precision_threshold::percentage); if(device_output_b) { - test_utils::assert_near(output_b, expected_b, 0.01); + test_utils::assert_near(output_b, expected_b, test_utils::precision_threshold::percentage); } } }; diff --git a/test/rocprim/test_utils.hpp b/test/rocprim/test_utils.hpp index 9e0dc0d60..66ab000a3 100644 --- a/test/rocprim/test_utils.hpp +++ b/test/rocprim/test_utils.hpp @@ -43,8 +43,20 @@ std::ostream& operator<<(std::ostream& stream, const rocprim::half& value) namespace test_utils { -static constexpr uint32_t random_data_generation_segments = 8; -static constexpr uint32_t random_data_generation_repeat_strides = 2; +static constexpr uint32_t random_data_generation_segments = 32; +static constexpr uint32_t random_data_generation_repeat_strides = 4; + +template +struct precision_threshold +{ + static constexpr float percentage = 0.01f; +}; + +template<> +struct precision_threshold +{ + static constexpr float percentage = 0.05f; +}; // Support half operators on host side @@ -213,24 +225,31 @@ inline auto get_random_data(size_t size, T min, T max, int seed_value) std::uniform_int_distribution distribution(min, max); std::vector data(size); uint32_t segment_size = size / random_data_generation_segments; - for(uint32_t segment_index = 0; segment_index < random_data_generation_segments; segment_index++) + if(segment_size != 0) { - if(segment_index % random_data_generation_repeat_strides == 0) + for(uint32_t segment_index = 0; segment_index < random_data_generation_segments; segment_index++) { - T repeated_value = distribution(gen); - std::fill( - data.begin() + segment_size * segment_index, - data.begin() + segment_size * (segment_index + 1), - repeated_value); + if(segment_index % random_data_generation_repeat_strides == 0) + { + T repeated_value = distribution(gen); + std::fill( + data.begin() + segment_size * segment_index, + data.begin() + segment_size * (segment_index + 1), + repeated_value); + } + else + { + std::generate( + data.begin() + segment_size * segment_index, + data.begin() + segment_size * (segment_index + 1), + [&]() { return distribution(gen); }); + } } - else - { - std::generate( - data.begin() + segment_size * segment_index, - data.begin() + segment_size * (segment_index + 1), - [&]() { return distribution(gen); }); - } + } + else + { + std::generate(data.begin(), data.end(), [&]() { return distribution(gen); }); } return data; } @@ -247,24 +266,32 @@ inline auto get_random_data(size_t size, T min, T max, int seed_value) std::uniform_real_distribution distribution(min, max); std::vector data(size); uint32_t segment_size = size / random_data_generation_segments; - for(uint32_t segment_index = 0; segment_index < random_data_generation_segments; segment_index++) + if(segment_size != 0) { - if(segment_index % random_data_generation_repeat_strides == 0) + for(uint32_t segment_index = 0; segment_index < random_data_generation_segments; segment_index++) { - T repeated_value = distribution(gen); - std::fill( - data.begin() + segment_size * segment_index, - data.begin() + segment_size * (segment_index + 1), - repeated_value); + if(segment_index % random_data_generation_repeat_strides == 0) + { + T repeated_value = distribution(gen); + std::fill( + data.begin() + segment_size * segment_index, + data.begin() + segment_size * (segment_index + 1), + repeated_value); + } + else + { + std::generate( + data.begin() + segment_size * segment_index, + data.begin() + segment_size * (segment_index + 1), + [&]() { return distribution(gen); }); + } } - else - { - std::generate( - data.begin() + segment_size * segment_index, - data.begin() + segment_size * (segment_index + 1), - [&]() { return distribution(gen); }); - } + } + else + { + std::generate(data.begin(), data.end(), [&]() { return distribution(gen); }); + } return data; } @@ -756,7 +783,33 @@ inline auto get_random_data(size_t size, typename T::value_type min, typename T: gen.seed(seed_value); std::uniform_int_distribution distribution(min, max); std::vector data(size); - std::generate(data.begin(), data.end(), [&]() { return T(distribution(gen), distribution(gen)); }); + uint32_t segment_size = size / random_data_generation_segments; + if(segment_size != 0) + { + for(uint32_t segment_index = 0; segment_index < random_data_generation_segments; segment_index++) + { + if(segment_index % random_data_generation_repeat_strides == 0) + { + T repeated_value = T(distribution(gen), distribution(gen)); + std::fill( + data.begin() + segment_size * segment_index, + data.begin() + segment_size * (segment_index + 1), + repeated_value); + + } + else + { + std::generate( + data.begin() + segment_size * segment_index, + data.begin() + segment_size * (segment_index + 1), + [&]() { return T(distribution(gen), distribution(gen)); }); + } + } + } + else + { + std::generate(data.begin(), data.end(), [&]() { return T(distribution(gen), distribution(gen)); }); + } return data; } @@ -772,7 +825,33 @@ inline auto get_random_data(size_t size, typename T::value_type min, typename T: gen.seed(seed_value); std::uniform_real_distribution distribution(min, max); std::vector data(size); - std::generate(data.begin(), data.end(), [&]() { return T(distribution(gen), distribution(gen)); }); + uint32_t segment_size = size / random_data_generation_segments; + if(segment_size != 0) + { + for(uint32_t segment_index = 0; segment_index < random_data_generation_segments; segment_index++) + { + if(segment_index % random_data_generation_repeat_strides == 0) + { + T repeated_value = T(distribution(gen), distribution(gen)); + std::fill( + data.begin() + segment_size * segment_index, + data.begin() + segment_size * (segment_index + 1), + repeated_value); + + } + else + { + std::generate( + data.begin() + segment_size * segment_index, + data.begin() + segment_size * (segment_index + 1), + [&]() { return T(distribution(gen), distribution(gen)); }); + } + } + } + else + { + std::generate(data.begin(), data.end(), [&]() { return T(distribution(gen), distribution(gen)); }); + } return data; } @@ -807,7 +886,7 @@ template inline auto get_random_value(typename T::value_type min, typename T::value_type max, int seed_value) -> typename std::enable_if::value || is_custom_test_array_type::value, T>::type { - return get_random_data(1, min, max, seed_value)[0]; + return get_random_data(random_data_generation_segments, min, max, seed_value)[0]; } template From ecc8fd388c969ff311c06727e3e750bdb512a623 Mon Sep 17 00:00:00 2001 From: Adel Johar Date: Mon, 30 Mar 2020 13:04:59 +0200 Subject: [PATCH 3/3] Minor fix to flat_block_id --- rocprim/include/rocprim/intrinsics/thread.hpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/rocprim/include/rocprim/intrinsics/thread.hpp b/rocprim/include/rocprim/intrinsics/thread.hpp index 15fb48023..353414bba 100644 --- a/rocprim/include/rocprim/intrinsics/thread.hpp +++ b/rocprim/include/rocprim/intrinsics/thread.hpp @@ -151,7 +151,7 @@ ROCPRIM_DEVICE inline auto flat_block_id() -> typename std::enable_if<(BlockSizeY > 1 && BlockSizeZ == 1), unsigned int>::type { - return hipThreadIdx_x + (hipBlockIdx_y * hipGridDim_x); + return hipBlockIdx_x + (hipBlockIdx_y * hipGridDim_x); } template @@ -159,7 +159,7 @@ ROCPRIM_DEVICE inline auto flat_block_id() -> typename std::enable_if<(BlockSizeY > 1 && BlockSizeZ > 1), unsigned int>::type { - return hipThreadIdx_x + (hipBlockIdx_y * hipGridDim_x) + + return hipBlockIdx_x + (hipBlockIdx_y * hipGridDim_x) + (hipBlockIdx_z * hipGridDim_y * hipGridDim_x); }