Method exec

dnn/src/cuda/deformable_conv/bwd_data/algo_matmul.cpp:122–187

}

void Algo::exec(const ExecArgs& args) const {
    auto&& opr = args.opr;
    auto&& handle = concrete_handle(opr->handle());
    auto&& param = opr->param();
    auto p = create_param(args, param, handle->cublas_handle(), handle->stream());
    auto bundle = get_bundle(args);
    bundle.set(args.workspace.raw_ptr);

    float* dev_im = args.im_tensor.ptr<float>();
    float* dev_filter = args.filter_tensor.ptr<float>();
    float* dev_offset = args.offset_tensor.ptr<float>();
    float* dev_mask = args.mask_tensor.ptr<float>();
    float* dev_out_grad = args.out_grad_tensor.ptr<float>();

    float* dev_im_grad = args.im_grad_tensor.ptr<float>();
    float* dev_offset_grad = args.offset_grad_tensor.ptr<float>();
    float* dev_mask_grad = args.mask_grad_tensor.ptr<float>();

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

    // clear out grad
    {
        size_t im_sz = p.batch_sz * p.IC * p.IH * p.IW * sizeof(float);
        size_t offset_sz = p.batch_sz * 2 * p.deformable_group * p.FH * p.FW * p.OH *
                           p.OW * sizeof(float);
        size_t mask_sz = p.batch_sz * p.deformable_group * p.FH * p.FW * p.OH * p.OW *
                         sizeof(float);

        cudaMemsetAsync(dev_im_grad, 0, im_sz, p.stream);
        cudaMemsetAsync(dev_offset_grad, 0, offset_sz, p.stream);
        cudaMemsetAsync(dev_mask_grad, 0, mask_sz, p.stream);
    }

    // relayout out_grad to [oc, N, OH, OW]
    {
        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(relayout_ws1, C3l);

        args.handle->relayout_opr()->exec(C2, C3);
    }
    // matmul [g, icpg, FH, FW, ocpg] * [g, ocpg, N, OH, OW] =>
    //        => [g, icpg, FH, FW, N, OH, OW]
    {
        auto config = prepare_sub_opr(args);

        TensorND A(static_cast<void*>(dev_filter), config.first[0]),
                B(static_cast<void*>(relayout_ws1), config.first[1]),
                C(static_cast<void*>(result_ws), config.first[2]);

        size_t bmm_ws_size = bundle.get_size(0);