Implemented runtime error for unsupported algorithms. Disabled tests for the same
MODIFY
test/test_partition.cpp
test/test_partition.cpp
@@ -1917,7 +1917,7 @@ TEST(PartitionTests,TestPartitionCopyDevice)
// thrust::device_ptr<int> begin(array);
// thrust::device_ptr<int> end(array + N);
// //TODO: The thrust::hip::par throw exception, we should fix it
-// thrust::stable_partition(thrust::seq, begin, end,is_even<int>());
+// thrust::stable_partition(thrust::hip::par, begin, end,is_even<int>());
//
// }
// }

MODIFY
test/test_sort.cpp
test/test_sort.cpp
@@ -437,44 +437,44 @@ TEST(SortTests, TestSortBoolDescending)
}
//TODO: refactor this test into a different set of tests
-__global__
-THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
-void SortKernel(int const N, int* array)
-{
- if(threadIdx.x == 0)
- {
- thrust::device_ptr<int> begin(array);
- thrust::device_ptr<int> end(array + N);
- thrust::sort(thrust::hip::par, begin, end);
- }
-}
-
-TEST(SortTests, TestSortDevice)
-{
- 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, size, seed);
-
- thrust::device_vector<int> d_data = h_data;
-
- thrust::sort(h_data.begin(), h_data.end());
- hipLaunchKernelGGL(SortKernel,
- dim3(1, 1, 1),
- dim3(128, 1, 1),
- 0,
- 0,
- size,
- thrust::raw_pointer_cast(&d_data[0]));
-
- ASSERT_EQ(h_data, d_data);
- }
- }
-}
+// __global__
+// THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
+// void SortKernel(int const N, int* array)
+// {
+// if(threadIdx.x == 0)
+// {
+// thrust::device_ptr<int> begin(array);
+// thrust::device_ptr<int> end(array + N);
+// thrust::sort(thrust::hip::par, begin, end);
+// }
+// }
+//
+// TEST(SortTests, TestSortDevice)
+// {
+// 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, size, seed);
+//
+// thrust::device_vector<int> d_data = h_data;
+//
+// thrust::sort(h_data.begin(), h_data.end());
+// hipLaunchKernelGGL(SortKernel,
+// 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_sort_by_key.cpp
test/test_sort_by_key.cpp
@@ -266,8 +266,7 @@ TEST(SortByKeyTests, TestSortByKeyBoolDescending)
// thrust::device_ptr<int> keys_begin(keys);
// thrust::device_ptr<int> keys_end(keys + N);
// thrust::device_ptr<short> val(values);
-// //TODO: The thrust::hip::par throw exception, we should fix it
-// thrust::sort_by_key(thrust::seq, keys_begin, keys_end, val);
+// thrust::sort_by_key(thrust::hip::par, keys_begin, keys_end, val);
// }
// }
//

MODIFY
test/test_stable_sort.cpp
test/test_stable_sort.cpp
@@ -190,44 +190,44 @@ TYPED_TEST(StableSortVectorTests, TestStableSortWithIndirection)
}
-__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 : 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, 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);
- }
- }
-}
+// __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 : 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, 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
@@ -193,7 +193,7 @@ TYPED_TEST(StableSortByKeyVectorPrimitiveTests, TestStableSortByKey)
// thrust::device_ptr<int> keys_end(keys + N);
// thrust::device_ptr<short> val(values);
// //TODO: The thrust::hip::par throw exception, we should fix it
-// thrust::stable_sort_by_key(thrust::seq, keys_begin, keys_end, val);
+// thrust::stable_sort_by_key(thrust::hip::par, keys_begin, keys_end, val);
// }
// }
//

MODIFY
thrust/system/detail/sequential/partition.h
thrust/system/detail/sequential/partition.h
@@ -175,6 +175,11 @@ __host__ __device__
typedef thrust::detail::temporary_array<T,DerivedPolicy> TempRange;
typedef typename TempRange::iterator TempIterator;
+ #ifdef __HIP_DEVICE_COMPILE__
+ THRUST_HIP_PRINTF("ERROR:Stable Partition is not currently supported on the device. \n");
+ abort();
+ #endif
+
TempRange temp(exec, first, last);
for(TempIterator iter = temp.begin(); iter != temp.end(); ++iter)
@@ -336,4 +341,3 @@ __host__ __device__
} // end namespace detail
} // end namespace system
} // end namespace thrust
-

MODIFY
thrust/system/detail/sequential/stable_merge_sort.inl
thrust/system/detail/sequential/stable_merge_sort.inl
@@ -46,6 +46,11 @@ void inplace_merge(sequential::execution_policy<DerivedPolicy> &exec,
{
typedef typename thrust::iterator_value<RandomAccessIterator>::type value_type;
+ #ifdef __HIP_DEVICE_COMPILE__
+ THRUST_HIP_PRINTF("ERROR:In Place Merge is not currently supported on the device. \n");
+ abort();
+ #endif
+
thrust::detail::temporary_array<value_type, DerivedPolicy> a(exec, first, middle);
thrust::detail::temporary_array<value_type, DerivedPolicy> b(exec, middle, last);
@@ -71,6 +76,11 @@ void inplace_merge_by_key(sequential::execution_policy<DerivedPolicy> &exec,
RandomAccessIterator2 middle2 = first2 + (middle1 - first1);
RandomAccessIterator2 last2 = first2 + (last1 - first1);
+ #ifdef __HIP_DEVICE_COMPILE__
+ THRUST_HIP_PRINTF("ERROR:Inplace Merge by Key is not currently supported on the device. \n");
+ abort();
+ #endif
+
thrust::detail::temporary_array<value_type1, DerivedPolicy> lhs1(exec, first1, middle1);
thrust::detail::temporary_array<value_type1, DerivedPolicy> rhs1(exec, middle1, last1);
thrust::detail::temporary_array<value_type2, DerivedPolicy> lhs2(exec, first2, middle2);
@@ -219,6 +229,10 @@ void iterative_stable_merge_sort(sequential::execution_policy<DerivedPolicy> &ex
}
else
{
+ #ifdef __HIP_DEVICE_COMPILE__
+ THRUST_HIP_PRINTF("ERROR:Iterative Stable Merge Sort is not currently supported on the device. \n");
+ abort();
+ #endif
thrust::detail::temporary_array<value_type, DerivedPolicy> temp(exec, n);
// insertion sort each 32 element partition
@@ -277,6 +291,10 @@ void iterative_stable_merge_sort_by_key(sequential::execution_policy<DerivedPoli
}
else
{
+ #ifdef __HIP_DEVICE_COMPILE__
+ THRUST_HIP_PRINTF("ERROR:Iterative Iterative Stable Merge Sort by Key is not currently supported on the device. \n");
+ abort();
+ #endif
thrust::detail::temporary_array<value_type1, DerivedPolicy> keys_temp(exec, n);
thrust::detail::temporary_array<value_type2, DerivedPolicy> values_temp(exec, n);

MODIFY
thrust/system/detail/sequential/stable_radix_sort.inl
thrust/system/detail/sequential/stable_radix_sort.inl
@@ -561,6 +561,11 @@ void stable_radix_sort(sequential::execution_policy<DerivedPolicy> &exec,
size_t N = last - first;
+ #ifdef __HIP_DEVICE_COMPILE__
+ THRUST_HIP_PRINTF("ERROR:Stable Radix Sort is not currently supported on the device. \n");
+ abort();
+ #endif
+
thrust::detail::temporary_array<KeyType, DerivedPolicy> temp(exec, N);
radix_sort_detail::radix_sort(exec, first, temp.begin(), N);
@@ -581,6 +586,11 @@ void stable_radix_sort_by_key(sequential::execution_policy<DerivedPolicy> &exec,
size_t N = last1 - first1;
+ #ifdef __HIP_DEVICE_COMPILE__
+ THRUST_HIP_PRINTF("ERROR: Stable Radix Sort by Key is not currently supported on the device. \n");
+ abort();
+ #endif
+
thrust::detail::temporary_array<KeyType, DerivedPolicy> temp1(exec, N);
thrust::detail::temporary_array<ValueType, DerivedPolicy> temp2(exec, N);
@@ -592,4 +602,3 @@ void stable_radix_sort_by_key(sequential::execution_policy<DerivedPolicy> &exec,
} // end namespace detail
} // end namespace system
} // end namespace thrust
-

Implemented runtime error for unsupported algorithms. Disabled tests for the same