Skip to content

Commit

Permalink
fix-trans-conv-issue-2459(02) [driver] Adjust tolerance for WrW trans…
Browse files Browse the repository at this point in the history 
…posed convolutions.
  • Loading branch information
@atamazov
atamazov committed Oct 26, 2023
1 parent 23b810f commit 3d3710c
Showing 1 changed file with 27 additions and 38 deletions.
65 changes: 27 additions & 38 deletions driver/conv_driver.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -360,6 +360,7 @@ class ConvDriver : public Driver

miopenConvolutionDescriptor_t convDesc;
miopenConvolutionDescriptor_t warmupConvDesc;
miopenConvolutionMode_t mode;

bool is_wrw = true, is_bwd = true, is_fwd = true;
bool is_wrw_winograd = false;
Expand Down Expand Up @@ -536,6 +537,19 @@ int ConvDriver<Tgpu, Tref>::ParseCmdLineArgs(int argc, char* argv[])
inflags.GetValueStr("out_layout") + "c" + std::to_string(vector_length));
}

if((inflags.GetValueStr("mode")) == "conv")
{
mode = miopenConvolution;
}
else if((inflags.GetValueStr("mode")) == "trans")
{
mode = miopenTranspose;
}
else
{
MIOPEN_THROW("Incorrect Convolution Mode\n");
}

num_iterations = inflags.GetValueInt("iter");
if(num_iterations < 1)
{
Expand Down Expand Up @@ -755,13 +769,12 @@ int ConvDriver<Tgpu, Tref>::GetandSetData()
std::vector<int> pads = {0, 0};
std::vector<int> conv_strides = {1, 1};
std::vector<int> conv_dilations = {1, 1};
miopenConvolutionMode_t mode = miopenConvolution;
miopenInitConvolutionNdDescriptor(warmupConvDesc,
spatial_dim,
pads.data(),
conv_strides.data(),
conv_dilations.data(),
mode);
miopenConvolution);
miopenSetConvolutionFindMode(warmupConvDesc, miopenConvolutionFindModeNormal);
miopenHiddenSetConvolutionFindMode(
warmupConvDesc,
Expand Down Expand Up @@ -980,20 +993,6 @@ std::vector<int> ConvDriver<Tgpu, Tref>::GetWeightTensorLengthsFromCmdLine()
}
}

miopenConvolutionMode_t mode;
if((inflags.GetValueStr("mode")) == "conv")
{
mode = miopenConvolution;
}
else if((inflags.GetValueStr("mode")) == "trans")
{
mode = miopenTranspose;
}
else
{
MIOPEN_THROW("Incorrect Convolution Mode\n");
}

if(mode == miopenTranspose)
{
wei_lens[0] = wei_c_len;
Expand Down Expand Up @@ -1087,21 +1086,6 @@ int ConvDriver<Tgpu, Tref>::SetConvDescriptorFromCmdLineArgs()
}
}

miopenConvolutionMode_t mode;
if((inflags.GetValueStr("mode")) == "conv")
{
mode = miopenConvolution;
}
else if((inflags.GetValueStr("mode")) == "trans")
{
mode = miopenTranspose;
}
else
{
printf("Incorrect Convolution Mode\n");
exit(0); // NOLINT (concurrency-mt-unsafe)
}

// adjust padding based on user-defined padding mode
if(mode == miopenConvolution &&
(miopen::all_of(conv_dilations, [](auto v) { return v == 1; }) ||
Expand Down Expand Up @@ -2252,7 +2236,7 @@ int ConvDriver<Tgpu, Tref>::RunForwardGpuImmed(const bool is_transform)
template <typename Tgpu, typename Tref>
int ConvDriver<Tgpu, Tref>::RunForwardCPU()
{
if(miopen::deref(convDesc).mode == miopenTranspose)
if(mode == miopenTranspose)
{
cpu_convolution_backward_data(miopen::deref(convDesc).GetSpatialDimension(),
outhost,
Expand Down Expand Up @@ -2308,7 +2292,7 @@ int ConvDriver<Tgpu, Tref>::RunForwardGPUReference()
std::cout << "gpu reference convolution does not support bias yet" << std::endl;
return -1;
}
auto ref_solution_id = miopen::deref(convDesc).mode == miopenTranspose
auto ref_solution_id = mode == miopenTranspose //
? miopen::solver::Id("ConvDirectNaiveConvBwd").Value()
: miopen::solver::Id("ConvDirectNaiveConvFwd").Value();
auto rc = miopenConvolutionForwardImmediate(handle,
Expand Down Expand Up @@ -3142,7 +3126,7 @@ int ConvDriver<Tgpu, Tref>::RunBackwardWrwGpuImmed()
template <typename Tgpu, typename Tref>
int ConvDriver<Tgpu, Tref>::RunBackwardWeightsCPU()
{
if(miopen::deref(convDesc).mode == miopenTranspose)
if(mode == miopenTranspose)
{
cpu_convolution_backward_weight(miopen::deref(convDesc).GetSpatialDimension(),
dout,
Expand Down Expand Up @@ -3178,7 +3162,7 @@ int ConvDriver<Tgpu, Tref>::RunBackwardWeightsCPU()
template <typename Tgpu, typename Tref>
int ConvDriver<Tgpu, Tref>::RunBackwardDataCPU()
{
if(miopen::deref(convDesc).mode == miopenTranspose)
if(mode == miopenTranspose)
{
cpu_convolution_forward(miopen::deref(convDesc).GetSpatialDimension(),
dout,
Expand Down Expand Up @@ -3275,7 +3259,7 @@ int ConvDriver<Tgpu, Tref>::RunBackwardDataGPUReference()
{
AutoPrepareForGpuReference naive_conv_enable;

auto ref_solution_id = miopen::deref(convDesc).mode == miopenTranspose
auto ref_solution_id = mode == miopenTranspose //
? miopen::solver::Id("ConvDirectNaiveConvFwd").Value()
: miopen::solver::Id("ConvDirectNaiveConvBwd").Value();
auto rc = miopenConvolutionBackwardDataImmediate(handle,
Expand Down Expand Up @@ -3324,7 +3308,7 @@ std::string ConvDriver<Tgpu, Tref>::GetVerificationCacheFileName(
{
std::ostringstream ss;

miopenConvolutionMode_t mode;
miopenConvolutionMode_t unused;

int spatial_dim = inflags.GetValueInt("spatial_dim");

Expand All @@ -3339,7 +3323,7 @@ std::string ConvDriver<Tgpu, Tref>::GetVerificationCacheFileName(
pads.data(),
conv_strides.data(),
conv_dilations.data(),
&mode);
&unused);

auto get_basename_string = [&]() {
switch(direction)
Expand Down Expand Up @@ -3538,6 +3522,11 @@ int ConvDriver<Tgpu, Tref>::VerifyBackward()
// WrW deviation is ~twice worse than Bwd due to more FP computations involved,
// which means more roundings, so GPU amd CPU computations diverge more.
auto tolerance = 2 * GetDefaultTolerance();

// fp32 transposed convolutions show worse precision.
if(mode == miopenTranspose && std::is_same<Tgpu, float>::value)
tolerance *= 2;

// Winograd and iGemm WrW algorithms reveal bigger deviation than other algos.
if(is_wrw_winograd && std::is_same<Tgpu, float>::value)
{
Expand Down

0 comments on commit 3d3710c

Please sign in to comment.