Skip to content

Add hipGetLastError calls to clear existing errors #686

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
wants to merge 5 commits into from
Closed

Add hipGetLastError calls to clear existing errors #686

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
341b93d
Test and benchmark updates for new hipGetLastError behaviour
umfranzw Feb 21, 2025
4fafa0a
Clear HIP error in public device API functions
umfranzw Feb 28, 2025
a6e932e
Improve accuracy of HIP_CHECK, HIP_CHECK_LAUNCH macros
umfranzw Mar 4, 2025
0fcf7fa
Add async version of HIP_CHECK_LAUNCH
umfranzw Mar 7, 2025
94719e3
Remove duplicate lines
umfranzw May 16, 2025
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
42 changes: 23 additions & 19 deletions benchmark/benchmark_block_adjacent_difference.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -282,15 +282,17 @@ auto run_benchmark(benchmark::State& state,
// Record start event
HIP_CHECK(hipEventRecord(start, stream));

hipLaunchKernelGGL(HIP_KERNEL_NAME(kernel<Benchmark, BlockSize, ItemsPerThread, WithTile>),
dim3(num_blocks),
dim3(BlockSize),
0,
stream,
d_input,
d_output,
Trials);
HIP_CHECK(hipGetLastError());
HIP_CHECK_LAUNCH(
hipLaunchKernelGGL(HIP_KERNEL_NAME(kernel<Benchmark, BlockSize, ItemsPerThread, WithTile>),
dim3(num_blocks),
dim3(BlockSize),
0,
stream,
d_input,
d_output,
Trials
)
);

// Record stop event and wait until it completes
HIP_CHECK(hipEventRecord(stop, stream));
Expand Down Expand Up @@ -368,16 +370,18 @@ auto run_benchmark(benchmark::State& state,
// Record start event
HIP_CHECK(hipEventRecord(start, stream));

hipLaunchKernelGGL(HIP_KERNEL_NAME(kernel<Benchmark, BlockSize, ItemsPerThread, WithTile>),
dim3(num_blocks),
dim3(BlockSize),
0,
stream,
d_input,
d_tile_sizes,
d_output,
Trials);
HIP_CHECK(hipGetLastError());
HIP_CHECK_LAUNCH(
hipLaunchKernelGGL(HIP_KERNEL_NAME(kernel<Benchmark, BlockSize, ItemsPerThread, WithTile>),
dim3(num_blocks),
dim3(BlockSize),
0,
stream,
d_input,
d_tile_sizes,
d_output,
Trials
)
);

// Record stop event and wait until it completes
HIP_CHECK(hipEventRecord(stop, stream));
Expand Down
20 changes: 11 additions & 9 deletions benchmark/benchmark_block_discontinuity.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -230,15 +230,17 @@ void run_benchmark(benchmark::State& state,
// Record start event
HIP_CHECK(hipEventRecord(start, stream));

hipLaunchKernelGGL(
HIP_KERNEL_NAME(kernel<Benchmark, T, BlockSize, ItemsPerThread, WithTile, Trials>),
dim3(size / items_per_block),
dim3(BlockSize),
0,
stream,
d_input,
d_output);
HIP_CHECK(hipGetLastError());
HIP_CHECK_LAUNCH(
hipLaunchKernelGGL(
HIP_KERNEL_NAME(kernel<Benchmark, T, BlockSize, ItemsPerThread, WithTile, Trials>),
dim3(size / items_per_block),
dim3(BlockSize),
0,
stream,
d_input,
d_output
)
);

// Record stop event and wait until it completes
HIP_CHECK(hipEventRecord(stop, stream));
Expand Down
20 changes: 11 additions & 9 deletions benchmark/benchmark_block_exchange.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -261,15 +261,17 @@ void run_benchmark(benchmark::State& state,
// Record start event
HIP_CHECK(hipEventRecord(start, stream));

hipLaunchKernelGGL(HIP_KERNEL_NAME(kernel<Benchmark, T, BlockSize, ItemsPerThread, Trials>),
dim3(size / items_per_block),
dim3(BlockSize),
0,
stream,
d_input,
d_ranks,
d_output);
HIP_CHECK(hipGetLastError());
HIP_CHECK_LAUNCH(
hipLaunchKernelGGL(HIP_KERNEL_NAME(kernel<Benchmark, T, BlockSize, ItemsPerThread, Trials>),
dim3(size / items_per_block),
dim3(BlockSize),
0,
stream,
d_input,
d_ranks,
d_output
)
);

// Record stop event and wait until it completes
HIP_CHECK(hipEventRecord(stop, stream));
Expand Down
20 changes: 11 additions & 9 deletions benchmark/benchmark_block_histogram.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -137,15 +137,17 @@ void run_benchmark(benchmark::State& state, hipStream_t stream, size_t bytes)
// Record start event
HIP_CHECK(hipEventRecord(start, stream));

hipLaunchKernelGGL(
HIP_KERNEL_NAME(kernel<Benchmark, T, BlockSize, ItemsPerThread, BinSize, Trials>),
dim3(size / items_per_block),
dim3(BlockSize),
0,
stream,
d_input,
d_output);
HIP_CHECK(hipGetLastError());
HIP_CHECK_LAUNCH(
hipLaunchKernelGGL(
HIP_KERNEL_NAME(kernel<Benchmark, T, BlockSize, ItemsPerThread, BinSize, Trials>),
dim3(size / items_per_block),
dim3(BlockSize),
0,
stream,
d_input,
d_output
)
);

// Record stop event and wait until it completes
HIP_CHECK(hipEventRecord(stop, stream));
Expand Down
42 changes: 23 additions & 19 deletions benchmark/benchmark_block_radix_sort.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -166,29 +166,33 @@ void run_benchmark(benchmark::State& state,

if(benchmark_kind == benchmark_kinds::sort_keys)
{
hipLaunchKernelGGL(
HIP_KERNEL_NAME(
sort_keys_kernel<T, BlockSize, RadixBitsPerPass, ItemsPerThread, Trials>),
dim3(size / items_per_block),
dim3(BlockSize),
0,
stream,
d_input,
d_output);
HIP_CHECK_LAUNCH(
hipLaunchKernelGGL(
HIP_KERNEL_NAME(
sort_keys_kernel<T, BlockSize, RadixBitsPerPass, ItemsPerThread, Trials>),
dim3(size / items_per_block),
dim3(BlockSize),
0,
stream,
d_input,
d_output
)
);
}
else if(benchmark_kind == benchmark_kinds::sort_pairs)
{
hipLaunchKernelGGL(
HIP_KERNEL_NAME(
sort_pairs_kernel<T, BlockSize, RadixBitsPerPass, ItemsPerThread, Trials>),
dim3(size / items_per_block),
dim3(BlockSize),
0,
stream,
d_input,
d_output);
HIP_CHECK_LAUNCH(
hipLaunchKernelGGL(
HIP_KERNEL_NAME(
sort_pairs_kernel<T, BlockSize, RadixBitsPerPass, ItemsPerThread, Trials>),
dim3(size / items_per_block),
dim3(BlockSize),
0,
stream,
d_input,
d_output)
);
}
HIP_CHECK(hipGetLastError());

// Record stop event and wait until it completes
HIP_CHECK(hipEventRecord(stop, stream));
Expand Down
18 changes: 10 additions & 8 deletions benchmark/benchmark_block_reduce.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -121,14 +121,16 @@ void run_benchmark(benchmark::State& state, hipStream_t stream, size_t bytes)
// Record start event
HIP_CHECK(hipEventRecord(start, stream));

hipLaunchKernelGGL(HIP_KERNEL_NAME(kernel<Benchmark, T, BlockSize, ItemsPerThread, Trials>),
dim3(size / items_per_block),
dim3(BlockSize),
0,
stream,
d_input,
d_output);
HIP_CHECK(hipGetLastError());
HIP_CHECK_LAUNCH(
hipLaunchKernelGGL(HIP_KERNEL_NAME(kernel<Benchmark, T, BlockSize, ItemsPerThread, Trials>),
dim3(size / items_per_block),
dim3(BlockSize),
0,
stream,
d_input,
d_output
)
);

// Record stop event and wait until it completes
HIP_CHECK(hipEventRecord(stop, stream));
Expand Down
18 changes: 10 additions & 8 deletions benchmark/benchmark_block_scan.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -154,14 +154,16 @@ void run_benchmark(benchmark::State& state, hipStream_t stream, size_t bytes)
// Record start event
HIP_CHECK(hipEventRecord(start, stream));

hipLaunchKernelGGL(HIP_KERNEL_NAME(kernel<Benchmark, T, BlockSize, ItemsPerThread, Trials>),
dim3(size / items_per_block),
dim3(BlockSize),
0,
stream,
d_input,
d_output);
HIP_CHECK(hipGetLastError());
HIP_CHECK_LAUNCH(
hipLaunchKernelGGL(HIP_KERNEL_NAME(kernel<Benchmark, T, BlockSize, ItemsPerThread, Trials>),
dim3(size / items_per_block),
dim3(BlockSize),
0,
stream,
d_input,
d_output
)
);

// Record stop event and wait until it completes
HIP_CHECK(hipEventRecord(stop, stream));
Expand Down
3 changes: 1 addition & 2 deletions benchmark/benchmark_block_sort.parallel.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -273,9 +273,8 @@ struct block_sort_benchmark : public config_autotune_interface

for(size_t i = 0; i < batch_size; ++i)
{
dispatch_block_sort(stable_tag, size, stream, d_input, d_output);
HIP_CHECK_LAUNCH(dispatch_block_sort(stable_tag, size, stream, d_input, d_output));
}
HIP_CHECK(hipGetLastError());

// Record stop event and wait until it completes
HIP_CHECK(hipEventRecord(stop, stream));
Expand Down
4 changes: 1 addition & 3 deletions benchmark/benchmark_config_dispatch.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -63,10 +63,8 @@ static void BM_kernel_launch(benchmark::State& state)

for(auto _ : state)
{
hipLaunchKernelGGL(empty_kernel, dim3(1), dim3(1), 0, stream);
HIP_CHECK(hipGetLastError());
HIP_CHECK_LAUNCH_SYNC(hipLaunchKernelGGL(empty_kernel, dim3(1), dim3(1), 0, stream));
}
HIP_CHECK(hipStreamSynchronize(stream));
}

#define CREATE_BENCHMARK(ST, SK) \
Expand Down
3 changes: 0 additions & 3 deletions benchmark/benchmark_device_adjacent_find.parallel.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -188,9 +188,6 @@ struct device_adjacent_find_benchmark : public config_autotune_interface
HIP_CHECK(hipEventRecord(stop, stream));
HIP_CHECK(hipEventSynchronize(stop));

HIP_CHECK(hipGetLastError());
HIP_CHECK(hipDeviceSynchronize());

float elapsed_mseconds;
HIP_CHECK(hipEventElapsedTime(&elapsed_mseconds, start, stop));
state.SetIterationTime(elapsed_mseconds / 1000);
Expand Down
1 change: 0 additions & 1 deletion benchmark/benchmark_utils.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -50,7 +50,6 @@
#include <memory>
#include <random>
#include <regex>
#include <sstream>
#include <stdexcept>
#include <stdint.h>
#include <string>
Expand Down
38 changes: 21 additions & 17 deletions benchmark/benchmark_warp_reduce.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -95,14 +95,16 @@ inline auto execute_warp_reduce_kernel(
T* input, T* output, Flag* /* flags */, size_t size, hipStream_t stream) ->
typename std::enable_if<!Segmented>::type
{
hipLaunchKernelGGL(HIP_KERNEL_NAME(warp_reduce_kernel<AllReduce, T, WarpSize, Trials>),
dim3(size / BlockSize),
dim3(BlockSize),
0,
stream,
input,
output);
HIP_CHECK(hipGetLastError());
HIP_CHECK_LAUNCH(
hipLaunchKernelGGL(HIP_KERNEL_NAME(warp_reduce_kernel<AllReduce, T, WarpSize, Trials>),
dim3(size / BlockSize),
dim3(BlockSize),
0,
stream,
input,
output
)
);
}

template<bool AllReduce,
Expand All @@ -116,15 +118,17 @@ inline auto
execute_warp_reduce_kernel(T* input, T* output, Flag* flags, size_t size, hipStream_t stream) ->
typename std::enable_if<Segmented>::type
{
hipLaunchKernelGGL(HIP_KERNEL_NAME(segmented_warp_reduce_kernel<T, Flag, WarpSize, Trials>),
dim3(size / BlockSize),
dim3(BlockSize),
0,
stream,
input,
flags,
output);
HIP_CHECK(hipGetLastError());
HIP_CHECK_LAUNCH(
hipLaunchKernelGGL(HIP_KERNEL_NAME(segmented_warp_reduce_kernel<T, Flag, WarpSize, Trials>),
dim3(size / BlockSize),
dim3(BlockSize),
0,
stream,
input,
flags,
output
)
);
}

template<bool AllReduce,
Expand Down
45 changes: 25 additions & 20 deletions benchmark/benchmark_warp_sort.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -151,35 +151,40 @@ void run_benchmark(benchmark::State& state,
ROCPRIM_NO_UNROLL
for(unsigned int trial = 0; trial < Trials; ++trial)
{
hipLaunchKernelGGL(
HIP_KERNEL_NAME(
warp_sort_by_key_kernel<Key, Value, BlockSize, WarpSize, ItemsPerThread>),
dim3(size / items_per_block),
dim3(BlockSize),
0,
stream,
d_input_key,
d_input_value,
d_output_key,
d_output_value);
HIP_CHECK_LAUNCH(
hipLaunchKernelGGL(
HIP_KERNEL_NAME(
warp_sort_by_key_kernel<Key, Value, BlockSize, WarpSize, ItemsPerThread>),
dim3(size / items_per_block),
dim3(BlockSize),
0,
stream,
d_input_key,
d_input_value,
d_output_key,
d_output_value
)
);
}
}
else
{
ROCPRIM_NO_UNROLL
for(unsigned int trial = 0; trial < Trials; ++trial)
{
hipLaunchKernelGGL(
HIP_KERNEL_NAME(warp_sort_kernel<Key, BlockSize, WarpSize, ItemsPerThread>),
dim3(size / items_per_block),
dim3(BlockSize),
0,
stream,
d_input_key,
d_output_key);
HIP_CHECK_LAUNCH(
hipLaunchKernelGGL(
HIP_KERNEL_NAME(warp_sort_kernel<Key, BlockSize, WarpSize, ItemsPerThread>),
dim3(size / items_per_block),
dim3(BlockSize),
0,
stream,
d_input_key,
d_output_key
)
);
}
}
HIP_CHECK(hipGetLastError());

// Record stop event and wait until it completes
HIP_CHECK(hipEventRecord(stop, stream));
Expand Down
Loading