Skip to content

Add settable seed #96

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
neon60 merged 2 commits into from
Feb 28, 2020
Merged

Add settable seed #96

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
bf2279f
Add settable seed
Feb 25, 2020
c60688f
Update readme
neon60 Feb 26, 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
1 change: 1 addition & 0 deletions test/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -33,6 +33,7 @@ endfunction()
# Tests
# ****************************************************************************


add_rocthrust_test("thrust.hip.adjacent_difference" test_adjacent_difference.cpp)
add_rocthrust_test("thrust.hip.advance" test_advance.cpp)
add_rocthrust_test("thrust.hip.allocator" test_allocator.cpp)
Expand Down
182 changes: 104 additions & 78 deletions test/test_adjacent_difference.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -70,42 +70,52 @@ TYPED_TEST(AdjacentDifferenceVariableTests, TestAdjacentDifference)
const std::vector<size_t> sizes = get_sizes();
for(auto size : sizes)
{
thrust::host_vector<T> h_input = get_random_data<T>(
size, std::numeric_limits<T>::min(), std::numeric_limits<T>::max());
thrust::device_vector<T> d_input = h_input;

thrust::host_vector<T> h_output(size);
thrust::device_vector<T> d_output(size);

typename thrust::host_vector<T>::iterator h_result;
typename thrust::device_vector<T>::iterator d_result;

h_result = thrust::adjacent_difference(h_input.begin(), h_input.end(), h_output.begin());
d_result = thrust::adjacent_difference(d_input.begin(), d_input.end(), d_output.begin());

ASSERT_EQ(h_result - h_output.begin(), size);
ASSERT_EQ(d_result - d_output.begin(), size);
ASSERT_EQ_QUIET(h_output, d_output);

h_result = thrust::adjacent_difference(
h_input.begin(), h_input.end(), h_output.begin(), thrust::plus<T>());
d_result = thrust::adjacent_difference(
d_input.begin(), d_input.end(), d_output.begin(), thrust::plus<T>());

ASSERT_EQ(h_result - h_output.begin(), size);
ASSERT_EQ(d_result - d_output.begin(), size);
ASSERT_EQ_QUIET(h_output, d_output);

// in-place operation
h_result = thrust::adjacent_difference(
h_input.begin(), h_input.end(), h_input.begin(), thrust::plus<T>());
d_result = thrust::adjacent_difference(
d_input.begin(), d_input.end(), d_input.begin(), thrust::plus<T>());

ASSERT_EQ(h_result - h_input.begin(), size);
ASSERT_EQ(d_result - d_input.begin(), size);
ASSERT_EQ_QUIET(h_input, h_output); //computed previously
ASSERT_EQ_QUIET(d_input, d_output); //computed previously
SCOPED_TRACE(testing::Message() << "with size= " << size);
for(size_t seed_index = 0; seed_index < random_seeds_count + seed_size; seed_index++)
{
unsigned int seed_value
= seed_index < random_seeds_count ? rand() : seeds[seed_index - random_seeds_count];
SCOPED_TRACE(testing::Message() << "with seed= " << seed_value);

thrust::host_vector<T> h_input = get_random_data<T>(
size, std::numeric_limits<T>::min(), std::numeric_limits<T>::max(), seed_value);
thrust::device_vector<T> d_input = h_input;

thrust::host_vector<T> h_output(size);
thrust::device_vector<T> d_output(size);

typename thrust::host_vector<T>::iterator h_result;
typename thrust::device_vector<T>::iterator d_result;

h_result
= thrust::adjacent_difference(h_input.begin(), h_input.end(), h_output.begin());
d_result
= thrust::adjacent_difference(d_input.begin(), d_input.end(), d_output.begin());

ASSERT_EQ(h_result - h_output.begin(), size);
ASSERT_EQ(d_result - d_output.begin(), size);
ASSERT_EQ_QUIET(h_output, d_output);

h_result = thrust::adjacent_difference(
h_input.begin(), h_input.end(), h_output.begin(), thrust::plus<T>());
d_result = thrust::adjacent_difference(
d_input.begin(), d_input.end(), d_output.begin(), thrust::plus<T>());

ASSERT_EQ(h_result - h_output.begin(), size);
ASSERT_EQ(d_result - d_output.begin(), size);
ASSERT_EQ_QUIET(h_output, d_output);

// in-place operation
h_result = thrust::adjacent_difference(
h_input.begin(), h_input.end(), h_input.begin(), thrust::plus<T>());
d_result = thrust::adjacent_difference(
d_input.begin(), d_input.end(), d_input.begin(), thrust::plus<T>());

ASSERT_EQ(h_result - h_input.begin(), size);
ASSERT_EQ(d_result - d_input.begin(), size);
ASSERT_EQ_QUIET(h_input, h_output); //computed previously
ASSERT_EQ_QUIET(d_input, d_output); //computed previously
}
}
}

Expand All @@ -116,31 +126,39 @@ TYPED_TEST(AdjacentDifferenceVariableTests, TestAdjacentDifferenceInPlaceWithRel
const std::vector<size_t> sizes = get_sizes();
for(auto size : sizes)
{
thrust::host_vector<T> h_input = get_random_data<T>(
size, std::numeric_limits<T>::min(), std::numeric_limits<T>::max());
thrust::device_vector<T> d_input = h_input;

thrust::host_vector<T> h_output(size);
thrust::device_vector<T> d_output(size);

typename thrust::host_vector<T>::iterator h_result;
typename thrust::device_vector<T>::iterator d_result;

h_result = thrust::adjacent_difference(
h_input.begin(), h_input.end(), h_output.begin(), thrust::plus<T>());
d_result = thrust::adjacent_difference(
d_input.begin(), d_input.end(), d_output.begin(), thrust::plus<T>());

// in-place operation with different iterator types
h_result = thrust::adjacent_difference(
h_input.cbegin(), h_input.cend(), h_input.begin(), thrust::plus<T>());
d_result = thrust::adjacent_difference(
d_input.cbegin(), d_input.cend(), d_input.begin(), thrust::plus<T>());

ASSERT_EQ(h_result - h_input.begin(), size);
ASSERT_EQ(d_result - d_input.begin(), size);
ASSERT_EQ_QUIET(h_output, h_input); // reference computed previously
ASSERT_EQ_QUIET(d_output, d_input); // reference computed previously
SCOPED_TRACE(testing::Message() << "with size= " << size);
for(size_t seed_index = 0; seed_index < random_seeds_count + seed_size; seed_index++)
{
unsigned int seed_value
= seed_index < random_seeds_count ? rand() : seeds[seed_index - random_seeds_count];
SCOPED_TRACE(testing::Message() << "with seed= " << seed_value);

thrust::host_vector<T> h_input = get_random_data<T>(
size, std::numeric_limits<T>::min(), std::numeric_limits<T>::max(), seed_value);
thrust::device_vector<T> d_input = h_input;

thrust::host_vector<T> h_output(size);
thrust::device_vector<T> d_output(size);

typename thrust::host_vector<T>::iterator h_result;
typename thrust::device_vector<T>::iterator d_result;

h_result = thrust::adjacent_difference(
h_input.begin(), h_input.end(), h_output.begin(), thrust::plus<T>());
d_result = thrust::adjacent_difference(
d_input.begin(), d_input.end(), d_output.begin(), thrust::plus<T>());

// in-place operation with different iterator types
h_result = thrust::adjacent_difference(
h_input.cbegin(), h_input.cend(), h_input.begin(), thrust::plus<T>());
d_result = thrust::adjacent_difference(
d_input.cbegin(), d_input.cend(), d_input.begin(), thrust::plus<T>());

ASSERT_EQ(h_result - h_input.begin(), size);
ASSERT_EQ(d_result - d_input.begin(), size);
ASSERT_EQ_QUIET(h_output, h_input); // reference computed previously
ASSERT_EQ_QUIET(d_output, d_input); // reference computed previously
}
}
}

Expand All @@ -151,28 +169,36 @@ TYPED_TEST(AdjacentDifferenceVariableTests, TestAdjacentDifferenceDiscardIterato
const std::vector<size_t> sizes = get_sizes();
for(auto size : sizes)
{
thrust::host_vector<T> h_input = get_random_data<T>(
size, std::numeric_limits<T>::min(), std::numeric_limits<T>::max());
thrust::device_vector<T> d_input = h_input;

thrust::discard_iterator<> h_result;
thrust::discard_iterator<> d_result;

h_result = thrust::adjacent_difference(
h_input.begin(), h_input.end(), thrust::make_discard_iterator());
d_result = thrust::adjacent_difference(
d_input.begin(), d_input.end(), thrust::make_discard_iterator());

thrust::discard_iterator<> reference(size);

ASSERT_EQ_QUIET(reference, h_result);
ASSERT_EQ_QUIET(reference, d_result);
SCOPED_TRACE(testing::Message() << "with size= " << size);
for(size_t seed_index = 0; seed_index < random_seeds_count + seed_size; seed_index++)
{
unsigned int seed_value
= seed_index < random_seeds_count ? rand() : seeds[seed_index - random_seeds_count];
SCOPED_TRACE(testing::Message() << "with seed= " << seed_value);

thrust::host_vector<T> h_input = get_random_data<T>(
size, std::numeric_limits<T>::min(), std::numeric_limits<T>::max(), seed_value);
thrust::device_vector<T> d_input = h_input;

thrust::discard_iterator<> h_result;
thrust::discard_iterator<> d_result;

h_result = thrust::adjacent_difference(
h_input.begin(), h_input.end(), thrust::make_discard_iterator());
d_result = thrust::adjacent_difference(
d_input.begin(), d_input.end(), thrust::make_discard_iterator());

thrust::discard_iterator<> reference(size);

ASSERT_EQ_QUIET(reference, h_result);
ASSERT_EQ_QUIET(reference, d_result);
}
}
}

template <typename InputIterator, typename OutputIterator>
OutputIterator
adjacent_difference(my_system& system, InputIterator, InputIterator, OutputIterator result)
adjacent_difference(my_system& system, InputIterator, InputIterator, OutputIterator result)
{
system.validate_dispatch();
return result;
Expand Down
95 changes: 41 additions & 54 deletions test/test_binary_search.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -27,19 +27,15 @@ TESTS_DEFINE(BinarySearchTestsInKernel, NumericalTestsParams);

struct custom_less
{
template<class T>
__device__ inline
bool operator()(T a, T b)
template <class T>
__device__ inline bool operator()(T a, T b)
{
return a < b;
}
};

template<class T>
__global__
void lower_bound_kernel(size_t n,
T* input,
ptrdiff_t* output)
template <class T>
__global__ void lower_bound_kernel(size_t n, T* input, ptrdiff_t* output)
{
output[0] = thrust::lower_bound(thrust::device, input, input + n, T(0), custom_less()) - input;
output[1] = thrust::lower_bound(thrust::device, input, input + n, T(1)) - input;
Expand All @@ -66,13 +62,14 @@ TYPED_TEST(BinarySearchTestsInKernel, TestLowerBound)

thrust::device_vector<ptrdiff_t> d_output(10);

hipLaunchKernelGGL(
HIP_KERNEL_NAME(lower_bound_kernel),
dim3(1), dim3(1), 0, 0,
size_t(d_input.size()),
thrust::raw_pointer_cast(d_input.data()),
thrust::raw_pointer_cast(d_output.data())
);
hipLaunchKernelGGL(HIP_KERNEL_NAME(lower_bound_kernel),
dim3(1),
dim3(1),
0,
0,
size_t(d_input.size()),
thrust::raw_pointer_cast(d_input.data()),
thrust::raw_pointer_cast(d_output.data()));

thrust::host_vector<ptrdiff_t> output = d_output;
ASSERT_EQ(output[0], 0);
Expand All @@ -87,11 +84,8 @@ TYPED_TEST(BinarySearchTestsInKernel, TestLowerBound)
ASSERT_EQ(output[9], 5);
}

template<class T>
__global__
void upper_bound_kernel(size_t n,
T* input,
ptrdiff_t* output)
template <class T>
__global__ void upper_bound_kernel(size_t n, T* input, ptrdiff_t* output)
{
output[0] = thrust::upper_bound(thrust::device, input, input + n, T(0)) - input;
output[1] = thrust::upper_bound(thrust::device, input, input + n, T(1)) - input;
Expand All @@ -118,13 +112,14 @@ TYPED_TEST(BinarySearchTestsInKernel, TestUpperBound)

thrust::device_vector<ptrdiff_t> d_output(10);

hipLaunchKernelGGL(
HIP_KERNEL_NAME(upper_bound_kernel),
dim3(1), dim3(1), 0, 0,
size_t(d_input.size()),
thrust::raw_pointer_cast(d_input.data()),
thrust::raw_pointer_cast(d_output.data())
);
hipLaunchKernelGGL(HIP_KERNEL_NAME(upper_bound_kernel),
dim3(1),
dim3(1),
0,
0,
size_t(d_input.size()),
thrust::raw_pointer_cast(d_input.data()),
thrust::raw_pointer_cast(d_output.data()));

thrust::host_vector<ptrdiff_t> output = d_output;
ASSERT_EQ(output[0], 1);
Expand All @@ -139,11 +134,8 @@ TYPED_TEST(BinarySearchTestsInKernel, TestUpperBound)
ASSERT_EQ(output[9], 5);
}

template<class T>
__global__
void binary_search_kernel(size_t n,
T* input,
bool* output)
template <class T>
__global__ void binary_search_kernel(size_t n, T* input, bool* output)
{
output[0] = thrust::binary_search(thrust::device, input, input + n, T(0));
output[1] = thrust::binary_search(thrust::device, input, input + n, T(1));
Expand All @@ -170,13 +162,14 @@ TYPED_TEST(BinarySearchTestsInKernel, TestBinarySearch)

thrust::device_vector<bool> d_output(10);

hipLaunchKernelGGL(
HIP_KERNEL_NAME(binary_search_kernel),
dim3(1), dim3(1), 0, 0,
size_t(d_input.size()),
thrust::raw_pointer_cast(d_input.data()),
thrust::raw_pointer_cast(d_output.data())
);
hipLaunchKernelGGL(HIP_KERNEL_NAME(binary_search_kernel),
dim3(1),
dim3(1),
0,
0,
size_t(d_input.size()),
thrust::raw_pointer_cast(d_input.data()),
thrust::raw_pointer_cast(d_output.data()));

thrust::host_vector<bool> output = d_output;
ASSERT_EQ(output[0], true);
Expand Down Expand Up @@ -220,10 +213,8 @@ TYPED_TEST(BinarySearchTests, TestScalarLowerBoundSimple)
}

template <typename ForwardIterator, typename LessThanComparable>
ForwardIterator lower_bound(my_system& system,
ForwardIterator first,
ForwardIterator,
const LessThanComparable&)
ForwardIterator
lower_bound(my_system& system, ForwardIterator first, ForwardIterator, const LessThanComparable&)
{
system.validate_dispatch();
return first;
Expand All @@ -241,7 +232,7 @@ TEST(BinarySearchTests, TestScalarLowerBoundDispatchExplicit)

template <typename ForwardIterator, typename LessThanComparable>
ForwardIterator
lower_bound(my_tag, ForwardIterator first, ForwardIterator, const LessThanComparable&)
lower_bound(my_tag, ForwardIterator first, ForwardIterator, const LessThanComparable&)
{
*first = 13;
return first;
Expand Down Expand Up @@ -281,10 +272,8 @@ TYPED_TEST(BinarySearchTests, TestScalarUpperBoundSimple)
}

template <typename ForwardIterator, typename LessThanComparable>
ForwardIterator upper_bound(my_system& system,
ForwardIterator first,
ForwardIterator,
const LessThanComparable&)
ForwardIterator
upper_bound(my_system& system, ForwardIterator first, ForwardIterator, const LessThanComparable&)
{
system.validate_dispatch();
return first;
Expand All @@ -302,7 +291,7 @@ TEST(BinarySearchTests, TestScalarUpperBoundDispatchExplicit)

template <typename ForwardIterator, typename LessThanComparable>
ForwardIterator
upper_bound(my_tag, ForwardIterator first, ForwardIterator, const LessThanComparable&)
upper_bound(my_tag, ForwardIterator first, ForwardIterator, const LessThanComparable&)
{
*first = 13;
return first;
Expand Down Expand Up @@ -410,10 +399,8 @@ TYPED_TEST(BinarySearchTests, TestScalarEqualRangeSimple)
}

template <typename ForwardIterator, typename LessThanComparable>
thrust::pair<ForwardIterator, ForwardIterator> equal_range(my_system& system,
ForwardIterator first,
ForwardIterator,
const LessThanComparable&)
thrust::pair<ForwardIterator, ForwardIterator>
equal_range(my_system& system, ForwardIterator first, ForwardIterator, const LessThanComparable&)
{
system.validate_dispatch();
return thrust::make_pair(first, first);
Expand All @@ -431,7 +418,7 @@ TEST(BinarySearchTests, TestScalarEqualRangeDispatchExplicit)

template <typename ForwardIterator, typename LessThanComparable>
thrust::pair<ForwardIterator, ForwardIterator>
equal_range(my_tag, ForwardIterator first, ForwardIterator, const LessThanComparable&)
equal_range(my_tag, ForwardIterator first, ForwardIterator, const LessThanComparable&)
{
*first = 13;
return thrust::make_pair(first, first);
Expand Down
Loading