Skip to content

[rocPRIM] Test PR (do not merge) #519

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Closed
umfranzw wants to merge 1 commit into from
Closed

[rocPRIM] Test PR (do not merge) #519

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
15 changes: 15 additions & 0 deletions projects/rocprim/rocprim/include/rocprim/common.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -62,6 +62,21 @@ namespace detail
while(0)
#endif // ROCPRIM_RETURN_ON_ERROR

#ifndef ROCPRIM_RETURN_ON_ERROR2
#define ROCPRIM_RETURN_ON_ERROR2(line, ...) \
do \
{ \
hipError_t error = (__VA_ARGS__); \
if(error != hipSuccess) \
{ \
std::cout << "ROCPRIM_RETURN_ON_ERROR2 encountered error: " << error << std::endl; \
std::cout << "Called from line " << line << std::endl; \
return error; \
} \
} \
while(0)
#endif // ROCPRIM_RETURN_ON_ERROR

} // namespace detail

#endif // ROCPRIM_COMMON_HPP_
32 changes: 16 additions & 16 deletions projects/rocprim/rocprim/include/rocprim/device/device_merge_inplace.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -720,7 +720,7 @@ inline hipError_t merge_inplace(void* temporary_storage,

size_t num_divisions = impl::get_num_global_divisions(left_size, right_size);

ROCPRIM_RETURN_ON_ERROR(detail::temp_storage::partition(
ROCPRIM_RETURN_ON_ERROR2(__LINE__, detail::temp_storage::partition(
temporary_storage,
storage_size,
detail::temp_storage::make_linear_partition(
Expand Down Expand Up @@ -757,12 +757,12 @@ inline hipError_t merge_inplace(void* temporary_storage,
bool is_large_merge = left_size + right_size > impl::block_merge_items_per_block;
typename impl::offset_t iteration = 0;

ROCPRIM_RETURN_ON_ERROR(hipMemcpyAsync(&work_storage[1],
ROCPRIM_RETURN_ON_ERROR2(__LINE__, hipMemcpyAsync(&work_storage[1],
&left_size,
sizeof(decltype(*work_storage)),
hipMemcpyHostToDevice,
stream));
ROCPRIM_RETURN_ON_ERROR(hipMemcpyAsync(&scratch_storage->iteration,
ROCPRIM_RETURN_ON_ERROR2(__LINE__, hipMemcpyAsync(&scratch_storage->iteration,
&iteration,
sizeof(decltype(scratch_storage->iteration)),
hipMemcpyHostToDevice,
Expand All @@ -780,9 +780,9 @@ inline hipError_t merge_inplace(void* temporary_storage,
while(is_large_merge)
{
// clear this flag
ROCPRIM_RETURN_ON_ERROR(hipMemset(&scratch_storage->is_large_merge, false, sizeof(bool)));
ROCPRIM_RETURN_ON_ERROR2(__LINE__, hipMemset(&scratch_storage->is_large_merge, false, sizeof(bool)));

ROCPRIM_RETURN_ON_ERROR(detail::grid_dim_for_max_active_blocks(grid_dim,
ROCPRIM_RETURN_ON_ERROR2(__LINE__, detail::grid_dim_for_max_active_blocks(grid_dim,
global_block_size,
impl::find_pivots_kernel,
stream));
Expand All @@ -793,7 +793,7 @@ inline hipError_t merge_inplace(void* temporary_storage,
work_storage,
pivot_storage,
scratch_storage);
ROCPRIM_RETURN_ON_ERROR(hipGetLastError());
ROCPRIM_RETURN_ON_ERROR2(__LINE__, hipGetLastError());
if(debug_synchronous)
{
ROCPRIM_RETURN_ON_ERROR(hipStreamSynchronize(stream));
Expand All @@ -802,7 +802,7 @@ inline hipError_t merge_inplace(void* temporary_storage,
std::cout << " find_pivots_kernel: " << delta << "ms\n";
}

ROCPRIM_RETURN_ON_ERROR(detail::grid_dim_for_max_active_blocks(grid_dim,
ROCPRIM_RETURN_ON_ERROR2(__LINE__, detail::grid_dim_for_max_active_blocks(grid_dim,
global_block_size,
impl::rotate_p1_kernel,
stream));
Expand All @@ -813,7 +813,7 @@ inline hipError_t merge_inplace(void* temporary_storage,
work_storage,
pivot_storage,
scratch_storage);
ROCPRIM_RETURN_ON_ERROR(hipGetLastError());
ROCPRIM_RETURN_ON_ERROR2(__LINE__, hipGetLastError());
if(debug_synchronous)
{
ROCPRIM_RETURN_ON_ERROR(hipStreamSynchronize(stream));
Expand All @@ -822,7 +822,7 @@ inline hipError_t merge_inplace(void* temporary_storage,
std::cout << " rotate_p1_kernel: " << delta << "ms\n";
}

ROCPRIM_RETURN_ON_ERROR(detail::grid_dim_for_max_active_blocks(grid_dim,
ROCPRIM_RETURN_ON_ERROR2(__LINE__, detail::grid_dim_for_max_active_blocks(grid_dim,
global_block_size,
impl::rotate_p2_kernel,
stream));
Expand All @@ -833,7 +833,7 @@ inline hipError_t merge_inplace(void* temporary_storage,
work_storage,
pivot_storage,
scratch_storage);
ROCPRIM_RETURN_ON_ERROR(hipGetLastError());
ROCPRIM_RETURN_ON_ERROR2(__LINE__, hipGetLastError());
if(debug_synchronous)
{
ROCPRIM_RETURN_ON_ERROR(hipStreamSynchronize(stream));
Expand All @@ -842,7 +842,7 @@ inline hipError_t merge_inplace(void* temporary_storage,
std::cout << " rotate_p2_kernel: " << delta << "ms\n";
}

ROCPRIM_RETURN_ON_ERROR(
ROCPRIM_RETURN_ON_ERROR2(__LINE__,
detail::grid_dim_for_max_active_blocks(grid_dim,
global_block_size,
impl::update_work_tree_kernel,
Expand All @@ -854,20 +854,20 @@ inline hipError_t merge_inplace(void* temporary_storage,
work_storage,
pivot_storage,
scratch_storage);
ROCPRIM_RETURN_ON_ERROR(hipGetLastError());
ROCPRIM_RETURN_ON_ERROR2(__LINE__, hipGetLastError());
if(debug_synchronous)
{
float delta = rocprim::detail::update_time_point(t_time);
std::cout << " update_work_tree_kernel: " << delta << "ms\n";
}

ROCPRIM_RETURN_ON_ERROR(hipMemcpy(&is_large_merge,
ROCPRIM_RETURN_ON_ERROR2(__LINE__, hipMemcpy(&is_large_merge,
&scratch_storage->is_large_merge,
sizeof(bool),
hipMemcpyDeviceToHost));

++iteration;
ROCPRIM_RETURN_ON_ERROR(hipMemcpy(&scratch_storage->iteration,
ROCPRIM_RETURN_ON_ERROR2(__LINE__, hipMemcpy(&scratch_storage->iteration,
&iteration,
sizeof(typename impl::offset_t),
hipMemcpyHostToDevice));
Expand All @@ -881,7 +881,7 @@ inline hipError_t merge_inplace(void* temporary_storage,
// since we potentially may have more blocks than is legally
// allowed on, we simplify launching by doing grid stride per block
int block_merge_grid_size = 0;
ROCPRIM_RETURN_ON_ERROR(detail::grid_dim_for_max_active_blocks(block_merge_grid_size,
ROCPRIM_RETURN_ON_ERROR2(__LINE__, detail::grid_dim_for_max_active_blocks(block_merge_grid_size,
block_block_size,
impl::block_merge_kernel,
stream));
Expand All @@ -903,7 +903,7 @@ inline hipError_t merge_inplace(void* temporary_storage,
compare_function,
work_storage,
scratch_storage);
ROCPRIM_RETURN_ON_ERROR(hipGetLastError());
ROCPRIM_RETURN_ON_ERROR2(__LINE__, hipGetLastError());
if(debug_synchronous)
{
ROCPRIM_RETURN_ON_ERROR(hipStreamSynchronize(stream));
Expand Down
2 changes: 1 addition & 1 deletion projects/rocprim/test/rocprim/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -284,7 +284,7 @@ add_rocprim_test("rocprim.device_adjacent_find" test_device_adjacent_find.cpp)
add_rocprim_test("rocprim.device_find_end" test_device_find_end.cpp)
add_rocprim_test("rocprim.device_histogram" test_device_histogram.cpp)
add_rocprim_test("rocprim.device_merge" test_device_merge.cpp)
add_rocprim_test("rocprim.device_merge_inplace" test_device_merge_inplace.cpp)
add_rocprim_test("rocprim.device_debug_inplace" test_device_debug_inplace.cpp)
add_rocprim_test("rocprim.device_merge_sort" test_device_merge_sort.cpp)
add_rocprim_test("rocprim.nth_element" test_device_nth_element.cpp)
add_rocprim_test("rocprim.device_partial_sort" test_device_partial_sort.cpp)
Expand Down
99 changes: 50 additions & 49 deletions projects/rocprim/test/rocprim/test_device_merge_inplace.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -48,7 +48,7 @@
#include <type_traits>
#include <vector>

TEST(RocprimDeviceMergeInplaceTests, Basic)
TEST(RocprimDeviceDebugInplaceTests, Basic)
{
using value_type = int;
size_t storage_size = 0;
Expand Down Expand Up @@ -238,7 +238,7 @@ template<typename ValueType,
typename Sizes = small_sizes,
typename CompareOp = ::rocprim::less<ValueType>,
bool UseGraphs = false>
struct DeviceMergeInplaceParams
struct DeviceDebugInplaceParams
{
using value_type = ValueType;
using gen_a_type = GenAType;
Expand All @@ -248,64 +248,65 @@ struct DeviceMergeInplaceParams

typedef ::testing::Types<
// linear even-odd
DeviceMergeInplaceParams<int64_t,
linear_data_generator<int64_t, 0, 2>,
linear_data_generator<int64_t, 1, 2>>,
DeviceMergeInplaceParams<int32_t,
linear_data_generator<int32_t, 0, 2>,
linear_data_generator<int32_t, 1, 2>>,
DeviceMergeInplaceParams<int16_t,
linear_data_generator<int16_t, 0, 2>,
linear_data_generator<int16_t, 1, 2>>,
// linear edge cases
DeviceMergeInplaceParams<int32_t,
linear_data_generator<int32_t, 0, 1>,
linear_data_generator<int32_t, 0, 4>>,
DeviceMergeInplaceParams<int32_t,
linear_data_generator<int32_t, 0, 4>,
linear_data_generator<int32_t, 0, 1>>,
DeviceMergeInplaceParams<int32_t,
// DeviceDebugInplaceParams<int64_t,
// linear_data_generator<int64_t, 0, 2>,
// linear_data_generator<int64_t, 1, 2>>,
// DeviceDebugInplaceParams<int32_t,
// linear_data_generator<int32_t, 0, 2>,
// linear_data_generator<int32_t, 1, 2>>,
// DeviceDebugInplaceParams<int16_t,
// linear_data_generator<int16_t, 0, 2>,
// linear_data_generator<int16_t, 1, 2>>,
// // linear edge cases
// DeviceDebugInplaceParams<int32_t,
// linear_data_generator<int32_t, 0, 1>,
// linear_data_generator<int32_t, 0, 4>>,
// DeviceDebugInplaceParams<int32_t,
// linear_data_generator<int32_t, 0, 4>,
// linear_data_generator<int32_t, 0, 1>>,
DeviceDebugInplaceParams<int32_t,
linear_data_generator<int32_t, 128, 0>,
linear_data_generator<int32_t, 0, 1>>,
DeviceMergeInplaceParams<int32_t,
linear_data_generator<int32_t, 0, 1>,
linear_data_generator<int32_t, 128, 0>>,
// random data
DeviceMergeInplaceParams<int64_t,
random_data_generator<int64_t, 2, 2>,
random_data_generator<int64_t, 2, 2>>,
DeviceMergeInplaceParams<int32_t,
random_data_generator<int32_t, 2, 2>,
random_data_generator<int32_t, 2, 2>>,
DeviceMergeInplaceParams<int16_t,
random_data_generator<int16_t, 2, 2>,
random_data_generator<int16_t, 2, 2>>,
DeviceMergeInplaceParams<float,
random_data_generator<float, 2, 2>,
random_data_generator<float, 2, 2>>,
// large input sizes
DeviceMergeInplaceParams<int32_t,
random_data_generator<int32_t, 2, 4>,
random_data_generator<int32_t, 2, 4>,
large_sizes>,
DeviceMergeInplaceParams<int64_t,
random_data_generator<int64_t, 4, 4>,
random_data_generator<int64_t, 4, 4>,
large_sizes>>
DeviceMergeInplaceTestsParams;
linear_data_generator<int32_t, 0, 1>>//,
// DeviceDebugInplaceParams<int32_t,
// linear_data_generator<int32_t, 0, 1>,
// linear_data_generator<int32_t, 128, 0>>,
// // random data
// DeviceDebugInplaceParams<int64_t,
// random_data_generator<int64_t, 2, 2>,
// random_data_generator<int64_t, 2, 2>>,
// DeviceDebugInplaceParams<int32_t,
// random_data_generator<int32_t, 2, 2>,
// random_data_generator<int32_t, 2, 2>>,
// DeviceDebugInplaceParams<int16_t,
// random_data_generator<int16_t, 2, 2>,
// random_data_generator<int16_t, 2, 2>>,
// DeviceDebugInplaceParams<float,
// random_data_generator<float, 2, 2>,
// random_data_generator<float, 2, 2>>,
// // large input sizes
// DeviceDebugInplaceParams<int32_t,
// random_data_generator<int32_t, 2, 4>,
// random_data_generator<int32_t, 2, 4>,
// large_sizes>,
// DeviceDebugInplaceParams<int64_t,
// random_data_generator<int64_t, 4, 4>,
// random_data_generator<int64_t, 4, 4>,
// large_sizes>
>
DeviceDebugInplaceTestsParams;

template<typename Params>
struct DeviceMergeInplaceTests : public testing::Test
struct DeviceDebugInplaceTests : public testing::Test
{
using value_type = typename Params::value_type;
using gen_a_type = typename Params::gen_a_type;
using gen_b_type = typename Params::gen_b_type;
using sizes = typename Params::sizes;
};

TYPED_TEST_SUITE(DeviceMergeInplaceTests, DeviceMergeInplaceTestsParams);
TYPED_TEST_SUITE(DeviceDebugInplaceTests, DeviceDebugInplaceTestsParams);

TYPED_TEST(DeviceMergeInplaceTests, MergeInplace)
TYPED_TEST(DeviceDebugInplaceTests, DebugInplace)
{
using value_type = typename TestFixture::value_type;
using gen_a_type = typename TestFixture::gen_a_type;
Expand Down