| 96 | } |
| 97 | |
| 98 | ConvolutionBackwardDataImpl::Algorithm* ConvolutionBackwardDataImpl:: |
| 99 | get_algorithm_heuristic( |
| 100 | const TensorLayout& filter, const TensorLayout& diff, |
| 101 | const TensorLayout& grad, size_t workspace_limit_in_bytes, |
| 102 | const AlgoAttribute& positive_attr, |
| 103 | const AlgoAttribute& negative_attr) { |
| 104 | AlgoBase::SizeArgs args(this, filter, diff, grad); |
| 105 | |
| 106 | //! choose for large kernel cases |
| 107 | size_t fh = args.filter_meta.spatial[0], fw = args.filter_meta.spatial[1]; |
| 108 | size_t ho = diff[2], wo = diff[3]; |
| 109 | const bool prefer_dnn_lk_implbmm = args.filter_meta.format == Param::Format::NCHW && |
| 110 | ho <= 2 * fh && wo <= 2 * fw; |
| 111 | //! filter size > 9, choose large kernel cases |
| 112 | const bool prefer_direct_lk = |
| 113 | args.filter_meta.format == Param::Format::NCHW && fh > 9 && fw > 9; |
| 114 | if (prefer_dnn_lk_implbmm) { |
| 115 | #if CUDA_VERSION >= 10020 |
| 116 | if (sm_algo_pack.implbmm_nchw_hmma[0].is_available_attribute( |
| 117 | args, positive_attr, negative_attr, workspace_limit_in_bytes)) |
| 118 | return &sm_algo_pack.implbmm_nchw_hmma[0]; |
| 119 | #endif |
| 120 | if (sm_algo_pack.implbmm_nchw_fma[0].is_available_attribute( |
| 121 | args, positive_attr, negative_attr, workspace_limit_in_bytes)) |
| 122 | return &sm_algo_pack.implbmm_nchw_fma[0]; |
| 123 | } |
| 124 | |
| 125 | if (prefer_direct_lk && |
| 126 | sm_algo_pack.depthwise_large_filter.is_available_attribute( |
| 127 | args, positive_attr, negative_attr, workspace_limit_in_bytes)) { |
| 128 | return &sm_algo_pack.depthwise_large_filter; |
| 129 | } |
| 130 | |
| 131 | if (args.filter_meta.group > 1 && |
| 132 | sm_algo_pack.chanwise.is_available_attribute( |
| 133 | args, positive_attr, negative_attr, workspace_limit_in_bytes)) { |
| 134 | // prefer special chanwise impl |
| 135 | return &sm_algo_pack.chanwise; |
| 136 | } |
| 137 | |
| 138 | if (args.filter_layout->dtype.enumv() == DTypeTrait<dtype::QuantizedS8>::enumv) { |
| 139 | return megdnn::get_algo_match_attribute<ConvolutionBackwardDataImpl>( |
| 140 | sm_algo_pack.int8_algos, args, workspace_limit_in_bytes, |
| 141 | "cuda conv bwd_data", positive_attr, negative_attr); |
| 142 | } |
| 143 | |
| 144 | auto get_cudnn_algo = [this, &args, workspace_limit_in_bytes, positive_attr, |
| 145 | negative_attr]() -> ConvolutionBackwardDataImpl::AlgoBase* { |
| 146 | auto cudnn_handle = cuda::cudnn_handle(this->handle()); |
| 147 | CUDNNBwdDataDescs desc; |
| 148 | args.init_desc(desc); |
| 149 | |
| 150 | #if CUDNN_MAJOR >= 7 |
| 151 | MEGDNN_MARK_USED_VAR(negative_attr); |
| 152 | int max_count = 0; |
| 153 | cudnn_check(cudnnGetConvolutionBackwardDataAlgorithmMaxCount( |
| 154 | cudnn_handle, &max_count)); |
| 155 | SmallVector<cudnnConvolutionBwdDataAlgoPerf_t> algo_perf(max_count); |
nothing calls this directly
no test coverage detected