Implemented additional test cases for device side algorithms
MODIFY
test/test_find.cpp
test/test_find.cpp
@@ -396,6 +396,7 @@ void FindKernel(int const N, int* in_array, int value, int *out_array)
thrust::device_ptr<int> in_end(in_array + N);
auto x = thrust::find(thrust::hip::par, in_begin, in_end,value);
+ out_array[0] = x - in_begin;
}
}
TEST(FindTests, TestFindDevice)

MODIFY
test/test_remove.cpp
test/test_remove.cpp
@@ -877,3 +877,49 @@ TYPED_TEST(RemoveVariableTests, TestRemoveCopyIfStencilToDiscardIterator)
}
}
}
+
+
+__global__
+THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
+void RemoveKernel(int const N, int* array, int remove_value)
+{
+ if(threadIdx.x == 0)
+ {
+ thrust::device_ptr<int> begin(array);
+ thrust::device_ptr<int> end(array + N);
+ int size = thrust::remove(thrust::hip::par, begin, end,remove_value) - begin;
+ }
+}
+
+TEST(RemoveTests, TestRemoveDevice)
+{
+ 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;
+
+ int remove_value = get_random_data<int>(1,0,size,seed)[0];
+ SCOPED_TRACE(testing::Message() << "with remove_value= " <<remove_value);
+
+ thrust::remove(h_data.begin(), h_data.end(),remove_value);
+ hipLaunchKernelGGL(RemoveKernel,
+ dim3(1, 1, 1),
+ dim3(128, 1, 1),
+ 0,
+ 0,
+ size,
+ thrust::raw_pointer_cast(&d_data[0]),
+ remove_value);
+
+ ASSERT_EQ(h_data, d_data);
+ }
+ }
+}

MODIFY
test/test_replace.cpp
test/test_replace.cpp
@@ -830,3 +830,52 @@ TYPED_TEST(PrimitiveReplaceTests, ReplaceCopyIfStencilToDiscardIteratorRandomDat
}
}
}
+
+
+__global__
+THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
+void ReplaceKernel(int const N, int* array, int old_value,int new_value)
+{
+ if(threadIdx.x == 0)
+ {
+ thrust::device_ptr<int> begin(array);
+ thrust::device_ptr<int> end(array + N);
+ thrust::replace(thrust::hip::par, begin, end,old_value,new_value);
+ }
+}
+
+TEST(ReplaceTests, TestReplaceDevice)
+{
+ 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;
+
+ int old_value = get_random_data<int>(1,0,size,seed)[0];
+ int new_value = get_random_data<int>(1,-size,size,seed)[0];
+
+ SCOPED_TRACE(testing::Message() << "with old_value= " <<old_value);
+
+ thrust::replace(h_data.begin(), h_data.end(),old_value,new_value);
+ hipLaunchKernelGGL(ReplaceKernel,
+ dim3(1, 1, 1),
+ dim3(128, 1, 1),
+ 0,
+ 0,
+ size,
+ thrust::raw_pointer_cast(&d_data[0]),
+ old_value,
+ new_value);
+
+ ASSERT_EQ(h_data, d_data);
+ }
+ }
+}

MODIFY
test/test_scan.cpp
test/test_scan.cpp
@@ -686,3 +686,95 @@ TYPED_TEST(ScanVectorTests, TestInclusiveScanWithIndirection)
ASSERT_EQ(data[5], T(0));
ASSERT_EQ(data[6], T(1));
}
+
+
+__global__
+THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
+void InclusiveScanKernel(int const N, int* in_array, int *out_array)
+{
+ 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> out_begin(out_array);
+
+ thrust::inclusive_scan(thrust::hip::par, in_begin, in_end,out_begin);
+ }
+}
+
+TEST(ScanTests, TestInclusiveScanDevice)
+{
+ 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::host_vector<int> h_output(size);
+
+ thrust::inclusive_scan(h_data.begin(),h_data.end(),h_output.begin());
+ thrust::device_vector<int> d_data = h_data;
+ thrust::device_vector<int> d_output(size);
+ hipLaunchKernelGGL(InclusiveScanKernel,
+ 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);
+ }
+ }
+}
+
+
+__global__
+THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
+void ExclusiveScanKernel(int const N, int* in_array, int *out_array)
+{
+ 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> out_begin(out_array);
+
+ thrust::exclusive_scan(thrust::hip::par, in_begin, in_end,out_begin);
+ }
+}
+
+TEST(ScanTests, TestExclusiveScanDevice)
+{
+ 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::host_vector<int> h_output(size);
+
+ thrust::exclusive_scan(h_data.begin(),h_data.end(),h_output.begin());
+ thrust::device_vector<int> d_data = h_data;
+ thrust::device_vector<int> d_output(size);
+ hipLaunchKernelGGL(ExclusiveScanKernel,
+ 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);
+ }
+ }
+}

MODIFY
test/test_scan_by_key.cpp
test/test_scan_by_key.cpp
@@ -779,3 +779,132 @@ TEST(ScanByKeyTests, TestScanByKeyWithLargeTypes)
//_TestScanByKeyWithLargeTypes<int, 512>();
//_TestScanByKeyWithLargeTypes<int, 1024>();
}
+
+__global__
+THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
+void InclusiveScanByKeyKernel(int const N, int* in_array, int* keys_array, int *out_array)
+{
+ 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(keys_array);
+ thrust::device_ptr<int> keys_end(keys_array+N);
+ thrust::device_ptr<int> out_begin(out_array);
+
+ thrust::inclusive_scan_by_key(thrust::hip::par, keys_begin, keys_end, in_begin, out_begin);
+ }
+}
+
+TEST(ScanByKeyTests, TestInclusiveScanByKeyDevice)
+{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
+ for(auto size : get_sizes())
+ {
+ SCOPED_TRACE(testing::Message() << "with size= " << size);
+
+ thrust::host_vector<int> h_keys(size);
+ thrust::default_random_engine rng;
+ for(size_t i = 0, k = 0; i < size; i++)
+ {
+ h_keys[i] = k;
+ if(rng() % 10 == 0)
+ k++;
+ }
+ thrust::device_vector<int> d_keys = h_keys;
+
+ for(auto seed : get_seeds())
+ {
+ SCOPED_TRACE(testing::Message() << "with seed= " << seed);
+
+ thrust::host_vector<int> h_vals = get_random_data<int>(size, std::numeric_limits<int>::min(), std::numeric_limits<int>::max(), seed);
+ for(size_t i = 0; i < size; i++)
+ {
+ h_vals[i] = i % 10;
+ }
+ thrust::device_vector<int> d_vals = h_vals;
+
+ thrust::host_vector<int> h_output(size);
+ thrust::device_vector<int> d_output(size);
+
+ thrust::inclusive_scan_by_key(h_keys.begin(), h_keys.end(), h_vals.begin(), h_output.begin());
+
+ hipLaunchKernelGGL(InclusiveScanByKeyKernel,
+ dim3(1, 1, 1),
+ dim3(128, 1, 1),
+ 0,
+ 0,
+ size,
+ thrust::raw_pointer_cast(&d_vals[0]),
+ thrust::raw_pointer_cast(&d_keys[0]),
+ thrust::raw_pointer_cast(&d_output[0]));
+
+ ASSERT_EQ(d_output, h_output);
+ }
+ }
+}
+
+
+__global__
+THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
+void ExclusiveScanByKeyKernel(int const N, int* in_array, int* keys_array, int *out_array)
+{
+ 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(keys_array);
+ thrust::device_ptr<int> keys_end(keys_array+N);
+ thrust::device_ptr<int> out_begin(out_array);
+
+ thrust::exclusive_scan_by_key(thrust::hip::par, keys_begin, keys_end, in_begin, out_begin);
+ }
+}
+
+TEST(ScanByKeyTests, TestExclusiveScanByKeyDevice)
+{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
+ for(auto size : get_sizes())
+ {
+ SCOPED_TRACE(testing::Message() << "with size= " << size);
+
+ thrust::host_vector<int> h_keys(size);
+ thrust::default_random_engine rng;
+ for(size_t i = 0, k = 0; i < size; i++)
+ {
+ h_keys[i] = k;
+ if(rng() % 10 == 0)
+ k++;
+ }
+ thrust::device_vector<int> d_keys = h_keys;
+
+ for(auto seed : get_seeds())
+ {
+ SCOPED_TRACE(testing::Message() << "with seed= " << seed);
+
+ thrust::host_vector<int> h_vals = get_random_data<int>(size, std::numeric_limits<int>::min(), std::numeric_limits<int>::max(), seed);
+ for(size_t i = 0; i < size; i++)
+ {
+ h_vals[i] = i % 10;
+ }
+ thrust::device_vector<int> d_vals = h_vals;
+
+ thrust::host_vector<int> h_output(size);
+ thrust::device_vector<int> d_output(size);
+
+ thrust::exclusive_scan_by_key(h_keys.begin(), h_keys.end(), h_vals.begin(), h_output.begin());
+
+ hipLaunchKernelGGL(ExclusiveScanByKeyKernel,
+ dim3(1, 1, 1),
+ dim3(128, 1, 1),
+ 0,
+ 0,
+ size,
+ thrust::raw_pointer_cast(&d_vals[0]),
+ thrust::raw_pointer_cast(&d_keys[0]),
+ thrust::raw_pointer_cast(&d_output[0]));
+
+ ASSERT_EQ(d_output, h_output);
+ }
+ }
+}

MODIFY
test/test_stable_sort.cpp
test/test_stable_sort.cpp
@@ -188,3 +188,46 @@ TYPED_TEST(StableSortVectorTests, TestStableSortWithIndirection)
ASSERT_EQ(data[5], T(5));
ASSERT_EQ(data[6], T(2));
}
+
+
+__global__
+THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
+void StableSortKernel(int const N, int* array)
+{
+ if(threadIdx.x == 0)
+ {
+ thrust::device_ptr<int> begin(array);
+ thrust::device_ptr<int> end(array + N);
+ thrust::stable_sort(thrust::hip::par, begin, end);
+ }
+}
+
+TEST(StableSortTests, TestStableSortDevice)
+{
+ 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::stable_sort(h_data.begin(), h_data.end());
+ hipLaunchKernelGGL(StableSortKernel,
+ dim3(1, 1, 1),
+ dim3(128, 1, 1),
+ 0,
+ 0,
+ size,
+ thrust::raw_pointer_cast(&d_data[0]));
+
+ ASSERT_EQ(h_data, d_data);
+ }
+ }
+}

MODIFY
test/test_stable_sort_by_key.cpp
test/test_stable_sort_by_key.cpp
@@ -180,3 +180,56 @@ TYPED_TEST(StableSortByKeyVectorPrimitiveTests, TestStableSortByKey)
}
}
}
+
+__global__
+THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
+void StableSortByKeyKernel(int const N, int* keys, short* values)
+{
+ if(threadIdx.x == 0)
+ {
+ thrust::device_ptr<int> keys_begin(keys);
+ thrust::device_ptr<int> keys_end(keys + N);
+ thrust::device_ptr<short> val(values);
+ thrust::stable_sort_by_key(thrust::hip::par, keys_begin, keys_end, val);
+ }
+}
+
+TEST(StableSortByKeyTests, TestStableSortByKeyDevice)
+{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
+
+ for (auto size: std::vector<size_t>{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_keys = get_random_data<int>(size, 0, size, seed);
+
+ thrust::host_vector<short> h_values
+ = get_random_data<short>(size,
+ std::numeric_limits<short>::min(),
+ std::numeric_limits<short>::max(),
+ seed);
+
+ thrust::device_vector<int> d_keys = h_keys;
+ thrust::device_vector<short> d_values = h_values;
+
+ thrust::stable_sort_by_key(h_keys.begin(), h_keys.end(), h_values.begin());
+ hipLaunchKernelGGL(StableSortByKeyKernel,
+ 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]));
+
+ ASSERT_EQ(h_keys, d_keys);
+ // Only keys are compared here, the sequential stable_merge_sort that's used in
+ // CUDA and HIP don't generate the correct value sorting
+ }
+ }
+}

MODIFY
test/test_transform.cpp
test/test_transform.cpp
@@ -1229,3 +1229,111 @@ TYPED_TEST(TransformVectorTests, TestTransformWithIndirection)
ASSERT_EQ(output[5], T(1));
ASSERT_EQ(output[6], T(1));
}
+
+
+__global__
+THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
+void UnaryTransformKernel(int const N, int* input, int *output)
+{
+ if(threadIdx.x == 0)
+ {
+ thrust::device_ptr<int> in_begin(input);
+ thrust::device_ptr<int> in_end(input + N);
+ thrust::device_ptr<int> out_begin(output);
+ thrust::transform(thrust::hip::par, in_begin, in_end, out_begin, unary_transform<int>());
+ }
+}
+TEST(TransformTests, TestUnaryTransformDevice)
+{
+
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
+
+ for(auto size : get_sizes())
+ {
+ SCOPED_TRACE(testing::Message() << "with size = " << size);
+
+ thrust::host_vector<int> h_input(size);
+ for(size_t i = 0; i < size; i++)
+ {
+ h_input[i] = i;
+ }
+
+ // Calculate expected results on host
+ thrust::host_vector<int> expected(size);
+ thrust::transform(h_input.begin(), h_input.end(), expected.begin(), unary_transform<int>());
+
+ thrust::device_vector<int> d_input(h_input);
+ thrust::device_vector<int> d_output(size);
+
+ hipLaunchKernelGGL(UnaryTransformKernel,
+ dim3(1, 1, 1),
+ dim3(128, 1, 1),
+ 0,
+ 0,
+ size,
+ thrust::raw_pointer_cast(&d_input[0]),
+ thrust::raw_pointer_cast(&d_output[0]));
+
+
+ ASSERT_EQ(expected,d_output);
+ }
+}
+
+
+__global__
+THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
+void BinaryTransformKernel(int const N, int* input1, int* input2, int *output)
+{
+ if(threadIdx.x == 0)
+ {
+ thrust::device_ptr<int> in1_begin(input1);
+ thrust::device_ptr<int> in1_end(input1 + N);
+ thrust::device_ptr<int> in2_begin(input2);
+
+ thrust::device_ptr<int> out_begin(output);
+ thrust::transform(thrust::hip::par, in1_begin, in1_end,in2_begin, out_begin, binary_transform<int>());
+ }
+}
+
+
+TEST(TransformTests, TestBinaryTransformDevice)
+{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
+
+ for(auto size : get_sizes())
+ {
+ SCOPED_TRACE(testing::Message() << "with size = " << size);
+
+ thrust::host_vector<int> h_input1(size);
+ thrust::host_vector<int> h_input2(size);
+ for(size_t i = 0; i < size; i++)
+ {
+ h_input1[i] = i * 3;
+ h_input2[i] = i;
+ }
+
+ // Calculate expected results on host
+ thrust::host_vector<int> expected(size);
+ thrust::transform(h_input1.begin(),
+ h_input1.end(),
+ h_input2.begin(),
+ expected.begin(),
+ binary_transform<int>());
+
+ thrust::device_vector<int> d_input1(h_input1);
+ thrust::device_vector<int> d_input2(h_input2);
+ thrust::device_vector<int> d_output(size);
+ hipLaunchKernelGGL(BinaryTransformKernel,
+ dim3(1, 1, 1),
+ dim3(128, 1, 1),
+ 0,
+ 0,
+ size,
+ thrust::raw_pointer_cast(&d_input1[0]),
+ thrust::raw_pointer_cast(&d_input2[0]),
+ thrust::raw_pointer_cast(&d_output[0]));
+
+ ASSERT_EQ(expected,d_output);
+
+ }
+}

Implemented additional test cases for device side algorithms