Added remaining test cases
MODIFY
test/test_binary_search.cpp
test/test_binary_search.cpp
@@ -470,4 +470,53 @@ TEST(BinarySearchTests, TestScalarEqualRangeDispatchImplicit)
ASSERT_EQ(13, vec.front());
}
+
+__global__
+THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
+void BinarySearchKernel(int const N, int* in_array, int*result_array, int search_value)
+{
+ if(threadIdx.x == 0)
+ {
+ thrust::device_ptr<int> begin(in_array);
+ thrust::device_ptr<int> end(in_array + N);
+ result_array[search_value]=thrust::binary_search(thrust::seq, begin,end,search_value);
+ }
+}
+
+TEST(BinarySearchTests, TestBinarySearchDevice)
+{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
+ for(auto size: {0, 1, 2, 4, 6, 12, 16, 24, 32, 64, 84, 128, 160, 256} )
+ {
+ SCOPED_TRACE(testing::Message() << "with size= " << size);
+
+ for(auto seed : get_seeds())
+ {
+ SCOPED_TRACE(testing::Message() << "with seed= " << seed);
+
+ thrust::host_vector<int> h_data = get_random_data<int>(size, 0, size, seed);
+ thrust::device_vector<int> d_data = h_data;
+
+ thrust::host_vector<int> h_result(size*2,-1);
+ thrust::device_vector<int> d_result(size*2,-1);
+
+ for(int search_value = 0; search_value < size*2; search_value++)
+ {
+ SCOPED_TRACE(testing::Message() << "searching for " <<search_value);
+
+ h_result[search_value] = thrust::binary_search(h_data.begin(),h_data.end(),search_value);
+ hipLaunchKernelGGL(BinarySearchKernel,
+ dim3(1, 1, 1),
+ dim3(128, 1, 1),
+ 0,
+ 0,
+ size,
+ thrust::raw_pointer_cast(&d_data[0]),
+ thrust::raw_pointer_cast(&d_result[0]),
+ search_value);
+ }
+ ASSERT_EQ(h_result,d_result);
+ }
+ }
+}
THRUST_DISABLE_MSVC_POSSIBLE_LOSS_OF_DATA_WARNING_END

MODIFY
test/test_merge.cpp
test/test_merge.cpp
@@ -291,7 +291,6 @@ void MergeKernel(int const N, T* inA_array, T* inB_array, T *out_array)
thrust::device_ptr<int> out_begin(out_array);
thrust::merge(thrust::seq, inA_begin, inA_end, inB_begin, inB_end,out_begin);
- // thrust::merge_by_key(thrust::seq, inA_begin, inA_end, inB_begin, inB_end,inA_begin,inB_begin,out_begin,out_begin);
}
}

MODIFY
test/test_merge_by_key.cpp
test/test_merge_by_key.cpp
@@ -440,3 +440,107 @@ TYPED_TEST(PrimitiveMergeByKeyTests, MergeByKeyDescending)
}
}
}
+
+
+
+template<class T>
+__global__
+THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
+void MergeByKeyKernel(int const N, T* keys_A, T* keys_B, T * values_A, T* values_B, T* keys_result, T* values_result)
+{
+ if(threadIdx.x == 0)
+ {
+ thrust::device_ptr<int> keyA_begin(keys_A);
+ thrust::device_ptr<int> keyA_end(keys_A + N);
+ thrust::device_ptr<int> keyB_begin(keys_B);
+ thrust::device_ptr<int> keyB_end(keys_B + N);
+ thrust::device_ptr<int> valuesA_begin(values_A);
+ thrust::device_ptr<int> valuesB_begin(values_B);
+ thrust::device_ptr<int> keys_result_begin(keys_result);
+ thrust::device_ptr<int> values_result_begin(values_result);
+
+
+ thrust::merge_by_key(thrust::seq, keyA_begin, keyA_end,
+ keyB_begin, keyB_end,
+ valuesA_begin, valuesB_begin,
+ keys_result_begin,
+ values_result_begin);
+
+ }
+}
+TEST(MergeByKeyTests, TestMergeByKeyDevice)
+{
+ using T = int;
+
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
+
+ for(auto size : get_sizes())
+ {
+ SCOPED_TRACE(testing::Message() << "with size= " << size);
+
+ for(auto seed : get_seeds())
+ {
+ SCOPED_TRACE(testing::Message() << "with seed= " << seed);
+
+ thrust::host_vector<T> h_keys_a = get_random_data<T>(
+ size, std::numeric_limits<T>::min(), std::numeric_limits<T>::max(), seed);
+ thrust::host_vector<T> h_keys_b = get_random_data<T>(
+ size,
+ std::numeric_limits<T>::min(),
+ std::numeric_limits<T>::max(),
+ seed + seed_value_addition
+ );
+
+ thrust::host_vector<T> h_values_a = get_random_data<T>(
+ size, std::numeric_limits<T>::min(), std::numeric_limits<T>::max(), seed);
+ thrust::host_vector<T> h_values_b = get_random_data<T>(
+ size,
+ std::numeric_limits<T>::min(),
+ std::numeric_limits<T>::max(),
+ seed + seed_value_addition
+ );
+
+
+ thrust::stable_sort(h_keys_a.begin(), h_keys_a.end(), thrust::greater<T>());
+ thrust::stable_sort(h_keys_b.begin(), h_keys_b.end(), thrust::greater<T>());
+
+ thrust::device_vector<T> d_keys_a = h_keys_a;
+ thrust::device_vector<T> d_keys_b = h_keys_b;
+
+ thrust::device_vector<T> d_values_a = h_values_a;
+ thrust::device_vector<T> d_values_b = h_values_b;
+
+ thrust::host_vector<T> h_keys_result(h_keys_a.size() + h_keys_b.size());
+ thrust::device_vector<T> d_keys_result(d_keys_a.size() + d_keys_b.size());
+
+ thrust::host_vector<T> h_values_result(h_values_a.size() + h_values_b.size());
+ thrust::device_vector<T> d_values_result(d_values_a.size() + d_values_b.size());
+
+ typename thrust::host_vector<T>::iterator h_end;
+ typename thrust::device_vector<T>::iterator d_end;
+
+ thrust::merge_by_key(h_keys_a.begin(),h_keys_a.end(),
+ h_keys_b.begin(),h_keys_b.end(),
+ h_values_a.begin(),h_values_b.begin(),
+ h_keys_result.begin(),
+ h_values_result.begin());
+
+ hipLaunchKernelGGL(MergeByKeyKernel,
+ dim3(1, 1, 1),
+ dim3(128, 1, 1),
+ 0,
+ 0,
+ size,
+ thrust::raw_pointer_cast(&d_keys_a[0]),
+ thrust::raw_pointer_cast(&d_keys_b[0]),
+ thrust::raw_pointer_cast(&d_values_a[0]),
+ thrust::raw_pointer_cast(&d_values_b[0]),
+ thrust::raw_pointer_cast(&d_keys_result[0]),
+ thrust::raw_pointer_cast(&d_values_result[0]));
+
+ ASSERT_EQ(h_keys_result, d_keys_result);
+ ASSERT_EQ(h_values_result, d_values_result);
+
+ }
+ }
+}

MODIFY
test/test_reduce.cpp
test/test_reduce.cpp
@@ -255,3 +255,53 @@ TYPED_TEST(ReducePrimitiveTests, TestReduceCountingIterator)
}
}
}
+
+
+__global__
+THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
+void ReduceKernel(int const N, int *in_array, int *result)
+{
+ if(threadIdx.x == 0)
+ {
+ thrust::device_ptr<int> in_begin(in_array);
+ thrust::device_ptr<int> in_end(in_array + N);
+
+ result[0]= thrust::reduce(thrust::seq, in_begin, in_end);
+ }
+}
+
+
+TEST(ReduceTests, TestReduceDevice)
+{
+
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
+
+ for(auto size : get_sizes())
+ {
+ SCOPED_TRACE(testing::Message() << "with size= " << size);
+
+ for(auto seed : get_seeds())
+ {
+ SCOPED_TRACE(testing::Message() << "with seed= " << seed);
+
+ thrust::host_vector<int> h_data = get_random_data<int>(
+ size, 0, 15, seed);
+ thrust::device_vector<int> d_data = h_data;
+ thrust::device_vector<int> d_output(1,0);
+
+
+ int h_output = thrust::reduce(h_data.begin(), h_data.end());
+
+ hipLaunchKernelGGL(ReduceKernel,
+ dim3(1, 1, 1),
+ dim3(128, 1, 1),
+ 0,
+ 0,
+ size,
+ thrust::raw_pointer_cast(&d_data[0]),
+ thrust::raw_pointer_cast(&d_output[0]));
+
+ ASSERT_EQ(h_output, d_output[0]);
+ }
+ }
+}

MODIFY
test/test_reduce_by_key.cpp
test/test_reduce_by_key.cpp
@@ -341,3 +341,88 @@ TEST(ReduceByKeysTests, TestReduceByKeyDispatchImplicit)
ASSERT_EQ(13, vec.front());
}
+
+
+__global__
+THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
+void ReduceByKeyKernel(int const N, int *in_keys, int * in_values, int * out_keys, int * out_values, int * out_sizes)
+{
+ if(threadIdx.x == 0)
+ {
+ thrust::device_ptr<int> in_keys_begin(in_keys);
+ thrust::device_ptr<int> in_keys_end(in_keys + N);
+ thrust::device_ptr<int> in_values_begin(in_values);
+ thrust::device_ptr<int> out_keys_begin(out_keys);
+ thrust::device_ptr<int> out_values_begin(out_values);
+
+ auto end_pair = thrust::reduce_by_key(thrust::seq, in_keys_begin,in_keys_end,
+ in_values_begin,
+ out_keys_begin,
+ out_values_begin);
+ out_sizes[0] = end_pair.first - out_keys_begin;
+ out_sizes[1] = end_pair.second - out_values_begin;
+ }
+}
+
+
+TEST(ReduceByKeyTests, TestReduceByKeyDevice)
+{
+
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
+
+ for(auto size : get_sizes())
+ {
+ SCOPED_TRACE(testing::Message() << "with size= " << size);
+
+ for(auto seed : get_seeds())
+ {
+ SCOPED_TRACE(testing::Message() << "with seed= " << seed);
+
+ thrust::host_vector<int> h_values = get_random_data<int>(
+ size,
+ std::numeric_limits<int>::min(),
+ std::numeric_limits<int>::max(),
+ seed);
+ thrust::host_vector<int> h_keys = get_random_data<int>(
+ size, 0, 13, seed);
+ thrust::device_vector<int> d_values = h_values;
+ thrust::device_vector<int> d_keys = h_keys;
+
+ thrust::host_vector<int> h_keys_result(size);
+ thrust::host_vector<int> h_values_result(size);
+
+ thrust::device_vector<int> d_keys_result(size);
+ thrust::device_vector<int> d_values_result(size);
+ thrust::device_vector<int> d_output_sizes(2);
+
+
+
+
+ auto end_pair = thrust::reduce_by_key(h_keys.begin(), h_keys.end(),
+ h_values.begin(),
+ h_keys_result.begin(),
+ h_values_result.begin());
+
+ hipLaunchKernelGGL(ReduceByKeyKernel,
+ dim3(1, 1, 1),
+ dim3(128, 1, 1),
+ 0,
+ 0,
+ size,
+ thrust::raw_pointer_cast(&d_keys[0]),
+ thrust::raw_pointer_cast(&d_values[0]),
+ thrust::raw_pointer_cast(&d_keys_result[0]),
+ thrust::raw_pointer_cast(&d_values_result[0]),
+ thrust::raw_pointer_cast(&d_output_sizes[0]));
+
+ h_keys_result.resize(end_pair.first - h_keys_result.begin());
+ h_values_result.resize(end_pair.second - h_values_result.begin());
+ d_keys_result.resize(d_output_sizes[0]);
+ d_values_result.resize(d_output_sizes[1]);
+
+ ASSERT_EQ(h_keys_result,d_keys_result);
+ ASSERT_EQ(h_values_result,d_values_result);
+
+ }
+ }
+}

MODIFY
test/test_remove.cpp
test/test_remove.cpp
@@ -887,7 +887,7 @@ void RemoveKernel(int const N, int* array, int remove_value)
{
thrust::device_ptr<int> begin(array);
thrust::device_ptr<int> end(array + N);
- int size = thrust::remove(thrust::seq, begin, end,remove_value) - begin;
+ thrust::remove(thrust::seq, begin, end,remove_value) - begin;
}
}

MODIFY
test/test_unique.cpp
test/test_unique.cpp
@@ -319,7 +319,7 @@ void UniqueKernel(int const N, int *in_array, int *out_size)
}
-TEST(UniqueIntegralTests, TestUniqueDevice)
+TEST(UniqueTests, TestUniqueDevice)
{
SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());

MODIFY
test/test_unique_by_key.cpp
test/test_unique_by_key.cpp
@@ -514,3 +514,80 @@ TYPED_TEST(UniqueByKeyIntegralTests, TestUniqueCopyByKeyToDiscardIterator)
}
}
}
+
+__global__
+THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
+void UniqueByKeyKernel(int const N, int *in_array, int *in_keys, int *out_size)
+{
+ if(threadIdx.x == 0)
+ {
+ thrust::device_ptr<int> in_begin(in_array);
+ thrust::device_ptr<int> in_end(in_array + N);
+
+ thrust::device_ptr<int> keys_begin(in_keys);
+ thrust::device_ptr<int> keys_end(in_keys + N);
+
+ thrust::pair<thrust::device_vector<int>::iterator,thrust::device_vector<int>::iterator> d_result = thrust::unique_by_key(thrust::seq,keys_begin, keys_end, in_begin);
+ out_size[0] = d_result.second - thrust::device_vector<int>::iterator(in_begin);
+ out_size[1] = d_result.first - thrust::device_vector<int>::iterator(keys_begin);
+
+ }
+}
+
+
+TEST(UniqueIntegralTests, TestUniqueDevice)
+{
+
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
+
+ for(auto size : get_sizes())
+ {
+ SCOPED_TRACE(testing::Message() << "with size= " << size);
+
+ for(auto seed : get_seeds())
+ {
+ SCOPED_TRACE(testing::Message() << "with seed= " << seed);
+
+ thrust::host_vector<int> h_data = get_random_data<int>(
+ size, 0, 15, seed);
+ thrust::host_vector<int> h_keys = get_random_data<int>(
+ size, 0, 15, seed);
+ thrust::device_vector<int> d_data = h_data;
+ thrust::device_vector<int> d_keys = h_keys;
+
+ thrust::device_vector<int> d_output_size(2,0);
+
+ thrust::pair<thrust::host_vector<int>::iterator,thrust::host_vector<int>::iterator> h_new_last;
+ // thrust::pair<thrust::device_vector<int>::iterator,thrust::device_vector<int>::iterator> d_new_last;
+
+ h_new_last = thrust::unique_by_key(h_keys.begin(),h_keys.end(),h_data.begin());
+ auto h_values_last = h_new_last.second;
+ auto h_values_keys = h_new_last.first;
+
+ hipLaunchKernelGGL(UniqueByKeyKernel,
+ dim3(1, 1, 1),
+ dim3(128, 1, 1),
+ 0,
+ 0,
+ size,
+ thrust::raw_pointer_cast(&d_data[0]),
+ thrust::raw_pointer_cast(&d_keys[0]),
+ thrust::raw_pointer_cast(&d_output_size[0]));
+
+
+ ASSERT_EQ(h_values_last - h_data.begin(),d_output_size[0]);
+ ASSERT_EQ(h_values_keys - h_keys.begin(),d_output_size[1]);
+
+ h_data.resize(h_values_last - h_data.begin());
+ h_keys.resize(h_values_last - h_data.begin());
+
+ d_data.resize(d_output_size[0]);
+ d_keys.resize(d_output_size[1]);
+
+
+ ASSERT_EQ(h_data, d_data);
+ ASSERT_EQ(h_keys, d_keys);
+
+ }
+ }
+}

Added remaining test cases