Method exec

dnn/src/cuda/deformable_conv/bwd_flt/algo_matmul.cpp:122–165

}

void Algo::exec(const ExecArgs& args) const {
    auto&& opr = args.opr;
    auto&& param = opr->param();
    auto&& handle = concrete_handle(opr->handle());

    auto p = create_param(args, param, handle->cublas_handle(), handle->stream());

    auto bundle = get_bundle(args);
    bundle.set(args.workspace.raw_ptr);

    const float* dev_im = args.im_tensor.ptr<float>();
    const float* dev_offset = args.offset_tensor.ptr<float>();
    const float* dev_mask = args.mask_tensor.ptr<float>();
    float* dev_out_grad = args.out_grad_tensor.ptr<float>();
    float* dev_filter_grad = args.filter_grad_tensor.ptr<float>();

    float* col_ws = static_cast<float*>(bundle.get(0));
    float* out_grad_ws = static_cast<float*>(bundle.get(1));
    void* bmm_ws = bundle.get(2);

    // im2col
    deformable_conv::im2col(dev_im, dev_offset, dev_mask, col_ws, p);
    // relayout
    auto&& dt = args.im_layout.dtype;
    size_t dim0 = p.batch_sz, dim1 = p.OC, dim2 = p.OH * p.OW;
    TensorLayout C2l({dim0, dim1, dim2}, dt), C3l = C2l;
    C3l.stride[0] = dim2;
    C3l.stride[1] = dim0 * dim2;
    C3l.stride[2] = 1;
    TensorND C2(dev_out_grad, C2l);
    TensorND C3(out_grad_ws, C3l);

    args.handle->relayout_opr()->exec(C2, C3);
    // matmul
    auto config = prepare_sub_opr(args);

    TensorND A(static_cast<void*>(out_grad_ws), config.first[0]),
            B(static_cast<void*>(col_ws), config.first[1]),
            C(static_cast<void*>(dev_filter_grad), config.first[2]);

    size_t bmm_ws_size = bundle.get_size(2);
    config.second->exec(
            A, B, C, Workspace(static_cast<megdnn::dt_byte*>(bmm_ws), bmm_ws_size));
}
// vim: syntax=cpp.doxygen