Fix sequential stable_merge_sort when called from kernel
MODIFY
test/test_sort.cpp
test/test_sort.cpp
@@ -343,3 +343,38 @@ TEST(SortTests, TestSortBoolDescending)
ASSERT_EQ(h_data, d_data);
}
+
+//TODO: refactor this test into a different set of tests
+__global__ 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)
+{
+ std::vector<size_t> sizes = {
+ 0, 1, 2, 4, 6, 12, 16, 24, 32,
+ 64, 84, 128, 160, 256
+ };
+
+ for(auto size : sizes)
+ {
+ thrust::host_vector<int> h_data = get_random_data<int>(
+ size, 0, size);
+
+ 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
@@ -206,3 +206,46 @@ TEST(SortByKeyTests, TestSortByKeyBoolDescending)
ASSERT_EQ(h_keys, d_keys);
ASSERT_EQ(h_values, d_values);
}
+
+//TODO: refactor this test into a different set of tests
+__global__ void SortByKeyKernel(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::sort_by_key(thrust::hip::par, keys_begin, keys_end, val);
+ }
+}
+
+TEST(SortByKeyTests, TestSortByKeyDevice)
+{
+ std::vector<size_t> sizes = {
+ 0, 1, 2, 4, 6, 12, 16, 24, 32,
+ 64, 84, 128, 160, 256
+ };
+
+ for(auto size : sizes)
+ {
+ thrust::host_vector<int> h_keys = get_random_data<int>(
+ size, 0, size);
+
+ thrust::host_vector<short> h_values = get_random_data<short>(
+ size, std::numeric_limits<short>::min(), std::numeric_limits<short>::max());
+
+ thrust::device_vector<int> d_keys = h_keys;
+ thrust::device_vector<short> d_values = h_values;
+
+ thrust::sort_by_key(h_keys.begin(), h_keys.end(), h_values.begin());
+ hipLaunchKernelGGL(
+ SortByKeyKernel, 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
thrust/system/detail/sequential/stable_merge_sort.inl
thrust/system/detail/sequential/stable_merge_sort.inl
@@ -209,11 +209,19 @@ void iterative_stable_merge_sort(sequential::execution_policy<DerivedPolicy> &ex
typedef typename thrust::iterator_difference<RandomAccessIterator>::type difference_type;
difference_type n = last - first;
+ difference_type partition_size = 32;
+ // WORKAROUND: Memory access fault occurs on HIP when the size of the temp array is less than 16.
+ // This workaround is based on the check that's used in the recursive version.
+ if(n <= partition_size)
+ {
+ insertion_sort_each(first, last, partition_size, comp);
+ }
+ else
+ {
thrust::detail::temporary_array<value_type, DerivedPolicy> temp(exec, n);
// insertion sort each 32 element partition
- difference_type partition_size = 32;
insertion_sort_each(first, last, partition_size, comp);
// ping indicates whether or not the latest data is in the source range [first, last)
@@ -238,6 +246,7 @@ void iterative_stable_merge_sort(sequential::execution_policy<DerivedPolicy> &ex
{
thrust::copy(exec, temp.begin(), temp.end(), first);
} // end if
+ }
} // end iterative_stable_merge_sort()
@@ -257,12 +266,21 @@ void iterative_stable_merge_sort_by_key(sequential::execution_policy<DerivedPoli
typedef typename thrust::iterator_difference<RandomAccessIterator1>::type difference_type;
difference_type n = keys_last - keys_first;
+ difference_type partition_size = 32;
+ // WORKAROUND: Memory access fault occurs on HIP when the size of the temp array is less than 16.
+ // This workaround is based on the check that's used in the recursive version.
+ if(n <= partition_size)
+ {
+ // insertion sort each 32 element partition
+ insertion_sort_each_by_key(keys_first, keys_last, values_first, partition_size, comp);
+ }
+ else
+ {
thrust::detail::temporary_array<value_type1, DerivedPolicy> keys_temp(exec, n);
thrust::detail::temporary_array<value_type2, DerivedPolicy> values_temp(exec, n);
// insertion sort each 32 element partition
- difference_type partition_size = 32;
insertion_sort_each_by_key(keys_first, keys_last, values_first, partition_size, comp);
// ping indicates whether or not the latest data is in the source range [first, last)
@@ -288,6 +306,7 @@ void iterative_stable_merge_sort_by_key(sequential::execution_policy<DerivedPoli
thrust::copy(exec, keys_temp.begin(), keys_temp.end(), keys_first);
thrust::copy(exec, values_temp.begin(), values_temp.end(), values_first);
} // end if
+ }
} // end iterative_stable_merge_sort()

Fix sequential stable_merge_sort when called from kernel