From 4bd61bbb66ab8dd9b5383ded93c7e7dbc54db66d Mon Sep 17 00:00:00 2001 From: bibek <108366729+bghimireamd@users.noreply.github.com> Date: Fri, 15 Nov 2024 14:48:40 -0600 Subject: [PATCH 1/2] Cherry pick bn 3d for rocm rel 6.3 (#3387) * fix bn 3d issue * fix review comments * fix typo --- src/batch_norm.cpp | 19 +++-- src/batch_norm_api.cpp | 160 +++++++++++++++++++++-------------------- 2 files changed, 94 insertions(+), 85 deletions(-) diff --git a/src/batch_norm.cpp b/src/batch_norm.cpp index 938809d81c..2c5486f307 100644 --- a/src/batch_norm.cpp +++ b/src/batch_norm.cpp @@ -66,27 +66,34 @@ void DeriveBNTensorDescriptor(TensorDescriptor& derivedBnDesc, TensorDescriptor BuildReshaped4DTensorDescriptor(const miopen::TensorDescriptor& tDesc) { + std::vector dims(tDesc.GetLengths()); + auto dataType = tDesc.GetType(); auto layout = tDesc.GetLayout_t(); if(layout == miopenTensorNCDHW) { layout = miopenTensorNCHW; + + // NxCxDxHxW -> NxCx(D*H)xW + dims[2] *= dims[3]; + dims[3] = dims[4]; + dims.pop_back(); } else if(layout == miopenTensorNDHWC) { layout = miopenTensorNHWC; + + // NxDxHxWxC -> Nx(D*H)xWxC + dims[1] *= dims[2]; + dims[2] = dims[3]; + dims[3] = dims[4]; + dims.pop_back(); } else { std::cout << "Cannot handle layout : " << layout << "\n"; exit(EXIT_FAILURE); // NOLINT (concurrency-mt-unsafe) } - std::vector dims(tDesc.GetLengths()); - - // NxCxDxHxW -> NxCx(D*H)xW - dims[2] *= dims[3]; - dims[3] = dims[4]; - dims.pop_back(); return {dataType, layout, dims}; } diff --git a/src/batch_norm_api.cpp b/src/batch_norm_api.cpp index d3b824cee0..cb6b7c0842 100644 --- a/src/batch_norm_api.cpp +++ b/src/batch_norm_api.cpp @@ -206,7 +206,7 @@ miopenBatchNormalizationForwardInference_V2(miopenHandle_t handle, const miopenTensorDescriptor_t yDesc, void* y, const miopenTensorDescriptor_t scaleDesc, - const miopenTensorDescriptor_t BiasDesc, + const miopenTensorDescriptor_t biasDesc, const miopenTensorDescriptor_t estMeanDesc, const miopenTensorDescriptor_t estVarianceDesc, void* bnScale, @@ -222,7 +222,7 @@ miopenBatchNormalizationForwardInference_V2(miopenHandle_t handle, yDesc, y, scaleDesc, - BiasDesc, + biasDesc, estMeanDesc, estVarianceDesc, bnScale, @@ -239,31 +239,31 @@ miopenBatchNormalizationForwardInference_V2(miopenHandle_t handle, nullptr, nullptr, miopen::debug::BatchNormDirection_t::ForwardInference); - - // In case of NxCxDxHxW int size{0}; miopenGetTensorDescriptorSize(xDesc, &size); + // In case of NxCxDxHxW + auto ReshapeIfNeeded = [size](const auto desc) { + return (size == 5) ? miopen::BuildReshaped4DTensorDescriptor(miopen::deref(desc)) + : miopen::deref(desc); + }; return miopen::try_([&] { - miopen::BatchNormForwardInference( - miopen::deref(handle), - bn_mode, - alpha, - beta, - (size == 5) ? miopen::BuildReshaped4DTensorDescriptor(miopen::deref(xDesc)) - : miopen::deref(xDesc), - DataCast(x), - (size == 5) ? miopen::BuildReshaped4DTensorDescriptor(miopen::deref(yDesc)) - : miopen::deref(yDesc), - DataCast(y), - miopen::deref(scaleDesc), - miopen::deref(BiasDesc), - miopen::deref(estMeanDesc), - miopen::deref(estVarianceDesc), - DataCast(bnScale), - DataCast(bnBias), - DataCast(estimatedMean), - DataCast(estimatedVariance), - epsilon); + miopen::BatchNormForwardInference(miopen::deref(handle), + bn_mode, + alpha, + beta, + ReshapeIfNeeded(xDesc), + DataCast(x), + ReshapeIfNeeded(yDesc), + DataCast(y), + ReshapeIfNeeded(scaleDesc), + ReshapeIfNeeded(biasDesc), + ReshapeIfNeeded(estMeanDesc), + ReshapeIfNeeded(estVarianceDesc), + DataCast(bnScale), + DataCast(bnBias), + DataCast(estimatedMean), + DataCast(estimatedVariance), + epsilon); }); } @@ -277,7 +277,7 @@ miopenBatchNormalizationForwardTraining_V2(miopenHandle_t handle, const miopenTensorDescriptor_t yDesc, void* y, const miopenTensorDescriptor_t scaleDesc, - const miopenTensorDescriptor_t BiasDesc, + const miopenTensorDescriptor_t biasDesc, const miopenTensorDescriptor_t savedMeanDesc, const miopenTensorDescriptor_t savedVarianceDesc, void* bnScale, @@ -296,7 +296,7 @@ miopenBatchNormalizationForwardTraining_V2(miopenHandle_t handle, yDesc, y, scaleDesc, - BiasDesc, + biasDesc, savedMeanDesc, savedVarianceDesc, bnScale, @@ -316,33 +316,35 @@ miopenBatchNormalizationForwardTraining_V2(miopenHandle_t handle, resultSaveMean, resultSaveInvVariance, miopen::debug::BatchNormDirection_t::ForwardTraining); - // In case of NxCxDxHxW + int size{0}; miopenGetTensorDescriptorSize(xDesc, &size); + // In case of NxCxDxHxW + auto ReshapeIfNeeded = [size](const auto desc) { + return (size == 5) ? miopen::BuildReshaped4DTensorDescriptor(miopen::deref(desc)) + : miopen::deref(desc); + }; return miopen::try_([&] { - miopen::BatchNormForwardTraining( - miopen::deref(handle), - bn_mode, - alpha, - beta, - (size == 5) ? miopen::BuildReshaped4DTensorDescriptor(miopen::deref(xDesc)) - : miopen::deref(xDesc), - DataCast(x), - (size == 5) ? miopen::BuildReshaped4DTensorDescriptor(miopen::deref(yDesc)) - : miopen::deref(yDesc), - DataCast(y), - miopen::deref(scaleDesc), - miopen::deref(BiasDesc), - miopen::deref(savedMeanDesc), - miopen::deref(savedVarianceDesc), - DataCast(bnScale), - DataCast(bnBias), - expAvgFactor, - DataCast(resultRunningMean), - DataCast(resultRunningVariance), - epsilon, - DataCast(resultSaveMean), - DataCast(resultSaveInvVariance)); + miopen::BatchNormForwardTraining(miopen::deref(handle), + bn_mode, + alpha, + beta, + ReshapeIfNeeded(xDesc), + DataCast(x), + ReshapeIfNeeded(yDesc), + DataCast(y), + ReshapeIfNeeded(scaleDesc), + ReshapeIfNeeded(biasDesc), + ReshapeIfNeeded(savedMeanDesc), + ReshapeIfNeeded(savedVarianceDesc), + DataCast(bnScale), + DataCast(bnBias), + expAvgFactor, + DataCast(resultRunningMean), + DataCast(resultRunningVariance), + epsilon, + DataCast(resultSaveMean), + DataCast(resultSaveInvVariance)); }); } @@ -360,7 +362,7 @@ miopenBatchNormalizationBackward_V2(miopenHandle_t handle, const miopenTensorDescriptor_t dxDesc, void* dx, const miopenTensorDescriptor_t scaleDesc, - const miopenTensorDescriptor_t BiasDesc, + const miopenTensorDescriptor_t biasDesc, const miopenTensorDescriptor_t savedMeanDesc, const miopenTensorDescriptor_t savedVarianceDesc, const void* bnScale, @@ -379,7 +381,7 @@ miopenBatchNormalizationBackward_V2(miopenHandle_t handle, dxDesc, dx, scaleDesc, - BiasDesc, + biasDesc, savedMeanDesc, savedVarianceDesc, bnScale, @@ -396,35 +398,35 @@ miopenBatchNormalizationBackward_V2(miopenHandle_t handle, savedMean, savedInvVariance, miopen::debug::BatchNormDirection_t::Backward); - // In case of NxCxDxHxW int size{0}; miopenGetTensorDescriptorSize(xDesc, &size); + // In case of NxCxDxHxW + auto ReshapeIfNeeded = [size](const auto desc) { + return (size == 5) ? miopen::BuildReshaped4DTensorDescriptor(miopen::deref(desc)) + : miopen::deref(desc); + }; return miopen::try_([&] { - miopen::BatchNormBackward( - miopen::deref(handle), - bn_mode, - alphaDataDiff, - betaDataDiff, - alphaParamDiff, - betaParamDiff, - (size == 5) ? miopen::BuildReshaped4DTensorDescriptor(miopen::deref(xDesc)) - : miopen::deref(xDesc), - DataCast(x), - (size == 5) ? miopen::BuildReshaped4DTensorDescriptor(miopen::deref(dyDesc)) - : miopen::deref(dyDesc), - DataCast(dy), - (size == 5) ? miopen::BuildReshaped4DTensorDescriptor(miopen::deref(dxDesc)) - : miopen::deref(dxDesc), - DataCast(dx), - miopen::deref(scaleDesc), - miopen::deref(BiasDesc), - miopen::deref(savedMeanDesc), - miopen::deref(savedVarianceDesc), - DataCast(bnScale), - DataCast(resultBnScaleDiff), - DataCast(resultBnBiasDiff), - epsilon, - DataCast(savedMean), - DataCast(savedInvVariance)); + miopen::BatchNormBackward(miopen::deref(handle), + bn_mode, + alphaDataDiff, + betaDataDiff, + alphaParamDiff, + betaParamDiff, + ReshapeIfNeeded(xDesc), + DataCast(x), + ReshapeIfNeeded(dyDesc), + DataCast(dy), + ReshapeIfNeeded(dxDesc), + DataCast(dx), + ReshapeIfNeeded(scaleDesc), + ReshapeIfNeeded(biasDesc), + ReshapeIfNeeded(savedMeanDesc), + ReshapeIfNeeded(savedVarianceDesc), + DataCast(bnScale), + DataCast(resultBnScaleDiff), + DataCast(resultBnBiasDiff), + epsilon, + DataCast(savedMean), + DataCast(savedInvVariance)); }); } From cd1a3f1344020aad0a5f5e1d9be7a747b964b5dc Mon Sep 17 00:00:00 2001 From: Vasilii Filippov Date: Fri, 15 Nov 2024 21:06:38 +0100 Subject: [PATCH 2/2] Revert "Fixed incorrect transpose in find 2.0 (#3285)" This reverts commit 2d69aebdcf1b22df1357a38b68e1b8ddefd8b424. --- src/problem.cpp | 22 +++++++++++----------- 1 file changed, 11 insertions(+), 11 deletions(-) diff --git a/src/problem.cpp b/src/problem.cpp index ba84856850..fed48dfe88 100644 --- a/src/problem.cpp +++ b/src/problem.cpp @@ -179,11 +179,7 @@ Problem::FindSolutions(Handle& handle, const FindOptions& options, std::size_t m auto ret = std::visit( boost::hof::match( [&](const ConvolutionDescriptor& op_desc) { - if(op_desc.mode == miopenTranspose) - return MakeTransposed().FindSolutionsImpl( - handle, options, max_solutions, buffers, op_desc); - else - return FindSolutionsImpl(handle, options, max_solutions, buffers, op_desc); + return FindSolutionsImpl(handle, options, max_solutions, buffers, op_desc); }, [&](const SoftmaxDescriptor& op_desc) { return FindSolutionsImpl(handle, options, max_solutions, buffers, op_desc); @@ -481,17 +477,21 @@ std::vector Problem::FindSolutionsImpl(Handle& handle, const auto& w = buffers.at(miopenTensorConvolutionW); auto y = buffers.at(miopenTensorConvolutionY); - if(conv_desc.mode == miopenTranspose) - std::swap(x, y); - - const auto conv_problem = AsConvolution(); - - ValidateGroupCount(x_desc, w_desc, conv_desc); + const auto conv_problem = + conv_desc.mode == miopenTranspose ? MakeTransposed().AsConvolution() : AsConvolution(); std::size_t workspace_size; Allocator::ManageDataPtr owned_workspace; Data_t workspace; + if(conv_desc.mode == miopenTranspose) + { + std::swap(x, y); + std::swap(x_desc, y_desc); + } + + ValidateGroupCount(x_desc, w_desc, conv_desc); + if(options.preallocated_workspace) { workspace = options.preallocated_workspace->buffer;