Skip to content
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

Fix tuning timing results #3460

Merged
merged 6 commits into from
Jan 11, 2025
+64 −13
Merged

Fix tuning timing results #3460

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
a183ef3
Fix solvers missing from results when tuning
BrianHarrisonAMD Dec 19, 2024
c4f1d5a
Fix CK solvers capturing kernel time differently than other solvers w…
BrianHarrisonAMD Dec 19, 2024
bc74eac
Fix formatting
BrianHarrisonAMD Dec 19, 2024
99f21ea
Merge branch 'develop' into bharriso/fix-tuning-timing-results
BrianHarrisonAMD Jan 8, 2025
5ee3129
Merge branch 'develop' into bharriso/fix-tuning-timing-results
BrianHarrisonAMD Jan 11, 2025
020fb2a
Rename profiler to avoid confusion
BrianHarrisonAMD Jan 11, 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
2 changes: 1 addition & 1 deletion src/conv/solver_finders.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -298,7 +298,7 @@ static std::vector<Solution> EvaluateInvokers(const Handle& handle,
best_invoker = invoker;
}

auto solution = Solution{solver::Id{selected.solver_id}, best, selected.workspace_sz};
auto solution = Solution{solver::Id{sol.solver_id}, elapsed, sol.workspace_sz};
if(force_attach_binary)
solution.SetInvoker(invoker, programs, selected.construction_params);
else
Expand Down
75 changes: 63 additions & 12 deletions src/include/miopen/solver/implicitgemm_ck_util.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -485,6 +485,8 @@ class TransposeInstance

void ZeroOutBuffer(const Handle& handle)
{
HipEventProfiler pfr(handle);

[[maybe_unused]] auto status =
hipMemsetAsync(buf_handle.get(), 0, tensor_sz, handle.GetStream());
assert(status == hipSuccess);
Expand All @@ -508,7 +510,7 @@ class TransposeInstance
kern_args[0] = out_ptr;
kern_args[1] = in_ptr;

auto save = handle.IsProfilingEnabled() ? 0.0f : handle.GetKernelTime();
auto save = handle.IsProfilingEnabled() ? handle.GetKernelTime() : 0.0f;
handle.Run(kernels[kern_idx])(kern_args);
if(handle.IsProfilingEnabled())
{
Expand Down Expand Up @@ -749,6 +751,8 @@ ZeroOutTensor(const Handle& handle, const TensorDescriptor& tensorDesc, Data_t t
// Use faster clear if possible.
if(tensorDesc.IsPacked())
{
HipEventProfiler pfr(handle);

auto status = hipMemsetAsync(tensorData, 0, tensorDesc.GetNumBytes(), handle.GetStream());
if(status != hipSuccess)
{
Expand Down Expand Up @@ -855,16 +859,20 @@ ConvSolution InitInvokerFactoryNCHW(const ExecutionContext& ctx,
std::swap(conv_tensors.x, conv_tensors.y);
std::swap(conv_tensors.xDesc, conv_tensors.yDesc);
}
WorkAroundHipEventProfiler pfr(handle);
input1_tr_inst.ConvertFrom(handle, kernels, conv_tensors);

input2_tr_inst.ConvertFrom(handle, kernels, conv_tensors);
float elapsed = 0.0f;

// ConvertFrom automatically keeps kernel time and accumulates
input1_tr_inst.ConvertFrom(handle, kernels, conv_tensors);
input2_tr_inst.ConvertFrom(handle, kernels, conv_tensors);
output_init_tr_inst.ConvertFrom(handle, kernels, conv_tensors);

/// \todo: Will need SetTensor() to properly zero out non-packed tensors
/// Note: Need to clear buffer memory for output since all values may not be set.
elapsed = handle.IsProfilingEnabled() ? handle.GetKernelTime() : 0.0f;
output_tr_inst.ZeroOutBuffer(handle);
if(handle.IsProfilingEnabled())
elapsed += handle.GetKernelTime();

std::array<internal::TransposeInstanceTagged*, 3> tr_ptrs = {
&input1_tr_inst, &input2_tr_inst, &output_tr_inst};
Expand All @@ -874,7 +882,6 @@ ConvSolution InitInvokerFactoryNCHW(const ExecutionContext& ctx,
return left->GetConvOperandTagAsInt() < right->GetConvOperandTagAsInt();
});

auto invoker_ptr = sh_conv_ptr->MakeInvokerPointer();
std::unique_ptr<ck::tensor_operation::device::BaseArgument> argument_ptr;
if constexpr(IsSplitKNeeded<DeviceOpType>())
{
Expand Down Expand Up @@ -907,7 +914,21 @@ ConvSolution InitInvokerFactoryNCHW(const ExecutionContext& ctx,
assert(buf_handle.get());
sh_conv_ptr->SetWorkSpacePointer(argument_ptr.get(), buf_handle.get());
}
invoker_ptr->Run(argument_ptr.get(), {handle.GetStream(), false});

auto invoker_ptr = sh_conv_ptr->MakeInvokerPointer();
{
WorkAroundHipEventProfiler prf(handle);
invoker_ptr->Run(argument_ptr.get(), {handle.GetStream(), false});
}

if(handle.IsProfilingEnabled())
{
elapsed += handle.GetKernelTime();
handle.ResetKernelTime();
handle.AccumKernelTime(elapsed);
}

// ConvertTo automatically keeps kernel time and accumulates
output_tr_inst.ConvertTo(handle, kernels, conv_tensors);
};
};
Expand Down Expand Up @@ -983,12 +1004,15 @@ ConvSolution InitInvokerFactoryNHWC(const ExecutionContext&,
data_ctx.beta.GetAsFloat());
}

auto invoker_ptr = sh_conv_ptr->MakeInvokerPointer();
HipEventProfiler pfr(handle);

float elapsed = 0.0f;
if(alpha_beta_case == DEFAULT)
{
ZeroOutTensor(handle, data_ctx.tensors.dwDesc, data_ctx.tensors.dw);

if(handle.IsProfilingEnabled())
{
elapsed += handle.GetKernelTime();
}
}
// use captured value, other wise getting warning
// "lambda capture is not used" since this variable is only used in assert.
Expand All @@ -999,7 +1023,19 @@ ConvSolution InitInvokerFactoryNHWC(const ExecutionContext&,
{
sh_conv_ptr->SetWorkSpacePointer(argument_ptr.get(), data_ctx.workSpace);
}
invoker_ptr->Run(argument_ptr.get(), {handle.GetStream(), false});

auto invoker_ptr = sh_conv_ptr->MakeInvokerPointer();
{
WorkAroundHipEventProfiler prf(handle);
invoker_ptr->Run(argument_ptr.get(), {handle.GetStream(), false});
}

if(handle.IsProfilingEnabled())
{
elapsed += handle.GetKernelTime();
handle.ResetKernelTime();
handle.AccumKernelTime(elapsed);
}
};
};
result.workspace_sz = GetWorkspaceSizeLayoutTransformConv(problem);
Expand All @@ -1020,16 +1056,31 @@ ConvSolution InitInvokerFactoryNHWC(const ExecutionContext&,
data_ctx.alpha.GetAsFloat(),
data_ctx.beta.GetAsFloat());
auto invoker_ptr = sh_conv_ptr->MakeInvokerPointer();
HipEventProfiler pfr(handle);

// Zero out the buffer for output data since it won't always write all output
// values.
float elapsed = 0.0f;
if constexpr(std::is_same_v<CastType, miopen::conv::DataInvokeParams>)
{
ZeroOutTensor(handle, data_ctx.tensors.outDesc, data_ctx.tensors.out);

if(handle.IsProfilingEnabled())
{
elapsed += handle.GetKernelTime();
}
}

invoker_ptr->Run(argument_ptr.get(), {handle.GetStream(), false});
{
WorkAroundHipEventProfiler prf(handle);
invoker_ptr->Run(argument_ptr.get(), {handle.GetStream(), false});
}

if(handle.IsProfilingEnabled())
{
elapsed += handle.GetKernelTime();
handle.ResetKernelTime();
handle.AccumKernelTime(elapsed);
}
};
};
return result;
Expand Down
Loading