Skip to content

Commit 18880c5

Browse files
NaraendaNaraenda
committed
style(rocprim): update style
1 parent 0c97184 commit 18880c5

15 files changed

+104
-106
lines changed

‎common/utils.hpp‎

Lines changed: 25 additions & 25 deletions
Original file line numberDiff line numberDiff line change
@@ -26,32 +26,32 @@
2626
#include <rocprim/intrinsics/thread.hpp>
2727

2828
#ifndef HIP_CHECK
29-
#ifdef USE_GTEST
30-
// GoogleTest-compatible HIP_CHECK macro. FAIL is called to log the Google Test trace.
31-
// The lambda is invoked immediately as assertions that generate a fatal failure can
32-
// only be used in void-returning functions.
33-
#define HIP_CHECK(condition) \
34-
{ \
35-
hipError_t error = condition; \
36-
if(error != hipSuccess) \
37-
{ \
38-
[error]() \
39-
{ FAIL() << "HIP error " << error << ": " << hipGetErrorString(error); }(); \
40-
exit(error); \
41-
} \
42-
}
43-
#else
44-
#define HIP_CHECK(condition) \
45-
{ \
46-
hipError_t error = condition; \
47-
if(error != hipSuccess) \
29+
#ifdef USE_GTEST
30+
// GoogleTest-compatible HIP_CHECK macro. FAIL is called to log the Google Test trace.
31+
// The lambda is invoked immediately as assertions that generate a fatal failure can
32+
// only be used in void-returning functions.
33+
#define HIP_CHECK(condition) \
4834
{ \
49-
std::cout << "HIP error: " << hipGetErrorString(error) << " file: " << __FILE__ \
50-
<< " line: " << __LINE__ << std::endl; \
51-
exit(error); \
52-
} \
53-
}
54-
#endif
35+
hipError_t error = condition; \
36+
if(error != hipSuccess) \
37+
{ \
38+
[error]() \
39+
{ FAIL() << "HIP error " << error << ": " << hipGetErrorString(error); }(); \
40+
exit(error); \
41+
} \
42+
}
43+
#else
44+
#define HIP_CHECK(condition) \
45+
{ \
46+
hipError_t error = condition; \
47+
if(error != hipSuccess) \
48+
{ \
49+
std::cout << "HIP error: " << hipGetErrorString(error) \
50+
<< " file: " << __FILE__ << " line: " << __LINE__ << std::endl; \
51+
exit(error); \
52+
} \
53+
}
54+
#endif
5555
#endif
5656

5757
namespace common

‎rocprim/include/rocprim/detail/various.hpp‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -453,15 +453,15 @@ struct lookback_variant_util
453453
{
454454
using T = std::integral_constant<bool, true>;
455455
using F = std::integral_constant<bool, false>;
456-
if (use_sleep)
456+
if(use_sleep)
457457
{
458-
if (use_atomic)
458+
if(use_atomic)
459459
{
460460
return f(T{}, T{});
461461
}
462462
return f(T{}, F{});
463463
}
464-
if (use_atomic)
464+
if(use_atomic)
465465
{
466466
return f(F{}, T{});
467467
}

‎rocprim/include/rocprim/device/detail/device_batch_memcpy.hpp‎

Lines changed: 3 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -909,7 +909,7 @@ struct batch_memcpy_impl
909909
static void init_tile_state_kernel(blev_buffer_scan_state_type buffer_scan_state,
910910
blev_block_scan_state_type block_scan_state,
911911
tile_offset_type num_tiles,
912-
ordered_bid_type ordered_bid)
912+
ordered_bid_type ordered_bid)
913913
{
914914
const uint32_t block_id = rocprim::detail::block_id<0>();
915915
const uint32_t block_size = rocprim::detail::block_size<0>();
@@ -1062,8 +1062,7 @@ ROCPRIM_INLINE static hipError_t batch_memcpy_func(void* temporary_
10621062
ROCPRIM_RETURN_ON_ERROR(lookback_variant_util(false, use_atomic_block_id)(
10631063
[&](auto /* use_sleepy_scan */, auto use_atomic_block_id)
10641064
{
1065-
using block_id_type
1066-
= detail::block_id_wrapper<uint32_t, use_atomic_block_id>;
1065+
using block_id_type = detail::block_id_wrapper<uint32_t, use_atomic_block_id>;
10671066

10681067
using batch_memcpy_impl_type = detail::batch_memcpy_impl<Config,
10691068
IsMemCpy,
@@ -1110,7 +1109,7 @@ ROCPRIM_INLINE static hipError_t batch_memcpy_func(void* temporary_
11101109
// The non-blev kernel will prepare blev copy. Communication between the two
11111110
// kernels is done via `blev_buffers`.
11121111
typename batch_memcpy_impl_type::copyable_blev_buffers blev_buffers{};
1113-
1112+
11141113
typename block_id_type::id_type* block_id_pool = nullptr;
11151114

11161115
// Partition `d_temp_storage`.

‎rocprim/include/rocprim/device/detail/device_nth_element.hpp‎

Lines changed: 2 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -436,7 +436,7 @@ ROCPRIM_DEVICE ROCPRIM_FORCE_INLINE void
436436
typename WrappedBlockId::storage_type ordered_bid;
437437
} storage;
438438

439-
auto block_id = ordered_bid.get(threadIdx.x, storage.ordered_bid);
439+
auto block_id = ordered_bid.get(threadIdx.x, storage.ordered_bid);
440440
uint8_t buckets[num_items_per_thread];
441441

442442
const size_t nth_element = nth_element_data->bucket_idx;
@@ -593,8 +593,7 @@ template<class config,
593593
class KeysIterator,
594594
class BinaryFunction,
595595
class WrappedBlockId>
596-
ROCPRIM_KERNEL
597-
__launch_bounds__(device_params<config>().kernel_config.block_size)
596+
ROCPRIM_KERNEL __launch_bounds__(device_params<config>().kernel_config.block_size)
598597
void kernel_copy_buckets(KeysIterator keys,
599598
typename std::iterator_traits<KeysIterator>::value_type* tree,
600599
const size_t size,

‎rocprim/include/rocprim/device/detail/device_radix_sort.hpp‎

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -602,7 +602,8 @@ ROCPRIM_DEVICE ROCPRIM_FORCE_INLINE void sort_single(KeysInputIterator keys_i
602602
using sort_single_helper
603603
= radix_sort_single_helper<BlockSize, ItemsPerThread, Descending, key_type, value_type>;
604604

605-
ROCPRIM_SHARED_MEMORY typename sort_single_helper::storage_type storage;
605+
ROCPRIM_SHARED_MEMORY
606+
typename sort_single_helper::storage_type storage;
606607

607608
sort_single_helper().template sort_single<>(keys_input,
608609
keys_output,
@@ -1385,7 +1386,8 @@ ROCPRIM_DEVICE ROCPRIM_FORCE_INLINE void
13851386
Decomposer,
13861387
BlockIdWrapper>;
13871388

1388-
ROCPRIM_SHARED_MEMORY typename onesweep_iteration_helper_type::storage_type storage;
1389+
ROCPRIM_SHARED_MEMORY
1390+
typename onesweep_iteration_helper_type::storage_type storage;
13891391

13901392
constexpr unsigned int items_per_block = BlockSize * ItemsPerThread;
13911393
const unsigned int thread_id = ::rocprim::detail::block_thread_id<0>();

‎rocprim/include/rocprim/device/detail/device_run_length_encode.hpp‎

Lines changed: 11 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -819,7 +819,8 @@ ROCPRIM_DEVICE ROCPRIM_FORCE_INLINE auto
819819
load_input_method,
820820
scan_algorithm>;
821821

822-
ROCPRIM_SHARED_MEMORY union {
822+
ROCPRIM_SHARED_MEMORY union
823+
{
823824
typename block_processor::storage_type_ block_processor_storage;
824825
typename WrappedBlockId::storage_type ordered_bid_storage;
825826
} storage;
@@ -845,14 +846,15 @@ ROCPRIM_DEVICE ROCPRIM_FORCE_INLINE auto
845846
}
846847
else if(valid_in_last_block > 0)
847848
{
848-
OffsetCountPairType total = block_processor{}.process_block(block_input,
849-
offsets_output,
850-
counts_output,
851-
scan_state,
852-
block_id,
853-
grid_size,
854-
size,
855-
storage.block_processor_storage);
849+
OffsetCountPairType total
850+
= block_processor{}.process_block(block_input,
851+
offsets_output,
852+
counts_output,
853+
scan_state,
854+
block_id,
855+
grid_size,
856+
size,
857+
storage.block_processor_storage);
856858
// First thread of last block sets the total number of non-trivial runs found and updates
857859
// the counts with the last run's length if necessary.
858860
if(threadIdx.x == 0)

‎rocprim/include/rocprim/device/detail/device_scan_by_key.hpp‎

Lines changed: 1 addition & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -336,9 +336,7 @@ namespace detail
336336
} storage;
337337

338338
const auto flat_thread_id = ::rocprim::detail::block_thread_id<0>();
339-
const auto flat_block_id
340-
= ordered_bid.get(flat_thread_id,
341-
storage.ordered_bid);
339+
const auto flat_block_id = ordered_bid.get(flat_thread_id, storage.ordered_bid);
342340

343341
// Load input
344342
wrapped_type wrapped_values[items_per_thread];

‎rocprim/include/rocprim/device/device_histogram.hpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -21,11 +21,11 @@
2121
#ifndef ROCPRIM_DEVICE_DEVICE_HISTOGRAM_HPP_
2222
#define ROCPRIM_DEVICE_DEVICE_HISTOGRAM_HPP_
2323

24+
#include <chrono>
2425
#include <cmath>
2526
#include <iostream>
2627
#include <iterator>
2728
#include <type_traits>
28-
#include <chrono>
2929

3030
#include "../config.hpp"
3131
#include "../common.hpp"

‎rocprim/include/rocprim/device/device_partition.hpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -22,9 +22,9 @@
2222
#define ROCPRIM_DEVICE_DEVICE_PARTITION_HPP_
2323

2424
#include <algorithm>
25+
#include <chrono>
2526
#include <iostream>
2627
#include <iterator>
27-
#include <chrono>
2828

2929
#include "../config.hpp"
3030
#include "../common.hpp"

‎rocprim/include/rocprim/device/device_radix_sort.hpp‎

Lines changed: 24 additions & 26 deletions
Original file line numberDiff line numberDiff line change
@@ -21,11 +21,11 @@
2121
#ifndef ROCPRIM_DEVICE_DEVICE_RADIX_SORT_HPP_
2222
#define ROCPRIM_DEVICE_DEVICE_RADIX_SORT_HPP_
2323

24+
#include <chrono>
2425
#include <iostream>
2526
#include <iterator>
2627
#include <type_traits>
2728
#include <utility>
28-
#include <chrono>
2929

3030
#include "../config.hpp"
3131
#include "../common.hpp"
@@ -211,21 +211,20 @@ template<class Config,
211211
class Offset,
212212
class Decomposer,
213213
class BlockIdWrapper>
214-
ROCPRIM_KERNEL
215-
__launch_bounds__(device_params<Config>().sort.block_size) void onesweep_iteration_kernel(
216-
KeysInputIterator keys_input,
217-
KeysOutputIterator keys_output,
218-
ValuesInputIterator values_input,
219-
ValuesOutputIterator values_output,
220-
const unsigned int size,
221-
Offset* global_digit_offsets_in,
222-
Offset* global_digit_offsets_out,
223-
onesweep_lookback_state* lookback_states,
224-
Decomposer decomposer,
225-
const unsigned int bit,
226-
const unsigned int current_radix_bits,
227-
const unsigned int full_blocks,
228-
BlockIdWrapper block_id)
214+
ROCPRIM_KERNEL __launch_bounds__(device_params<Config>().sort.block_size)
215+
void onesweep_iteration_kernel(KeysInputIterator keys_input,
216+
KeysOutputIterator keys_output,
217+
ValuesInputIterator values_input,
218+
ValuesOutputIterator values_output,
219+
const unsigned int size,
220+
Offset* global_digit_offsets_in,
221+
Offset* global_digit_offsets_out,
222+
onesweep_lookback_state* lookback_states,
223+
Decomposer decomposer,
224+
const unsigned int bit,
225+
const unsigned int current_radix_bits,
226+
const unsigned int full_blocks,
227+
BlockIdWrapper block_id)
229228
{
230229
static constexpr radix_sort_onesweep_config_params params = device_params<Config>();
231230
onesweep_iteration<params.sort.block_size,
@@ -468,8 +467,7 @@ hipError_t radix_sort_onesweep_impl(
468467
ROCPRIM_RETURN_ON_ERROR(lookback_variant_util(false, use_atomic_block_id)(
469468
[&](auto /* use_sleepy_scan */, auto use_atomic_block_id)
470469
{
471-
using block_id_type
472-
= detail::block_id_wrapper<uint32_t, use_atomic_block_id>;
470+
using block_id_type = detail::block_id_wrapper<uint32_t, use_atomic_block_id>;
473471

474472
using config = wrapped_radix_sort_onesweep_config<Config, key_type, value_type>;
475473

@@ -500,11 +498,11 @@ hipError_t radix_sort_onesweep_impl(
500498
constexpr bool with_values = !std::is_same<value_type, ::rocprim::empty_type>::value;
501499
const bool with_double_buffer = keys_tmp != nullptr;
502500

503-
offset_type* global_digit_offsets;
504-
offset_type* global_digit_offsets_tmp;
505-
onesweep_lookback_state* lookback_states;
506-
key_type* keys_tmp_storage;
507-
value_type* values_tmp_storage;
501+
offset_type* global_digit_offsets;
502+
offset_type* global_digit_offsets_tmp;
503+
onesweep_lookback_state* lookback_states;
504+
key_type* keys_tmp_storage;
505+
value_type* values_tmp_storage;
508506
typename block_id_type::id_type* block_id_storage;
509507

510508
const hipError_t partition_result = detail::temp_storage::partition(
@@ -520,9 +518,9 @@ hipError_t radix_sort_onesweep_impl(
520518
detail::temp_storage::ptr_aligned_array(
521519
&values_tmp_storage,
522520
!with_double_buffer && with_values ? size : 0),
523-
detail::temp_storage::make_partition(
524-
&block_id_storage,
525-
block_id_type::get_temp_storage_layout())));
521+
detail::temp_storage::make_partition(
522+
&block_id_storage,
523+
block_id_type::get_temp_storage_layout())));
526524

527525
if(partition_result != hipSuccess || temporary_storage == nullptr)
528526
{

0 commit comments

Comments
 (0)