Changed execution type to sequential for on device algorithm test cases
MODIFY
test/test_adjacent_difference.cpp
test/test_adjacent_difference.cpp
@@ -246,8 +246,8 @@ void AdjacentDifferenceKernel(int const N, int* array)
{
thrust::device_ptr<int> begin(array);
thrust::device_ptr<int> end(array + N);
- thrust::advance(begin,2);
- thrust::adjacent_difference(thrust::hip::par, begin,end,begin);
+ // thrust::advance(begin,2);
+ thrust::adjacent_difference(thrust::seq, begin,end,begin);
}
}

MODIFY
test/test_find.cpp
test/test_find.cpp
@@ -395,7 +395,7 @@ void FindKernel(int const N, int* in_array, int value, int *out_array)
thrust::device_ptr<int> in_begin(in_array);
thrust::device_ptr<int> in_end(in_array + N);
- auto x = thrust::find(thrust::hip::par, in_begin, in_end,value);
+ auto x = thrust::find(thrust::seq, in_begin, in_end,value);
out_array[0] = x - in_begin;
}
}

MODIFY
test/test_merge.cpp
test/test_merge.cpp
@@ -290,8 +290,8 @@ void MergeKernel(int const N, T* inA_array, T* inB_array, T *out_array)
thrust::device_ptr<int> inB_end(inB_array + N);
thrust::device_ptr<int> out_begin(out_array);
- thrust::merge(thrust::hip::par, inA_begin, inA_end, inB_begin, inB_end,out_begin);
- thrust::merge_by_key(thrust::hip::par, inA_begin, inA_end, inB_begin, inB_end,inA_begin,inB_begin,out_begin,out_begin);
+ 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);
}
}
@@ -334,8 +334,7 @@ TEST(PrimitiveMergeTests, TestMergeDevice)
h_a.end(),
h_b.begin(),
h_b.end(),
- h_result.begin(),
- thrust::greater<T>());
+ h_result.begin());
hipLaunchKernelGGL(MergeKernel,
dim3(1, 1, 1),

MODIFY
test/test_partition.cpp
test/test_partition.cpp
@@ -1805,7 +1805,7 @@ void PartitionKernel(int const N, int* array)
{
thrust::device_ptr<int> begin(array);
thrust::device_ptr<int> end(array + N);
- thrust::partition(thrust::hip::par, begin, end,is_even<int>());
+ thrust::partition(thrust::seq, begin, end,is_even<int>());
}
}

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::hip::par, begin, end,remove_value) - begin;
+ int size = thrust::remove(thrust::seq, begin, end,remove_value) - begin;
}
}

MODIFY
test/test_scan.cpp
test/test_scan.cpp
@@ -698,7 +698,7 @@ void InclusiveScanKernel(int const N, int* in_array, int *out_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);
+ thrust::inclusive_scan(thrust::seq, in_begin, in_end,out_begin);
}
}
@@ -744,7 +744,7 @@ void ExclusiveScanKernel(int const N, int* in_array, int *out_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);
+ thrust::exclusive_scan(thrust::seq, in_begin, in_end,out_begin);
}
}

MODIFY
test/test_scan_by_key.cpp
test/test_scan_by_key.cpp
@@ -792,7 +792,7 @@ void InclusiveScanByKeyKernel(int const N, int* in_array, int* keys_array, int *
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);
+ thrust::inclusive_scan_by_key(thrust::seq, keys_begin, keys_end, in_begin, out_begin);
}
}
@@ -857,7 +857,7 @@ void ExclusiveScanByKeyKernel(int const N, int* in_array, int* keys_array, int *
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);
+ thrust::exclusive_scan_by_key(thrust::seq, keys_begin, keys_end, in_begin, out_begin);
}
}

MODIFY
test/test_uninitialized_copy.cpp
test/test_uninitialized_copy.cpp
@@ -279,3 +279,47 @@ TEST(UninitializedCopyTests, TestUninitializedCopyNNonPODHost)
ASSERT_EQ(false, x.copy_constructed_on_device);
ASSERT_EQ(true, x.copy_constructed_on_host);
}
+
+
+__global__
+THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
+void UninitializedCopyKernel(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::uninitialized_copy(thrust::hip::par, in_begin, in_end,out_begin);
+ }
+}
+
+TEST(UninitializedCopyTests, TestUninitializedCopyDevice)
+{
+ 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::device_vector<int> d_output(size);
+ hipLaunchKernelGGL(UninitializedCopyKernel,
+ 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_data, d_output);
+ }
+ }
+}

MODIFY
test/test_uninitialized_fill.cpp
test/test_uninitialized_fill.cpp
@@ -311,3 +311,50 @@ TEST(UninitializedFillTests, TestUninitializedFillNNonPOD)
thrust::device_free(v);
}
*/
+
+
+
+__global__
+THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
+void UninitializedFillKernel(int const N, int* array, int fill_value)
+{
+ if(threadIdx.x == 0)
+ {
+ thrust::device_ptr<int> begin(array);
+ thrust::device_ptr<int> end(array + N);
+ thrust::uninitialized_fill(thrust::hip::par, begin, end,fill_value);
+ }
+}
+
+TEST(UninitializedFillTests, TestUninitializedFillDevice)
+{
+ 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 fill_value = get_random_data<int>(1,0,size,seed)[0];
+ SCOPED_TRACE(testing::Message() << "with fill_value= " <<fill_value);
+
+ thrust::uninitialized_fill(h_data.begin(), h_data.end(),fill_value);
+ hipLaunchKernelGGL(UninitializedFillKernel,
+ dim3(1, 1, 1),
+ dim3(128, 1, 1),
+ 0,
+ 0,
+ size,
+ thrust::raw_pointer_cast(&d_data[0]),
+ fill_value);
+
+ ASSERT_EQ(h_data, d_data);
+ }
+ }
+}

MODIFY
test/test_unique.cpp
test/test_unique.cpp
@@ -302,3 +302,62 @@ TYPED_TEST(UniqueIntegralTests, TestUniqueCopyToDiscardIterator)
}
}
}
+
+
+__global__
+THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
+void UniqueKernel(int const N, int *in_array, 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_vector<int>::iterator last = thrust::unique(thrust::seq, in_begin, in_end);
+ out_size[0] = last - thrust::device_vector<int>::iterator(in_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::device_vector<int> d_data = h_data;
+ thrust::device_vector<int> d_output_size(1,0);
+
+ typename thrust::host_vector<int>::iterator h_new_last;
+ typename thrust::device_vector<int>::iterator d_new_last;
+
+ h_new_last = thrust::unique(h_data.begin(), h_data.end());
+
+ hipLaunchKernelGGL(UniqueKernel,
+ dim3(1, 1, 1),
+ dim3(128, 1, 1),
+ 0,
+ 0,
+ size,
+ thrust::raw_pointer_cast(&d_data[0]),
+ thrust::raw_pointer_cast(&d_output_size[0]));
+
+
+ ASSERT_EQ(h_new_last - h_data.begin(),d_output_size[0]);
+
+ h_data.resize(h_new_last - h_data.begin());
+ d_data.resize(d_output_size[0]);
+
+ ASSERT_EQ(h_data, d_data);
+ }
+ }
+}

MODIFY
thrust/system/hip/detail/adjacent_difference.h
thrust/system/hip/detail/adjacent_difference.h
@@ -414,54 +414,18 @@ adjacent_difference(execution_policy<Derived>& policy,
OutputIt result,
BinaryOp binary_op)
{
- // struct workaround is required for HIP-clang
- // THRUST_HIP_PRESERVE_KERNELS_WORKAROUND is required for HCC
- struct workaround
- {
- __host__
- static void par(execution_policy<Derived>& policy,
- InputIt first,
- InputIt last,
- OutputIt& result,
- BinaryOp binary_op)
- {
- #if __HCC__ && __HIP_DEVICE_COMPILE__
+ OutputIt ret = result;
THRUST_HIP_PRESERVE_KERNELS_WORKAROUND(
(__adjacent_difference::adjacent_difference<Derived, InputIt, OutputIt, BinaryOp>)
);
- #else
- result = __adjacent_difference::adjacent_difference(
- policy,
- first,
- last,
- result,
- binary_op
- );
- #endif
- }
- __device__
- static void seq(execution_policy<Derived>& policy,
- InputIt first,
- InputIt last,
- OutputIt& result,
- BinaryOp binary_op)
- {
- result = thrust::adjacent_difference(
- cvt_to_seq(derived_cast(policy)),
- first,
- last,
- result,
- binary_op
- );
- }
- };
#if __THRUST_HAS_HIPRT__
- workaround::par(policy, first, last, result, binary_op);
- #else
- workaround::seq(policy, first, last, result, binary_op);
- #endif
+ ret = __adjacent_difference::adjacent_difference(policy, first, last, result, binary_op);
+ #else // __THRUST_HAS_HIPRT__
+ ret = thrust::adjacent_difference(
+ cvt_to_seq(derived_cast(policy)), first, last, result, binary_op);
+ #endif // __THRUST_HAS_HIPRT__
- return result;
+ return ret;
}
template <class Derived, class InputIt, class OutputIt>

Changed execution type to sequential for on device algorithm test cases