Method exec

dnn/src/cuda/local/backward_filter.cpp:22–46

namespace cuda {

void LocalBackwardFilterImpl::exec(
        _megdnn_tensor_in src, _megdnn_tensor_in diff, _megdnn_tensor_out grad,
        _megdnn_workspace workspace) {
    check_exec(src.layout, diff.layout, grad.layout, workspace.size);
    megdnn_assert(param().mode == Mode::CROSS_CORRELATION);
    auto N = src.layout.shape[0], IC = src.layout.shape[1], IH = src.layout.shape[2],
         IW = src.layout.shape[3];
    auto OC = diff.layout.shape[1], OH = diff.layout.shape[2],
         OW = diff.layout.shape[3];
    auto FH = grad.layout.shape[3], FW = grad.layout.shape[4];
    auto handle = concrete_handle(this->handle());
    auto stream = cuda_stream(this->handle());
    auto cublas = cublas_handle(this->handle());
    auto one = handle->one_device();
    auto zero = handle->zero_device();
    if (use_cuda_convnet(src.layout, diff.layout, grad.layout)) {
        local::backward_filter_proxy_convnet(
                src.ptr<dt_float32>(), diff.ptr<dt_float32>(), grad.ptr<dt_float32>(),
                reinterpret_cast<float*>(workspace.raw_ptr), N, IC, IH, IW, OC, OH, OW,
                FH, FW, IC * IH * IW, OC * OH * OW, param().pad_h, param().pad_w,
                param().stride_h, param().stride_w, cublas, stream, one, zero);
    } else {
        local::boom_backward_filter();
    }
}

size_t LocalBackwardFilterImpl::get_workspace_in_bytes(
        const TensorLayout& src, const TensorLayout& diff, const TensorLayout& grad) {