Refactored the hip implementations
MODIFY
thrust/system/hip/detail/adjacent_difference.h
thrust/system/hip/detail/adjacent_difference.h
@@ -366,34 +366,41 @@ adjacent_difference(Policy &policy,
hipStream_t stream = hip_rocprim::stream(policy);
bool debug_sync = THRUST_HIP_DEBUG_SYNC_FLAG;
- hipError_t status;
- status = doit_step(d_temp_storage,
+ if (num_items == 0)
+ return result;
+
+ // Determine temporary device storage requirements.
+ hip_rocprim::throw_on_error(
+ doit_step(d_temp_storage,
temp_storage_bytes,
first,
result,
binary_op,
num_items,
stream,
- debug_sync);
- hip_rocprim::throw_on_error(status, "adjacent_difference failed on 1st step");
+ debug_sync),
+ "adjacent_difference failed on 1st step");
+ // Allocate temporary storage.
d_temp_storage = hip_rocprim::get_memory_buffer(policy, temp_storage_bytes);
hip_rocprim::throw_on_error(hipGetLastError(),
"adjacent_difference failed to get memory buffer");
- status = doit_step(d_temp_storage,
+ hip_rocprim::throw_on_error(
+ doit_step(d_temp_storage,
temp_storage_bytes,
first,
result,
binary_op,
num_items,
stream,
- debug_sync);
- hip_rocprim::throw_on_error(status, "adjacent_difference failed on 2nd step");
+ debug_sync),
+ "adjacent_difference failed on 2nd step");
hip_rocprim::return_memory_buffer(policy, d_temp_storage);
hip_rocprim::throw_on_error(hipGetLastError(),
"adjacent_difference failed to return memory buffer");
+
return result + num_items;
}

MODIFY
thrust/system/hip/detail/binary_search.h
thrust/system/hip/detail/binary_search.h
@@ -67,8 +67,9 @@ lower_bound(Policy& policy,
hipStream_t stream = hip_rocprim::stream(policy);
bool debug_sync = THRUST_HIP_DEBUG_SYNC_FLAG;
- hipError_t status;
- status = rocprim::lower_bound(d_temp_storage,
+ // Determine temporary device storage requirements.
+ hip_rocprim::throw_on_error(
+ rocprim::lower_bound(d_temp_storage,
temp_storage_bytes,
haystack_begin,
needles_begin,
@@ -77,13 +78,16 @@ lower_bound(Policy& policy,
needles_size,
compare_op,
stream,
- debug_sync);
- hip_rocprim::throw_on_error(status, "lower_bound: failed on 1st call");
+ debug_sync),
+ "lower_bound: failed on 1st call");
+ // Allocate temporary storage.
d_temp_storage = hip_rocprim::get_memory_buffer(policy, temp_storage_bytes);
- hip_rocprim::throw_on_error(hipGetLastError(), "lower_bound: failed to get memory buffer");
+ hip_rocprim::throw_on_error(hipGetLastError(),
+ "lower_bound: failed to get memory buffer");
- status = rocprim::lower_bound(d_temp_storage,
+ hip_rocprim::throw_on_error(
+ rocprim::lower_bound(d_temp_storage,
temp_storage_bytes,
haystack_begin,
needles_begin,
@@ -92,14 +96,12 @@ lower_bound(Policy& policy,
needles_size,
compare_op,
stream,
- debug_sync);
- hip_rocprim::throw_on_error(status, "lower_bound: failed on 2nt call");
-
- status = hip_rocprim::synchronize(policy);
- hip_rocprim::throw_on_error(status, "lower_bound: failed to synchronize");
+ debug_sync),
+ "lower_bound: failed on 2nt call");
hip_rocprim::return_memory_buffer(policy, d_temp_storage);
- hip_rocprim::throw_on_error(hipGetLastError(), "lower_bound: failed to return memory buffer");
+ hip_rocprim::throw_on_error(hipGetLastError(),
+ "lower_bound: failed to return memory buffer");
return result + needles_size;
}
@@ -131,8 +133,9 @@ upper_bound(Policy& policy,
hipStream_t stream = hip_rocprim::stream(policy);
bool debug_sync = THRUST_HIP_DEBUG_SYNC_FLAG;
- hipError_t status;
- status = rocprim::upper_bound(d_temp_storage,
+ // Determine temporary device storage requirements.
+ hip_rocprim::throw_on_error(
+ rocprim::upper_bound(d_temp_storage,
temp_storage_bytes,
haystack_begin,
needles_begin,
@@ -141,13 +144,16 @@ upper_bound(Policy& policy,
needles_size,
compare_op,
stream,
- debug_sync);
- hip_rocprim::throw_on_error(status, "upper_bound: failed on 1st call");
+ debug_sync),
+ "upper_bound: failed on 1st call");
+ // Allocate temporary storage.
d_temp_storage = hip_rocprim::get_memory_buffer(policy, temp_storage_bytes);
- hip_rocprim::throw_on_error(hipGetLastError(), "upper_bound: failed to get memory buffer");
+ hip_rocprim::throw_on_error(hipGetLastError(),
+ "upper_bound: failed to get memory buffer");
- status = rocprim::upper_bound(d_temp_storage,
+ hip_rocprim::throw_on_error(
+ rocprim::upper_bound(d_temp_storage,
temp_storage_bytes,
haystack_begin,
needles_begin,
@@ -156,14 +162,12 @@ upper_bound(Policy& policy,
needles_size,
compare_op,
stream,
- debug_sync);
- hip_rocprim::throw_on_error(status, "upper_bound: failed on 2nt call");
-
- status = hip_rocprim::synchronize(policy);
- hip_rocprim::throw_on_error(status, "upper_bound: failed to synchronize");
+ debug_sync),
+ "upper_bound: failed on 2nt call");
hip_rocprim::return_memory_buffer(policy, d_temp_storage);
- hip_rocprim::throw_on_error(hipGetLastError(), "upper_bound: failed to return memory buffer");
+ hip_rocprim::throw_on_error(hipGetLastError(),
+ "upper_bound: failed to return memory buffer");
return result + needles_size;
}
@@ -195,8 +199,9 @@ binary_search(Policy& policy,
hipStream_t stream = hip_rocprim::stream(policy);
bool debug_sync = THRUST_HIP_DEBUG_SYNC_FLAG;
- hipError_t status;
- status = rocprim::binary_search(d_temp_storage,
+ // Determine temporary device storage requirements.
+ hip_rocprim::throw_on_error(
+ rocprim::binary_search(d_temp_storage,
temp_storage_bytes,
haystack_begin,
needles_begin,
@@ -205,13 +210,16 @@ binary_search(Policy& policy,
needles_size,
compare_op,
stream,
- debug_sync);
- hip_rocprim::throw_on_error(status, "binary_search: failed on 1st call");
+ debug_sync),
+ "binary_search: failed on 1st call");
+ // Allocate temporary storage.
d_temp_storage = hip_rocprim::get_memory_buffer(policy, temp_storage_bytes);
- hip_rocprim::throw_on_error(hipGetLastError(), "binary_search: failed to get memory buffer");
+ hip_rocprim::throw_on_error(hipGetLastError(),
+ "binary_search: failed to get memory buffer");
- status = rocprim::binary_search(d_temp_storage,
+ hip_rocprim::throw_on_error(
+ rocprim::binary_search(d_temp_storage,
temp_storage_bytes,
haystack_begin,
needles_begin,
@@ -220,14 +228,12 @@ binary_search(Policy& policy,
needles_size,
compare_op,
stream,
- debug_sync);
- hip_rocprim::throw_on_error(status, "binary_search: failed on 2nt call");
-
- status = hip_rocprim::synchronize(policy);
- hip_rocprim::throw_on_error(status, "binary_search: failed to synchronize");
+ debug_sync),
+ "binary_search: failed on 2nt call");
hip_rocprim::return_memory_buffer(policy, d_temp_storage);
- hip_rocprim::throw_on_error(hipGetLastError(), "binary_search: failed to return memory buffer");
+ hip_rocprim::throw_on_error(hipGetLastError(),
+ "binary_search: failed to return memory buffer");
return result + needles_size;
}

MODIFY
thrust/system/hip/detail/copy_if.h
thrust/system/hip/detail/copy_if.h
@@ -89,8 +89,9 @@ copy_if(Policy &policy,
if (num_items == 0)
return output;
- hipError_t status;
- status = rocprim::select(d_temp_storage,
+ // Determine temporary device storage requirements.
+ hip_rocprim::throw_on_error(
+ rocprim::select(d_temp_storage,
temp_storage_bytes,
first,
output,
@@ -98,10 +99,10 @@ copy_if(Policy &policy,
num_items,
predicate,
stream,
- debug_sync);
- hip_rocprim::throw_on_error(status, "copy_if failed on 1st step");
+ debug_sync),
+ "copy_if failed on 1st step");
- temp_storage_bytes = rocprim::detail::align_size(temp_storage_bytes);
+ // Allocate temporary storage.
d_temp_storage = hip_rocprim::get_memory_buffer(policy, temp_storage_bytes + sizeof(size_type));
hip_rocprim::throw_on_error(hipGetLastError(),
"copy_if failed to get memory buffer");
@@ -109,7 +110,8 @@ copy_if(Policy &policy,
d_num_selected_out = reinterpret_cast<size_type *>(
reinterpret_cast<char *>(d_temp_storage) + temp_storage_bytes);
- status = rocprim::select(d_temp_storage,
+ hip_rocprim::throw_on_error(
+ rocprim::select(d_temp_storage,
temp_storage_bytes,
first,
output,
@@ -117,8 +119,8 @@ copy_if(Policy &policy,
num_items,
predicate,
stream,
- debug_sync);
- hip_rocprim::throw_on_error(status, "copy_if failed on 2nd step");
+ debug_sync),
+ "copy_if failed on 2nd step");
size_type num_selected = get_value(policy, d_num_selected_out);
@@ -156,8 +158,9 @@ copy_if(Policy &policy,
auto flags = thrust::make_transform_iterator(stencil, predicate);
- hipError_t status;
- status = rocprim::select(d_temp_storage,
+ // Determine temporary device storage requirements.
+ hip_rocprim::throw_on_error(
+ rocprim::select(d_temp_storage,
temp_storage_bytes,
first,
flags,
@@ -165,10 +168,10 @@ copy_if(Policy &policy,
d_num_selected_out,
num_items,
stream,
- debug_sync);
- hip_rocprim::throw_on_error(status, "copy_if failed on 1st step");
+ debug_sync),
+ "copy_if failed on 1st step");
- temp_storage_bytes = rocprim::detail::align_size(temp_storage_bytes);
+ // Allocate temporary storage.
d_temp_storage = hip_rocprim::get_memory_buffer(policy, temp_storage_bytes + sizeof(size_type));
hip_rocprim::throw_on_error(hipGetLastError(),
"copy_if failed to get memory buffer");
@@ -176,7 +179,8 @@ copy_if(Policy &policy,
d_num_selected_out = reinterpret_cast<size_type *>(
reinterpret_cast<char *>(d_temp_storage) + temp_storage_bytes);
- status = rocprim::select(d_temp_storage,
+ hip_rocprim::throw_on_error(
+ rocprim::select(d_temp_storage,
temp_storage_bytes,
first,
flags,
@@ -184,8 +188,8 @@ copy_if(Policy &policy,
d_num_selected_out,
num_items,
stream,
- debug_sync);
- hip_rocprim::throw_on_error(status, "copy_if failed on 2nd step");
+ debug_sync),
+ "copy_if failed on 2nd step");
size_type num_selected = get_value(policy, d_num_selected_out);

MODIFY
thrust/system/hip/detail/extrema.h
thrust/system/hip/detail/extrema.h
@@ -162,44 +162,55 @@ namespace hip_rocprim {
Size num_items,
BinaryOp binary_op,
T *)
-
{
- hipStream_t stream = hip_rocprim::stream(policy);
+ if (num_items == 0)
+ hip_rocprim::throw_on_error(hipErrorInvalidValue,
+ "extrema number of items is zero");
+ void * d_temp_storage = NULL;
+ size_t temp_storage_bytes = 0;
+ hipStream_t stream = hip_rocprim::stream(policy);
+ T * d_ret_ptr = NULL;
+ bool debug_sync = THRUST_HIP_DEBUG_SYNC_FLAG;
// Determine temporary device storage requirements.
- T* ret_ptr = NULL;
- size_t tmp_size = 0;
hip_rocprim::throw_on_error(
- rocprim::reduce(nullptr, tmp_size,
- first, ret_ptr, num_items, binary_op,
- stream, THRUST_HIP_DEBUG_SYNC_FLAG),
- "after reduction step 1");
+ rocprim::reduce(d_temp_storage,
+ temp_storage_bytes,
+ first,
+ d_ret_ptr,
+ num_items,
+ binary_op,
+ stream,
+ debug_sync),
+ "extrema failed on 1st step");
// Allocate temporary storage.
+ d_temp_storage = hip_rocprim::get_memory_buffer(policy, sizeof(T) + temp_storage_bytes);
+ hip_rocprim::throw_on_error(hipGetLastError(),
+ "extrema failed to get memory buffer");
- detail::temporary_array<detail::uint8_t, Derived>
- tmp(policy, sizeof(T) + tmp_size);
-
- // Run reduction.
+ d_ret_ptr = reinterpret_cast<T*>(
+ reinterpret_cast<char *>(d_temp_storage) + temp_storage_bytes);
- // `tmp.begin()` yields a `normal_iterator`, which dereferences to a
- // `reference`, which has an `operator&` that returns a `pointer`, which
- // has a `.get` method that returns a raw pointer, which we can (finally)
- // `static_cast` to `void*`.
- //
- // The array was dynamically allocated, so we assume that it's suitably
- // aligned for any type of data. `malloc`/`cudaMalloc`/`new`/`std::allocator`
- // make this guarantee.
- ret_ptr = detail::aligned_reinterpret_cast<T*>((&*tmp.begin()).get());
- void* tmp_ptr = static_cast<void*>((&*(tmp.begin() + sizeof(T))).get());
hip_rocprim::throw_on_error(
- rocprim::reduce(tmp_ptr, tmp_size,
- first, ret_ptr, num_items, binary_op,
- stream, THRUST_HIP_DEBUG_SYNC_FLAG),
- "after reduction step 2");
+ rocprim::reduce(d_temp_storage,
+ temp_storage_bytes,
+ first,
+ d_ret_ptr,
+ num_items,
+ binary_op,
+ stream,
+ debug_sync),
+ "extrema failed on 2nd step");
+
+ T return_value = hip_rocprim::get_value(policy, d_ret_ptr);
- return hip_rocprim::get_value(policy,ret_ptr);
+ hip_rocprim::return_memory_buffer(policy, d_temp_storage);
+ hip_rocprim::throw_on_error(hipGetLastError(),
+ "extrema failed to return memory buffer");
+
+ return return_value;
}
template <template <class, class, class> class ArgFunctor,
@@ -263,14 +274,12 @@ namespace hip_rocprim {
first,
last,
binary_pred);
-
#else // __THRUST_HAS_HIPRT__
ret = thrust::min_element(cvt_to_seq(derived_cast(policy)),
first,
last,
binary_pred);
#endif // __THRUST_HAS_HIPRT__
-
return ret;
}
@@ -371,14 +380,12 @@ namespace hip_rocprim {
ret = thrust::make_pair(first + get<1>(get<0>(result)),
first + get<1>(get<1>(result)));
-
#else // __THRUST_HAS_HIPRT__
ret = thrust::minmax_element(cvt_to_seq(derived_cast(policy)),
first,
last,
binary_pred);
#endif // __THRUST_HAS_HIPRT__
-
return ret;
}

MODIFY
thrust/system/hip/detail/merge.h
thrust/system/hip/detail/merge.h
@@ -95,8 +95,9 @@ merge(execution_policy<Derived>& policy,
hipStream_t stream = hip_rocprim::stream(policy);
bool debug_sync = THRUST_HIP_DEBUG_SYNC_FLAG;
- hipError_t status;
- status = rocprim::merge( d_temp_storage,
+ // Determine temporary device storage requirements.
+ hip_rocprim::throw_on_error(
+ rocprim::merge(d_temp_storage,
temp_storage_bytes,
keys1_first,
keys2_first,
@@ -105,16 +106,16 @@ merge(execution_policy<Derived>& policy,
input2_size,
compare_op,
stream,
- debug_sync
- );
- hip_rocprim::throw_on_error(status, "merge failed on 1st step");
+ debug_sync),
+ "merge failed on 1st step");
- temp_storage_bytes = rocprim::detail::align_size(temp_storage_bytes);
+ // Allocate temporary storage.
d_temp_storage = hip_rocprim::get_memory_buffer(policy, temp_storage_bytes);
hip_rocprim::throw_on_error(hipGetLastError(),
"merge failed to get memory buffer");
- status = rocprim::merge( d_temp_storage,
+ hip_rocprim::throw_on_error(
+ rocprim::merge(d_temp_storage,
temp_storage_bytes,
keys1_first,
keys2_first,
@@ -123,9 +124,8 @@ merge(execution_policy<Derived>& policy,
input2_size,
compare_op,
stream,
- debug_sync
- );
- hip_rocprim::throw_on_error(status, "merge failed on 2nd step");
+ debug_sync),
+ "merge failed on 2nd step");
hip_rocprim::return_memory_buffer(policy, d_temp_storage);
hip_rocprim::throw_on_error(hipGetLastError(),
@@ -170,8 +170,9 @@ merge(Policy& policy,
hipStream_t stream = hip_rocprim::stream(policy);
bool debug_sync = THRUST_HIP_DEBUG_SYNC_FLAG;
- hipError_t status;
- status = rocprim::merge( d_temp_storage,
+ // Determine temporary device storage requirements.
+ hip_rocprim::throw_on_error(
+ rocprim::merge(d_temp_storage,
temp_storage_bytes,
rocprim::make_zip_iterator(rocprim::make_tuple(keys1_first, items1_first)),
rocprim::make_zip_iterator(rocprim::make_tuple(keys2_first, items2_first)),
@@ -180,16 +181,16 @@ merge(Policy& policy,
input2_size,
wrapped_binary_pred,
stream,
- debug_sync
- );
- hip_rocprim::throw_on_error(status, "merge_by_key failed on 1st step");
+ debug_sync),
+ "merge_by_key failed on 1st step");
- temp_storage_bytes = rocprim::detail::align_size(temp_storage_bytes);
+ // Allocate temporary storage.
d_temp_storage = hip_rocprim::get_memory_buffer(policy, temp_storage_bytes);
hip_rocprim::throw_on_error(hipGetLastError(),
"merge_by_key failed to get memory buffer");
- status = rocprim::merge( d_temp_storage,
+ hip_rocprim::throw_on_error(
+ rocprim::merge(d_temp_storage,
temp_storage_bytes,
rocprim::make_zip_iterator(rocprim::make_tuple(keys1_first, items1_first)),
rocprim::make_zip_iterator(rocprim::make_tuple(keys2_first, items2_first)),
@@ -198,9 +199,8 @@ merge(Policy& policy,
input2_size,
wrapped_binary_pred,
stream,
- debug_sync
- );
- hip_rocprim::throw_on_error(status, "merge_by_key failed on 2nd step");
+ debug_sync),
+ "merge_by_key failed on 2nd step");
hip_rocprim::return_memory_buffer(policy, d_temp_storage);
hip_rocprim::throw_on_error(hipGetLastError(),

MODIFY
thrust/system/hip/detail/partition.h
thrust/system/hip/detail/partition.h
@@ -189,7 +189,7 @@ namespace __partition
bool debug_sync = THRUST_HIP_DEBUG_SYNC_FLAG;
value_type* d_partition_out = NULL;
-
+ // Determine temporary device storage requirements.
hip_rocprim::throw_on_error(
rocprim::partition(d_temp_storage,
temp_storage_bytes,
@@ -203,7 +203,6 @@ namespace __partition
"partition failed on 1st step");
// Allocate temporary storage.
- temp_storage_bytes = rocprim::detail::align_size(temp_storage_bytes);
d_temp_storage = hip_rocprim::get_memory_buffer(policy, sizeof(size_type) +
temp_storage_bytes +
sizeof(value_type) * num_items );
@@ -234,9 +233,14 @@ namespace __partition
num_selected = get_value(policy, d_num_selected_out);
}
+ // fill the values
hip_rocprim::parallel_for(
policy,
- partition_fill<SINGLE_OUTPUT, SelectedOutIt, RejectedOutIt, size_type, value_type>(selected_result,
+ partition_fill<SINGLE_OUTPUT,
+ SelectedOutIt,
+ RejectedOutIt,
+ size_type,
+ value_type>(selected_result,
rejected_result,
d_partition_out,
num_selected,
@@ -282,6 +286,7 @@ namespace __partition
value_type* d_partition_out = NULL;
bool* d_flags = NULL;
+ // Determine temporary device storage requirements.
hip_rocprim::throw_on_error(
rocprim::partition(d_temp_storage,
temp_storage_bytes,
@@ -295,7 +300,6 @@ namespace __partition
"partition failed on 1st step");
// Allocate temporary storage.
- temp_storage_bytes = rocprim::detail::align_size(temp_storage_bytes);
d_temp_storage = hip_rocprim::get_memory_buffer(policy, sizeof(size_type) +
temp_storage_bytes +
sizeof(value_type) * num_items +
@@ -312,13 +316,11 @@ namespace __partition
d_flags = reinterpret_cast<bool *>(
reinterpret_cast<char *>(d_partition_out) + sizeof(value_type) * num_items);
- hip_rocprim::transform(
- policy,
+ hip_rocprim::transform(policy,
stencil,
stencil + num_items,
d_flags,
- predicate
- );
+ predicate);
hip_rocprim::throw_on_error(
rocprim::partition(d_temp_storage,
@@ -340,7 +342,11 @@ namespace __partition
hip_rocprim::parallel_for(
policy,
- partition_fill<SINGLE_OUTPUT, SelectedOutIt, RejectedOutIt, size_type, value_type>(selected_result,
+ partition_fill<SINGLE_OUTPUT,
+ SelectedOutIt,
+ RejectedOutIt,
+ size_type,
+ value_type>(selected_result,
rejected_result,
d_partition_out,
num_selected,

MODIFY
thrust/system/hip/detail/reduce.h
thrust/system/hip/detail/reduce.h
@@ -104,7 +104,7 @@ namespace __reduce {
// Allocate temporary storage.
d_temp_storage = hip_rocprim::get_memory_buffer(policy, sizeof(T) + temp_storage_bytes);
hip_rocprim::throw_on_error(hipGetLastError(),
- "reduce_by_key failed to get memory buffer");
+ "reduce failed to get memory buffer");
d_ret_ptr = reinterpret_cast<T*>(
reinterpret_cast<char *>(d_temp_storage) + temp_storage_bytes);
@@ -144,12 +144,6 @@ reduce_n(execution_policy<Derived> &policy,
T init,
BinaryOp binary_op)
{
- /*THRUST_HIP_PRESERVE_KERNELS_WORKAROUND((
- rocprim::reduce<rocprim::default_config, InputIt, T*, T, BinaryOp>
- ));
- THRUST_HIP_PRESERVE_KERNELS_WORKAROUND((
- hip_rocprim::get_value<Derived, T*>
- ));*/
THRUST_HIP_PRESERVE_KERNELS_WORKAROUND((
__reduce::reduce<Derived, InputIt, Size, T, BinaryOp>
));

MODIFY
thrust/system/hip/detail/reduce_by_key.h
thrust/system/hip/detail/reduce_by_key.h
@@ -119,7 +119,6 @@ namespace __reduce_by_key {
"reduce_by_key failed on 1st step");
// Allocate temporary storage.
- temp_storage_bytes = rocprim::detail::align_size(temp_storage_bytes);
d_temp_storage = hip_rocprim::get_memory_buffer(policy, sizeof(size_type) + temp_storage_bytes);
hip_rocprim::throw_on_error(hipGetLastError(),
"reduce_by_key failed to get memory buffer");

MODIFY
thrust/system/hip/detail/scan.h
thrust/system/hip/detail/scan.h
@@ -108,10 +108,7 @@ inclusive_scan(Policy &policy,
debug_sync),
"scan failed on 1st step");
- hipDeviceSynchronize();
-
// Allocate temporary storage.
- temp_storage_bytes = rocprim::detail::align_size(temp_storage_bytes);
d_temp_storage = hip_rocprim::get_memory_buffer(policy, temp_storage_bytes);
hip_rocprim::throw_on_error(hipGetLastError(),
"scan failed to get memory buffer");
@@ -170,10 +167,7 @@ exclusive_scan(Policy &policy,
debug_sync),
"scan failed on 1st step");
- hipDeviceSynchronize();
-
// Allocate temporary storage.
- temp_storage_bytes = rocprim::detail::align_size(temp_storage_bytes);
d_temp_storage = hip_rocprim::get_memory_buffer(policy, temp_storage_bytes);
hip_rocprim::throw_on_error(hipGetLastError(),
"scan failed to get memory buffer");

MODIFY
thrust/system/hip/detail/scan_by_key.h
thrust/system/hip/detail/scan_by_key.h
@@ -85,7 +85,6 @@ inclusive_scan_by_key(Policy &policy,
"scan_by_key failed on 1st step");
// Allocate temporary storage.
- temp_storage_bytes = rocprim::detail::align_size(temp_storage_bytes);
d_temp_storage = hip_rocprim::get_memory_buffer(policy, temp_storage_bytes);
hip_rocprim::throw_on_error(hipGetLastError(),
"scan_by_key failed to get memory buffer");
@@ -155,7 +154,6 @@ exclusive_scan_by_key(Policy &policy,
"scan_by_key failed on 1st step");
// Allocate temporary storage.
- temp_storage_bytes = rocprim::detail::align_size(temp_storage_bytes);
d_temp_storage = hip_rocprim::get_memory_buffer(policy, temp_storage_bytes);
hip_rocprim::throw_on_error(hipGetLastError(),
"scan_by_key failed to get memory buffer");

MODIFY
thrust/system/hip/detail/unique.h
thrust/system/hip/detail/unique.h
@@ -83,8 +83,9 @@ namespace __unique {
if (num_items == 0)
return items_result;
- hipError_t status;
- status = rocprim::unique(d_temp_storage,
+ // Determine temporary device storage requirements.
+ hip_rocprim::throw_on_error(
+ rocprim::unique(d_temp_storage,
temp_storage_bytes,
items_first,
items_result,
@@ -92,10 +93,10 @@ namespace __unique {
num_items,
binary_pred,
stream,
- debug_sync);
- hip_rocprim::throw_on_error(status, "unique failed on 1st step");
+ debug_sync),
+ "unique failed on 1st step");
- temp_storage_bytes = rocprim::detail::align_size(temp_storage_bytes);
+ // Allocate temporary storage.
d_temp_storage = hip_rocprim::get_memory_buffer(policy, temp_storage_bytes + sizeof(size_type));
hip_rocprim::throw_on_error(hipGetLastError(),
"unique failed to get memory buffer");
@@ -103,7 +104,8 @@ namespace __unique {
d_num_selected_out = reinterpret_cast<size_type *>(
reinterpret_cast<char *>(d_temp_storage) + temp_storage_bytes);
- status = rocprim::unique(d_temp_storage,
+ hip_rocprim::throw_on_error(
+ rocprim::unique(d_temp_storage,
temp_storage_bytes,
items_first,
items_result,
@@ -111,8 +113,8 @@ namespace __unique {
num_items,
binary_pred,
stream,
- debug_sync);
- hip_rocprim::throw_on_error(status, "unique failed on 2nd step");
+ debug_sync),
+ "unique failed on 2nd step");
size_type num_selected = get_value(policy, d_num_selected_out);

MODIFY
thrust/system/hip/detail/unique_by_key.h
thrust/system/hip/detail/unique_by_key.h
@@ -61,9 +61,6 @@ unique_by_key_copy(
namespace hip_rocprim {
-// XXX it should be possible to unify unique & unique_by_key into a single
-// agent with various specializations, similar to what is done
-// with partition
namespace __unique_by_key {
template <class KeyType,
@@ -117,10 +114,9 @@ namespace __unique_by_key {
if (num_items == 0)
return thrust::make_pair(keys_result, values_result);
-
-
- hipError_t status;
- status = rocprim::unique(d_temp_storage,
+ // Determine temporary device storage requirements.
+ hip_rocprim::throw_on_error(
+ rocprim::unique(d_temp_storage,
temp_storage_bytes,
rocprim::make_zip_iterator(rocprim::make_tuple(keys_first, values_first)),
rocprim::make_zip_iterator(rocprim::make_tuple(keys_result, values_result)),
@@ -128,10 +124,10 @@ namespace __unique_by_key {
num_items,
wrapped_binary_pred,
stream,
- debug_sync);
- hip_rocprim::throw_on_error(status, "unique_by_key failed on 1st step");
+ debug_sync),
+ "unique_by_key failed on 1st step");
- temp_storage_bytes = rocprim::detail::align_size(temp_storage_bytes);
+ // Allocate temporary storage.
d_temp_storage = hip_rocprim::get_memory_buffer(policy, temp_storage_bytes + sizeof(size_type));
hip_rocprim::throw_on_error(hipGetLastError(),
"unique_by_key failed to get memory buffer");
@@ -139,7 +135,8 @@ namespace __unique_by_key {
d_num_selected_out = reinterpret_cast<size_type *>(
reinterpret_cast<char *>(d_temp_storage) + temp_storage_bytes);
- status = rocprim::unique(d_temp_storage,
+ hip_rocprim::throw_on_error(
+ rocprim::unique(d_temp_storage,
temp_storage_bytes,
rocprim::make_zip_iterator(rocprim::make_tuple(keys_first, values_first)),
rocprim::make_zip_iterator(rocprim::make_tuple(keys_result, values_result)),
@@ -147,8 +144,8 @@ namespace __unique_by_key {
num_items,
wrapped_binary_pred,
stream,
- debug_sync);
- hip_rocprim::throw_on_error(status, "unique_by_key failed on 2nd step");
+ debug_sync),
+ "unique_by_key failed on 2nd step");
size_type num_selected = get_value(policy, d_num_selected_out);

Refactored the hip implementations