Skip to content

Develop stream 20200330 #144

New issue

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

Sign up for GitHub

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

Already on GitHub? Sign in to your account

Merged
saadrahim merged 5 commits into from
Apr 2, 2020
Merged

Develop stream 20200330 #144

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
34d1212
Adjusted get_random_data to include runs of same values
Mar 23, 2020
8da7479
Adding different threshold for half types. Adjust custom_test_type to…
Mar 25, 2020
ecc8fd3
Minor fix to flat_block_id
Mar 30, 2020
1bf8b68
Merge branch 'EXSWSTRHPC-19' into 'develop_stream'
Mar 30, 2020
a6a09be
Merge branch 'minor_fix' into 'develop_stream'
neon60 Mar 30, 2020
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 2 additions & 2 deletions rocprim/include/rocprim/intrinsics/thread.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -151,15 +151,15 @@ 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<unsigned int BlockSizeX, unsigned int BlockSizeY, unsigned int BlockSizeZ>
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);
}

Expand Down
4 changes: 2 additions & 2 deletions test/rocprim/test_block_scan.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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<T>::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<T>::percentage);
}
}
};
Expand Down
132 changes: 126 additions & 6 deletions test/rocprim/test_utils.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -43,6 +43,21 @@ std::ostream& operator<<(std::ostream& stream, const rocprim::half& value)
namespace test_utils
{

static constexpr uint32_t random_data_generation_segments = 32;
static constexpr uint32_t random_data_generation_repeat_strides = 4;

template<class T>
struct precision_threshold
{
static constexpr float percentage = 0.01f;
};

template<>
struct precision_threshold<rocprim::half>
{
static constexpr float percentage = 0.05f;
};

// Support half operators on host side

ROCPRIM_HOST inline
Expand Down Expand Up @@ -209,7 +224,33 @@ inline auto get_random_data(size_t size, T min, T max, int seed_value)
gen.seed(seed_value);
std::uniform_int_distribution<T> distribution(min, max);
std::vector<T> data(size);
std::generate(data.begin(), data.end(), [&]() { return 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 = 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(), data.end(), [&]() { return distribution(gen); });
}
return data;
}

Expand All @@ -224,7 +265,34 @@ inline auto get_random_data(size_t size, T min, T max, int seed_value)
using dis_type = typename std::conditional<std::is_same<rocprim::half, T>::value, float, T>::type;
std::uniform_real_distribution<dis_type> distribution(min, max);
std::vector<T> data(size);
std::generate(data.begin(), data.end(), [&]() { return 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 = 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(), data.end(), [&]() { return distribution(gen); });

}
return data;
}

Expand Down Expand Up @@ -252,7 +320,7 @@ template<class T>
inline auto get_random_value(T min, T max, int seed_value)
-> typename std::enable_if<rocprim::is_arithmetic<T>::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
Expand Down Expand Up @@ -715,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<typename T::value_type> distribution(min, max);
std::vector<T> 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;
}

Expand All @@ -731,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<typename T::value_type> distribution(min, max);
std::vector<T> 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;
}

Expand Down Expand Up @@ -766,7 +886,7 @@ template<class T>
inline auto get_random_value(typename T::value_type min, typename T::value_type max, int seed_value)
-> typename std::enable_if<is_custom_test_type<T>::value || is_custom_test_array_type<T>::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<class T>
Expand Down