Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
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
22 changes: 22 additions & 0 deletions projects/miopen/driver/conv_driver.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -2166,6 +2166,13 @@ int ConvDriver<Tgpu, Tref>::RunForwardGPU()
{
float alpha = static_cast<float>(1), beta = static_cast<float>(0);

const bool perf_log_enabled = miopen::IsPerformanceLoggingEnabled();
miopen::ScopedKernelPhase bias_phase_scope(miopen::KernelPhase::Execution);
if(perf_log_enabled)
{
miopen::AddPerformanceConfig("ConvolutionForwardBias", "");
}

int bias_return_code = CaptureKernel([&]() -> int {
miopenConvolutionForwardBias(GetHandle(),
&alpha,
Expand Down Expand Up @@ -2194,6 +2201,10 @@ int ConvDriver<Tgpu, Tref>::RunForwardGPU()
{
miopenGetKernelTime(GetHandle(), &time);
}
if(perf_log_enabled)
{
miopen::AddInvokerTimes(std::vector<float>{time});
}
printf("GPU Kernel Time Forward Conv. Bias Elapsed: %f ms\n", time);
}

Expand Down Expand Up @@ -2843,6 +2854,13 @@ int ConvDriver<Tgpu, Tref>::RunBackwardGPU()
{
float alpha = static_cast<float>(1), beta = static_cast<float>(0);

const bool perf_log_enabled = miopen::IsPerformanceLoggingEnabled();
miopen::ScopedKernelPhase bias_phase_scope(miopen::KernelPhase::Execution);
if(perf_log_enabled)
{
miopen::AddPerformanceConfig("ConvolutionBackwardBias", "");
}

int bias_return_code = CaptureKernel([&]() -> int {
return miopenConvolutionBackwardBias(GetHandle(),
&alpha,
Expand Down Expand Up @@ -2879,6 +2897,10 @@ int ConvDriver<Tgpu, Tref>::RunBackwardGPU()
{
miopenGetKernelTime(GetHandle(), &time);
}
if(perf_log_enabled)
{
miopen::AddInvokerTimes(std::vector<float>{time});
}
printf("GPU Kernel Time Backward Bias Conv. Elapsed: %f ms\n", time);
}

Expand Down
84 changes: 78 additions & 6 deletions projects/miopen/src/include/miopen/kernel_tuning_mode.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -504,32 +504,65 @@ AddKernelToJsonAccumulator(const std::string& kernel_name, float time_ms, bool i
data.current_config.kernels.push_back({kernel_name, time_ms, is_transform});
}

/// Thread-local override for the solver name reported by LogSolutionName.
///
/// Background: when a TransposingSolver wrapper delegates Search/GetSolution to
/// its inner NHWC-native solver, the inner's own search machinery calls
/// LogSolutionName with the inner's name (e.g. "ConvBinWinogradRxSf2x3").
/// Without an override the perf log credits the inner against the user's
/// original (e.g. NCHW) problem key, even though the actual candidate the
/// runtime would select is the wrapper. Downstream parsers can't then locate
/// the wrapper's perf row.
///
/// When non-empty, LogSolutionName substitutes these values for the ones the
/// caller passed. Use `ScopedSolverNameOverride` to set in an RAII scope.
struct SolverNameOverride
{
std::string solution_name;
uint64_t solver_id = 0;
bool active = false;
};

inline SolverNameOverride& GetSolverNameOverride()
{
thread_local SolverNameOverride override_state;
return override_state;
}

/// Log solution name if appropriate for the current log level
/// Only prints if the solution name has changed since the last call
/// Only prints if the solution name has changed since the last call.
/// Honors any active ScopedSolverNameOverride so wrapper solvers can be
/// credited correctly when their inner solver does the logging.
inline void
LogSolutionName(const std::string& solution_name, uint64_t solver_id, size_t workspace_bytes = 0)
{
const bool logging_enabled = IsLoggingKernel();
if(logging_enabled && !solution_name.empty())
{
// Apply override if active (wrapper solver delegating to inner).
const auto& override_state = GetSolverNameOverride();
const std::string& effective_name =
override_state.active ? override_state.solution_name : solution_name;
const uint64_t effective_id = override_state.active ? override_state.solver_id : solver_id;

auto& last_solution = GetLastPrintedSolutionName();
auto& last_solver_id = GetLastPrintedSolverId();

// Check if solution name or solver_id has changed
if(solution_name != last_solution || solver_id != last_solver_id)
if(effective_name != last_solution || effective_id != last_solver_id)
{
// Flush previous solution's JSON data
FlushJsonAccumulator();

// Set up new solution in accumulator
const KernelPhase current_phase = GetKernelPhase();
auto& data = GetJsonAccumulator();
data.solution_name = solution_name;
data.solver_id = solver_id;
data.solution_name = effective_name;
data.solver_id = effective_id;
data.workspace_bytes = workspace_bytes;
data.phase = KernelPhaseToString(current_phase);
last_solution = solution_name;
last_solver_id = solver_id;
last_solution = effective_name;
last_solver_id = effective_id;
}
}
}
Expand All @@ -553,6 +586,45 @@ class ScopedKernelPhase
KernelPhase prev_phase;
};

/// RAII helper that forces LogSolutionName to attribute its records to a
/// specific solver name + id for the duration of the scope, regardless of what
/// the caller passes. Used by TransposingSolver to credit the wrapper instead
/// of the inner NHWC-native solver when it delegates Search/GetSolution.
/// Restores the previous override (if any) on scope exit, so nesting is safe.
///
/// Gated on MIOPEN_PERFORMANCE_LOGS>0: when performance logging is disabled
/// the constructor and destructor are no-ops, so the thread-local is never
/// touched and the caller pays only an env-var read plus a branch. This keeps
/// the override fully zero-impact for normal production paths.
class ScopedSolverNameOverride
{
public:
ScopedSolverNameOverride(const std::string& wrapper_name, uint64_t wrapper_id)
: engaged(IsPerformanceLoggingEnabled()), prev_state()
{
if(!engaged)
return;
prev_state = GetSolverNameOverride();
auto& s = GetSolverNameOverride();
s.solution_name = wrapper_name;
s.solver_id = wrapper_id;
s.active = !wrapper_name.empty();
}

~ScopedSolverNameOverride()
{
if(engaged)
GetSolverNameOverride() = prev_state;
}

ScopedSolverNameOverride(const ScopedSolverNameOverride&) = delete;
ScopedSolverNameOverride& operator=(const ScopedSolverNameOverride&) = delete;

private:
bool engaged;
SolverNameOverride prev_state;
};

} // namespace miopen

#endif // GUARD_MIOPEN_KERNEL_TUNING_MODE_HPP
Original file line number Diff line number Diff line change
Expand Up @@ -5,7 +5,9 @@

#include <miopen/batched_transpose_sol.hpp>
#include <miopen/datatype.hpp>
#include <miopen/kernel_tuning_mode.hpp>
#include <miopen/op_kernel_args.hpp>
#include <miopen/solver_id.hpp>
#include <miopen/subbuffers.hpp>
#include <miopen/tensor_layout.hpp>

Expand Down Expand Up @@ -806,6 +808,17 @@ struct TransposingSolverGetSolution<Derived, Base, Problem, InvokeParams, Inner,
{
ConvSolution GetSolution(const ExecutionContext& ctx, const Problem& problem) const override
{
// Credit the wrapper (not the inner) in MIOPEN_PERFORMANCE_LOGS so the
// perf-eval DB attribution is correct. The inner's own logging path
// calls LogSolutionName with the inner's name; the override forces the
// wrapper's name onto the emitted JSON record. The override is gated on
// MIOPEN_PERFORMANCE_LOGS>0 inside its constructor, so the empty-string
// path below pays only the env-var read in production builds.
const auto wrapper_name =
IsPerformanceLoggingEnabled() ? this->SolverDbId() : std::string{};
const auto wrapper_id = wrapper_name.empty() ? 0 : miopen::solver::Id(wrapper_name).Value();
ScopedSolverNameOverride name_override(wrapper_name, wrapper_id);

auto transposed_problem = Derived::Transpose(problem);
ConvSolution sln = Inner{}.GetSolution(ctx, transposed_problem);
return static_cast<const Derived*>(this)->WrapSolutionWithTranspose(
Expand Down Expand Up @@ -839,6 +852,16 @@ struct TransposingSolverGetSolution<Derived, Base, Problem, InvokeParams, Inner,
const Problem& problem,
const AnyInvokeParams& invoke_ctx) const override
{
// Force the wrapper's name onto the per-config records emitted by the
// inner solver's GenericSearch loop. Without this the perf log credits
// the inner solver against the original (e.g. NCHW) problem key, which
// a NHWC-only inner cannot actually serve at runtime. Gated on the env
// var: in production builds we pay only the env-var read.
const auto wrapper_name =
IsPerformanceLoggingEnabled() ? this->SolverDbId() : std::string{};
const auto wrapper_id = wrapper_name.empty() ? 0 : miopen::solver::Id(wrapper_name).Value();
ScopedSolverNameOverride name_override(wrapper_name, wrapper_id);

auto transposed_problem = Derived::Transpose(problem);
return Inner{}.Search(ctx, transposed_problem, invoke_ctx);
}
Expand All @@ -847,6 +870,11 @@ struct TransposingSolverGetSolution<Derived, Base, Problem, InvokeParams, Inner,
const Problem& problem,
const PerformanceConfigType& config) const override
{
const auto wrapper_name =
IsPerformanceLoggingEnabled() ? this->SolverDbId() : std::string{};
const auto wrapper_id = wrapper_name.empty() ? 0 : miopen::solver::Id(wrapper_name).Value();
ScopedSolverNameOverride name_override(wrapper_name, wrapper_id);

auto transposed_problem = Derived::Transpose(problem);
ConvSolution sln = Inner{}.GetSolution(ctx, transposed_problem, config);
return static_cast<const Derived*>(this)->WrapSolutionWithTranspose(
Expand Down
Loading