Method try_transpose

dnn/src/cuda/relayout/opr_impl.cpp:19–63

}

bool RelayoutForwardImpl::Param::try_transpose() {
    if (m_dst.layout.dtype.is_low_bit())
        return false;
    relayout::TransposeParam transp;
    bool trans = relayout::is_transpose(m_src.layout, m_dst.layout, transp);
    if (!trans)
        return false;
    size_t dsize = transp.c * m_src.layout.dtype.size();
    if (dsize != 1 && dsize != 2 && dsize != 4)
        return false;

    if (m_src.layout.dtype == dtype::Float32() && transp.batch == 1 && transp.c == 1) {
        auto handle = concrete_handle(m_opr->handle());
        cublas_check(cublasSgeam(
                handle->cublas_handle(), CUBLAS_OP_T, CUBLAS_OP_T, transp.m, transp.n,
                handle->one_device(), m_src.ptr<dt_float32>(), transp.n,
                handle->zero_device(), m_src.ptr<dt_float32>(), transp.n,
                m_dst.ptr<dt_float32>(), transp.m));
        return true;
    }
    float square_ratio = static_cast<float>(transp.m) / static_cast<float>(transp.n);
    if (transp.m < 32 || transp.n < 32 || square_ratio < 0.5f || square_ratio > 2.f)
        return false;
    size_t batch = transp.batch, m = transp.m, n = transp.n;
    size_t lda = n, ldb = m, stride_A = m * n, stride_B = m * n;
    auto&& stream = m_opr->stream();
#define RUN(_dt)                                                                  \
    do {                                                                          \
        typedef DTypeTrait<dtype::_dt>::ctype ctype;                              \
        copy_by_transpose<ctype>(                                                 \
                reinterpret_cast<const ctype*>(m_src.raw_ptr()),                  \
                reinterpret_cast<ctype*>(m_dst.raw_ptr()), batch, m, n, lda, ldb, \
                stride_A, stride_B, stream);                                      \
        return true;                                                              \
    } while (0)
    switch (dsize) {
        case 1:
            RUN(Int8);
        case 2:
            RUN(Float16);
        case 4:
            RUN(Int32);
    }
    megdnn_assert(0, "bad dtype size");
}

bool RelayoutForwardImpl::Param::try_copy_contig() {
    auto &&lsrc = m_src.layout, &&ldst = m_dst.layout;