From 16390146d3e6bbf4049cc25afa25d346d2de96da Mon Sep 17 00:00:00 2001 From: Peter Yeh Date: Fri, 28 Sep 2018 20:24:19 -0700 Subject: [PATCH] Revert "[Caffe2] MIOpen dims change check (#229)" This reverts commit 1f7bb657a127fd0348324d6c6cd27400227938fb. --- caffe2/operators/hip/conv_op_miopen.cc | 222 +++++++++++-------------- 1 file changed, 93 insertions(+), 129 deletions(-) diff --git a/caffe2/operators/hip/conv_op_miopen.cc b/caffe2/operators/hip/conv_op_miopen.cc index 85da5fdccd31d2..13f2428bf9daa2 100644 --- a/caffe2/operators/hip/conv_op_miopen.cc +++ b/caffe2/operators/hip/conv_op_miopen.cc @@ -66,6 +66,19 @@ class MIOPENConvOpBase : public ConvPoolOpBase { dilation_h() == 1 && dilation_w() == 1, "MIOpen convolution does not support dilation for groups > 1."); } + + MIOPEN_ENFORCE(miopenInitConvolutionDescriptor( + conv_desc_, + mode_, + pad_t(), + pad_l(), + stride_h(), + stride_w(), + dilation_h(), + dilation_w())); + + MIOPEN_ENFORCE(miopenSetConvolutionGroupCount( + conv_desc_, group_)); } ~MIOPENConvOpBase() { @@ -78,8 +91,6 @@ class MIOPENConvOpBase : public ConvPoolOpBase { } protected: - vector mio_input_dims_; - vector mio_weight_dims_; MIOPENWrapper miopen_wrapper_; miopenTensorDescriptor_t bottom_desc_; miopenTensorDescriptor_t bias_desc_; @@ -246,59 +257,35 @@ bool MIOPENConvOp::DoRunWithType() { "If you set group, the number of output channels should be divisible " "by group."); - bool input_changed = (X.dims() != mio_input_dims_); - bool weight_changed = (Weight.dims() != mio_weight_dims_); - - if (input_changed || weight_changed) { - VLOG(1) << "Changing MIOpen descriptor configurations."; - if (input_changed) { - mio_input_dims_ = X.dims(); - MIOPEN_ENFORCE(miopenSet4dTensorDescriptor( - bottom_desc_, miopenTypeWrapper::type, N, C, H, W)); - } + MIOPEN_ENFORCE(miopenSet4dTensorDescriptor( + bottom_desc_, miopenTypeWrapper::type, N, C, H, W)); - if (weight_changed) { - mio_weight_dims_ = Weight.dims(); - MIOPEN_ENFORCE(miopenInitConvolutionDescriptor( - conv_desc_, - mode_, - pad_t(), - pad_l(), - stride_h(), - stride_w(), - dilation_h(), - dilation_w())); - - MIOPEN_ENFORCE(miopenSetConvolutionGroupCount( - conv_desc_, group_)); - - MIOPEN_ENFORCE(miopenSet4dTensorDescriptor( - weight_desc_, - miopenTypeWrapper::type, - M, - C / group_, - kernel_h(), - kernel_w())); - } + MIOPEN_ENFORCE(miopenSet4dTensorDescriptor( + weight_desc_, + miopenTypeWrapper::type, + M, + C / group_, + kernel_h(), + kernel_w())); - MIOPEN_ENFORCE(miopenGetConvolutionForwardOutputDim( - conv_desc_, - bottom_desc_, - weight_desc_, - &N_out, - &C_out, - &H_out, - &W_out)); + MIOPEN_ENFORCE(miopenGetConvolutionForwardOutputDim( + conv_desc_, + bottom_desc_, + weight_desc_, + &N_out, + &C_out, + &H_out, + &W_out)); - MIOPEN_ENFORCE(miopenSet4dTensorDescriptor( - top_desc_, miopenTypeWrapper::type, N_out, C_out, H_out, W_out)); + MIOPEN_ENFORCE(miopenSet4dTensorDescriptor( + top_desc_, miopenTypeWrapper::type, N_out, C_out, H_out, W_out)); - if (InputSize() == 3) { + if (InputSize() == 3) { MIOPEN_ENFORCE(miopenSet4dTensorDescriptor( bias_desc_, miopenTypeWrapper::type, 1, M, 1, 1)); - } + } - while (!bestAlgoFound_) { + while (!bestAlgoFound_) { miopenConvAlgoPerf_t perf; MIOPEN_ENFORCE(miopenConvolutionForwardGetWorkSpaceSize( @@ -331,8 +318,8 @@ bool MIOPENConvOp::DoRunWithType() { }); bestAlgoFound_ = true; fwdAlgo_ = perf.fwd_algo; - } } + miopen_wrapper_.with_miopen_state(miopen_state_, [&](MIOPENState* state) { MIOPEN_ENFORCE(miopenConvolutionForward( state->miopen_handle(), @@ -437,59 +424,36 @@ bool MIOPENConvGradientOp::DoRunWithType() { "by group."); bool doBwdDataComputation = (OutputSize() == 3 || (no_bias_ && (OutputSize() == 2))); - bool input_changed = (X.dims() != mio_input_dims_); - bool weight_changed = (Weight.dims() != mio_weight_dims_); - - if (input_changed || weight_changed) { - VLOG(1) << "Changing MIOpen descriptor configurations."; - if (input_changed) { - mio_input_dims_ = X.dims(); - MIOPEN_ENFORCE(miopenSet4dTensorDescriptor( - bottom_desc_, miopenTypeWrapper::type, N, C, H, W)); - } - if (weight_changed) { - mio_weight_dims_ = Weight.dims(); - MIOPEN_ENFORCE(miopenInitConvolutionDescriptor( - conv_desc_, - mode_, - pad_t(), - pad_l(), - stride_h(), - stride_w(), - dilation_h(), - dilation_w())); - - MIOPEN_ENFORCE(miopenSetConvolutionGroupCount( - conv_desc_, group_)); + MIOPEN_ENFORCE(miopenSet4dTensorDescriptor( + bottom_desc_, miopenTypeWrapper::type, N, C, H, W)); - MIOPEN_ENFORCE(miopenSet4dTensorDescriptor( - weight_desc_, - miopenTypeWrapper::type, - M, - C / group_, - kernel_h(), - kernel_w())); - } + MIOPEN_ENFORCE(miopenSet4dTensorDescriptor( + weight_desc_, + miopenTypeWrapper::type, + M, + C / group_, + kernel_h(), + kernel_w())); - MIOPEN_ENFORCE(miopenGetConvolutionForwardOutputDim( - conv_desc_, - bottom_desc_, - weight_desc_, - &N_out, - &C_out, - &H_out, - &W_out)); + MIOPEN_ENFORCE(miopenGetConvolutionForwardOutputDim( + conv_desc_, + bottom_desc_, + weight_desc_, + &N_out, + &C_out, + &H_out, + &W_out)); - MIOPEN_ENFORCE(miopenSet4dTensorDescriptor( - top_desc_, miopenTypeWrapper::type, N_out, C_out, H_out, W_out)); + MIOPEN_ENFORCE(miopenSet4dTensorDescriptor( + top_desc_, miopenTypeWrapper::type, N_out, C_out, H_out, W_out)); - if (!no_bias_) { - MIOPEN_ENFORCE(miopenSet4dTensorDescriptor( - bias_desc_, miopenTypeWrapper::type, 1, M, 1, 1)); - } + if (!no_bias_) { + MIOPEN_ENFORCE(miopenSet4dTensorDescriptor( + bias_desc_, miopenTypeWrapper::type, 1, M, 1, 1)); + } - while ((!bestDataAlgoFound_) && doBwdDataComputation) { + while ((!bestDataAlgoFound_) && doBwdDataComputation) { miopenConvAlgoPerf_t perf; MIOPEN_ENFORCE(miopenConvolutionBackwardDataGetWorkSpaceSize( @@ -523,43 +487,43 @@ bool MIOPENConvGradientOp::DoRunWithType() { bestDataAlgoFound_ = true; bwdDataAlgo_ = perf.bwd_data_algo; - } + } - while (!bestWeightAlgoFound_) { - miopenConvAlgoPerf_t perf; + while (!bestWeightAlgoFound_) { + miopenConvAlgoPerf_t perf; - MIOPEN_ENFORCE(miopenConvolutionBackwardWeightsGetWorkSpaceSize( - miopen_wrapper_.inline_miopen_handle(), - top_desc_, - bottom_desc_, - conv_desc_, - weight_desc_, - &bwdWeightWsSize_)); - if ((bwdWeightWsSize_ > 0) && (bwdWeightWs_ == nullptr)) { - HIP_CHECK(hipMalloc(&bwdWeightWs_, bwdWeightWsSize_)); - } + MIOPEN_ENFORCE(miopenConvolutionBackwardWeightsGetWorkSpaceSize( + miopen_wrapper_.inline_miopen_handle(), + top_desc_, + bottom_desc_, + conv_desc_, + weight_desc_, + &bwdWeightWsSize_)); + if ((bwdWeightWsSize_ > 0) && (bwdWeightWs_ == nullptr)) { + HIP_CHECK(hipMalloc(&bwdWeightWs_, bwdWeightWsSize_)); + } - miopen_wrapper_.with_miopen_state(miopen_state_, [&](MIOPENState* state) { - MIOPEN_ENFORCE(miopenFindConvolutionBackwardWeightsAlgorithm( - state->miopen_handle(), - top_desc_, - dY.template data(), - bottom_desc_, - X.template data(), - conv_desc_, - weight_desc_, - dW->template mutable_data(), - requestAlgoCount_, - &returnedAlgoCount_, - &perf, - bwdWeightWs_, - bwdWeightWsSize_, - false)); - }); - bestWeightAlgoFound_ = true; - bwdWeiAlgo_ = perf.bwd_weights_algo; - } + miopen_wrapper_.with_miopen_state(miopen_state_, [&](MIOPENState* state) { + MIOPEN_ENFORCE(miopenFindConvolutionBackwardWeightsAlgorithm( + state->miopen_handle(), + top_desc_, + dY.template data(), + bottom_desc_, + X.template data(), + conv_desc_, + weight_desc_, + dW->template mutable_data(), + requestAlgoCount_, + &returnedAlgoCount_, + &perf, + bwdWeightWs_, + bwdWeightWsSize_, + false)); + }); + bestWeightAlgoFound_ = true; + bwdWeiAlgo_ = perf.bwd_weights_algo; } + if (doBwdDataComputation) { miopen_wrapper_.with_miopen_state(miopen_state_, [&](MIOPENState* state) { MIOPEN_ENFORCE(miopenConvolutionBackwardData(