Skip to content

Commit

Permalink
Fix to prevent recording events when handle doesnt have profiling ena…
Browse files Browse the repository at this point in the history
…bled (#3232)
  • Loading branch information
BrianHarrisonAMD authored Sep 4, 2024
1 parent 0cf61d6 commit e275d9d
Showing 1 changed file with 17 additions and 9 deletions.
26 changes: 17 additions & 9 deletions src/include/miopen/solver/implicitgemm_ck_util.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -109,22 +109,30 @@ bool IsCKApplicable(const ProblemDescriptionType& problem)
struct HipEventProfiler
{
const Handle& handle;
float event_time;
HipEventPtr start;
HipEventPtr stop;

HipEventProfiler(const Handle& handle_)
: handle(handle_), event_time(0.0f), start(make_hip_event()), stop(make_hip_event())
HipEventProfiler(const Handle& handle_) : handle(handle_), start(nullptr), stop(nullptr)
{
hipEventRecord(start.get(), handle.GetStream());
if(handle.IsProfilingEnabled())
{
start = make_hip_event();
stop = make_hip_event();
hipEventRecord(start.get(), handle.GetStream());
}
}

~HipEventProfiler()
{
hipEventRecord(stop.get(), handle.GetStream());
hipEventSynchronize(stop.get());
hipEventElapsedTime(&event_time, start.get(), stop.get());
handle.ResetKernelTime();
handle.AccumKernelTime(event_time);
if(start)
{
hipEventRecord(stop.get(), handle.GetStream());
hipEventSynchronize(stop.get());
float event_time = 0.0f;
hipEventElapsedTime(&event_time, start.get(), stop.get());
handle.ResetKernelTime();
handle.AccumKernelTime(event_time);
}
}
};
#endif
Expand Down

0 comments on commit e275d9d

Please sign in to comment.