opencv
diff --git a/‎modules/dnn/include/opencv2/dnn/all_layers.hpp
Copy file name to clipboard
+5Lines changed: 5 additions & 0 deletions b/‎modules/dnn/include/opencv2/dnn/all_layers.hpp
Copy file name to clipboard
+5Lines changed: 5 additions & 0 deletions
diff --git a/‎modules/dnn/perf/perf_gemm.cpp
Copy file name to clipboard
+168-2Lines changed: 168 additions & 2 deletions b/‎modules/dnn/perf/perf_gemm.cpp
Copy file name to clipboard
+168-2Lines changed: 168 additions & 2 deletions
diff --git a/‎modules/dnn/src/cuda4dnn/csl/cublas.hpp
Copy file name to clipboard
+140Lines changed: 140 additions & 0 deletions b/‎modules/dnn/src/cuda4dnn/csl/cublas.hpp
Copy file name to clipboard
+140Lines changed: 140 additions & 0 deletions
@@ -1160,6 +1160,11 @@ CV__DNN_INLINE_NS_BEGIN
         static Ptr<GemmLayer> create(const LayerParams& params);
     };
 
+    class CV_EXPORTS MatMulLayer : public Layer {
+     public:
+        static Ptr<MatMulLayer> create(const LayerParams &params);
+    };
+
     class CV_EXPORTS ExpandLayer : public Layer
     {
     public:
 
@@ -5,6 +5,8 @@
 #include "perf_precomp.hpp"
 #include <opencv2/dnn/shape_utils.hpp>
 
+#include <numeric>
+
 namespace opencv_test {
 
 struct GemmParam_t {
@@ -71,6 +73,18 @@ static const GemmParam_t test_gemm_configs[] = {
 */
 };
 
+static const GemmParam_t test_matmul_configs[] = {
+    // vision transformer cases
+    { {12, 197, 197}, {12, 197, 64} },
+    { {12, 197, 64 }, {12, 64, 197} },
+    { {12, 50, 64}, {12, 64, 50} },
+    { {12, 50, 50}, {12, 50, 64} },
+    { {16, 197, 197}, {16, 197, 64} },
+    { {16, 197, 64 }, {16, 64, 197} },
+    { {16, 50, 64}, {16, 64, 50} },
+    { {16, 50, 50}, {16, 50, 64} },
+};
+
 struct GemmParamId
 {
     enum {
@@ -88,6 +102,21 @@ struct GemmParamId
     }
 };
 
+struct MatMulParamId {
+    enum {
+        MATMUL_0 = 0,
+        MATMUL_LAST = sizeof(test_matmul_configs) / sizeof(test_matmul_configs[0])
+    };
+    int val_;
+    MatMulParamId(int val = 0) : val_(val) {}
+    operator int() const { return val_; }
+    static ::testing::internal::ParamGenerator<MatMulParamId> all() {
+        enum { NUM = (int)MATMUL_LAST };
+        MatMulParamId v_[NUM]; for (int i = 0; i < NUM; i++) { v_[i] = MatMulParamId(i); }
+        return ::testing::ValuesIn(v_, v_ + NUM);
+    }
+};
+
 static inline void PrintTo(const GemmParamId& v, std::ostream* os)
 {
     CV_Assert((int)v >= 0); CV_Assert((int)v < GemmParamId::GEMM_LAST);
@@ -138,7 +167,7 @@ PERF_TEST_P_(Gemm, gemm)
     Mat A(static_cast<int>(a_shape.size()), a_shape.data(), CV_32F);
     randu(A, -1.0f, 1.0f);
     Mat B(static_cast<int>(b_shape.size()), b_shape.data(), CV_32F);
-    randu(A, -1.0f, 1.0f);
+    randu(B, -1.0f, 1.0f);
 
     LayerParams lp;
     lp.type = "Gemm";
@@ -197,7 +226,7 @@ PERF_TEST_P_(Gemm, innerproduct)
     Mat A(static_cast<int>(a_shape.size()), a_shape.data(), CV_32F);
     randu(A, -1.0f, 1.0f);
     Mat B(static_cast<int>(b_shape.size()), b_shape.data(), CV_32F);
-    randu(A, -1.0f, 1.0f);
+    randu(B, -1.0f, 1.0f);
 
     LayerParams lp;
     lp.type = "InnerProduct";
@@ -241,9 +270,146 @@ PERF_TEST_P_(Gemm, innerproduct)
     SANITY_CHECK_NOTHING();
 }
 
+static inline void PrintTo(const MatMulParamId& v, std::ostream* os)
+{
+    CV_Assert((int)v >= 0); CV_Assert((int)v < MatMulParamId::MATMUL_LAST);
+    const GemmParam_t& p = test_matmul_configs[(int)v];
+
+    auto print_shape = [os](const std::vector<int>& shape, const std::string tag) {
+        if (shape.empty()) {
+            return ;
+        }
+
+        *os << tag << "=[";
+        for (size_t i = 0; i < shape.size(); ++i) {
+            if (i == shape.size() - 1) {
+                *os << shape[i] << "]";
+                break;
+            }
+            *os << shape[i] << ", ";
+        }
+    };
+
+    print_shape(p.a_shape, "A");
+    print_shape(p.b_shape, ", B");
+    print_shape(p.c_shape, ", C");
+    *os << ", trans_a=" << p.trans_a << ", trans_b=" << p.trans_b;
+}
+
+using MatMulTestParam_t = tuple<MatMulParamId, tuple<Backend, Target>>;
+using MatMul = TestBaseWithParam<MatMulTestParam_t>;
+
+PERF_TEST_P_(MatMul, matmul)
+{
+    int test_id = (int)get<0>(GetParam());
+    ASSERT_GE(test_id, 0); ASSERT_LT(test_id, MatMulParamId::MATMUL_LAST);
+    const GemmParam_t& params = test_matmul_configs[test_id];
+    auto a_shape = params.a_shape;
+    auto b_shape = params.b_shape;
+    auto trans_a = params.trans_a;
+    auto trans_b = params.trans_b;
+    float alpha = 1.f;
+    float beta = 1.f;
+
+    Backend backend_id = get<0>(get<1>(GetParam()));
+    Target target_id = get<1>(get<1>(GetParam()));
+
+    Mat A(a_shape, CV_32F);
+    randu(A, -1.0f, 1.0f);
+    Mat B(b_shape, CV_32F);
+    randu(B, -1.0f, 1.0f);
+
+    LayerParams lp;
+    lp.type = "MatMul";
+    lp.name = "testLayer";
+    lp.set("transA", trans_a);
+    lp.set("transB", trans_b);
+    lp.set("alpha", alpha);
+    lp.set("beta", beta);
+    lp.blobs.push_back(B);
+
+    Net net;
+    net.addLayerToPrev(lp.name, lp.type, lp);
+    net.setPreferableBackend(backend_id);
+    net.setPreferableTarget(target_id);
+
+    // warmup
+    {
+        std::vector<std::string> input_names{"A"};
+        net.setInputsNames(input_names);
+        net.setInput(A, input_names[0]);
+        Mat out = net.forward();
+    }
+
+    TEST_CYCLE()
+    {
+        Mat res = net.forward();
+    }
+
+    SANITY_CHECK_NOTHING();
+}
+
+PERF_TEST_P_(MatMul, innerproduct)
+{
+    int test_id = (int)get<0>(GetParam());
+    ASSERT_GE(test_id, 0); ASSERT_LT(test_id, MatMulParamId::MATMUL_LAST);
+    const GemmParam_t& params = test_matmul_configs[test_id];
+    auto a_shape = params.a_shape;
+    auto b_shape = params.b_shape;
+
+    Backend backend_id = get<0>(get<1>(GetParam()));
+    Target target_id = get<1>(get<1>(GetParam()));
+
+    Mat A(a_shape, CV_32F);
+    randu(A, -1.0f, 1.0f);
+    Mat B(b_shape, CV_32F);
+    randu(B, -1.0f, 1.0f);
+
+    LayerParams lp;
+    lp.type = "InnerProduct";
+    lp.name = "testLayer";
+    lp.set("axis", (int)(a_shape.size() - 1));
+    lp.set("bias_term", false);
+
+    // pre-transpose
+    std::vector<int> order(b_shape.size());
+    std::iota(order.begin(), order.end(), 0);
+    std::swap(order.back(), order[b_shape.size() - 2]);
+    Mat B_transposed;
+    transposeND(B, order, B_transposed);
+    lp.blobs.push_back(B_transposed);
+    lp.set("num_output", int(B_transposed.total(0, b_shape.size() - 1)));
+    lp.set("is_matmul", true);
+
+    Net net;
+    net.addLayerToPrev(lp.name, lp.type, lp);
+    net.setPreferableBackend(backend_id);
+    net.setPreferableTarget(target_id);
+
+    // warmup
+    {
+        std::vector<std::string> input_names{"A"};
+        net.setInputsNames(input_names);
+        net.setInput(A, input_names[0]);
+        Mat out = net.forward();
+    }
+
+    TEST_CYCLE()
+    {
+        Mat res = net.forward();
+    }
+
+    SANITY_CHECK_NOTHING();
+}
+
 INSTANTIATE_TEST_CASE_P(/**/, Gemm, Combine(
     GemmParamId::all(),
     dnnBackendsAndTargets(false, false)  // defined in ../test/test_common.hpp
 ));
 
+INSTANTIATE_TEST_CASE_P(/**/, MatMul, Combine(
+    MatMulParamId::all(),
+    dnnBackendsAndTargets(false, false)  // defined in ../test/test_common.hpp
+));
+
 } // namespace
@@ -8,6 +8,7 @@
 #include "error.hpp"
 #include "stream.hpp"
 #include "pointer.hpp"
+#include "memory.hpp"
 
 #include <opencv2/core.hpp>
 
@@ -363,6 +364,145 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace csl { namespace cu
         );
     }
 
+    /** @brief Strided batched GEMM for colummn-major matrices
+     *
+     * \f$ C_i = \alpha A_i B_i + \beta C_i \f$ for a stack of matrices A, B and C indexed by i
+     *
+     * @tparam          T           matrix element type (must be `half` or `float`)
+     *
+     * @param           handle      valid cuBLAS Handle
+     * @param           trans_a     use transposed matrix of A_i for computation
+     * @param           trans_b     use transposed matrix of B_i for computation
+     * @param           M           number of rows in C
+     * @param           N           number of columns in C
+     * @param           K           common dimension of A (or trans A) and B (or trans B)
+     * @param           alpha       scale factor for A B
+     * @param[in]       A           pointer to stack of column-major matrices A in device memory
+     * @param           lda         leading dimension of matrix A
+     * @param           A_offsets   offsets to get A slices
+     * @param[in]       B           pointer to stack of column-major matrices B in device memory
+     * @param           ldb         leading dimension of matrix B
+     * @param           B_offsets   offsets to get B slices
+     * @param           beta        scale factor for C
+     * @param[in,out]   C           pointer to stack of column-major matrices C in device memory
+     * @param           ldc         leading dimension of matrix C
+     * @param           C_offsets   offsets to get C slices
+     * @param           batchCount  number of matrices in the batch
+     *
+     * Exception Guarantee: Basic
+     */
+    template <class T>
+    void gemmBatched(const Handle &handle,
+                     bool trans_a, bool trans_b,
+                     std::size_t M, std::size_t N, std::size_t K,
+                     T alpha,
+                     const DevicePtr<const T> A, std::size_t lda, std::vector<std::size_t> A_offsets,
+                     const DevicePtr<const T> B, std::size_t ldb, std::vector<std::size_t> B_offsets,
+                     T beta,
+                     const DevicePtr<T> C, std::size_t ldc, std::vector<std::size_t> C_offsets,
+                     std::size_t batchCount);
+
+    template <> inline
+    void gemmBatched<half>(const Handle &handle,
+                           bool trans_a, bool trans_b,
+                           std::size_t M, std::size_t N, std::size_t K,
+                           half alpha,
+                           const DevicePtr<const half> A, std::size_t lda, std::vector<std::size_t> A_offsets,
+                           const DevicePtr<const half> B, std::size_t ldb, std::vector<std::size_t> B_offsets,
+                           half beta,
+                           const DevicePtr<half> C, std::size_t ldc, std::vector<std::size_t> C_offsets,
+                           std::size_t batchCount) {
+        CV_Assert(handle);
+
+        const auto opa = trans_a ? CUBLAS_OP_T : CUBLAS_OP_N,
+                   opb = trans_b ? CUBLAS_OP_T : CUBLAS_OP_N;
+        const auto iM = static_cast<int>(M),
+                   iN = static_cast<int>(N),
+                   iK = static_cast<int>(K),
+                   ilda = static_cast<int>(lda),
+                   ildb = static_cast<int>(ldb),
+                   ildc = static_cast<int>(ldc);
+
+        const auto batch_count = static_cast<int>(batchCount);
+
+        AutoBuffer<half> buffer(3 * batch_count);
+        auto A_slices = (half**)(buffer.data());
+        auto B_slices = A_slices + batch_count;
+        auto C_slices = B_slices + batch_count;
+        // collect A, B and C slices
+        for (int i = 0; i < batch_count; i++) {
+            A_slices[i] = (half*)(A.get()) + A_offsets[i];
+            B_slices[i] = (half*)(B.get()) + B_offsets[i];
+            C_slices[i] = (half*)(C.get()) + C_offsets[i];
+        }
+
+        const half **dev_A_slices = 0, **dev_B_slices = 0;
+        half **dev_C_slices = 0;
+        cudaMalloc((void**)&dev_A_slices, batch_count * sizeof(half*));
+        cudaMalloc((void**)&dev_B_slices, batch_count * sizeof(half*));
+        cudaMalloc((void**)&dev_C_slices, batch_count * sizeof(half*));
+        cudaMemcpy(dev_A_slices, A_slices, batch_count * sizeof(half*), cudaMemcpyHostToDevice);
+        cudaMemcpy(dev_B_slices, B_slices, batch_count * sizeof(half*), cudaMemcpyHostToDevice);
+        cudaMemcpy(dev_C_slices, C_slices, batch_count * sizeof(half*), cudaMemcpyHostToDevice);
+
+        CUDA4DNN_CHECK_CUBLAS(cublasHgemmBatched(handle.get(), opa, opb, iM, iN, iK, &alpha, dev_A_slices, ilda, dev_B_slices, ildb, &beta, dev_C_slices, ildc, batch_count));
+
+        cudaFree(dev_A_slices);
+        cudaFree(dev_B_slices);
+        cudaFree(dev_C_slices);
+    }
+
+    template <> inline
+    void gemmBatched<float>(const Handle &handle,
+                           bool trans_a, bool trans_b,
+                           std::size_t M, std::size_t N, std::size_t K,
+                           float alpha,
+                           const DevicePtr<const float> A, std::size_t lda, std::vector<std::size_t> A_offsets,
+                           const DevicePtr<const float> B, std::size_t ldb, std::vector<std::size_t> B_offsets,
+                           float beta,
+                           const DevicePtr<float> C, std::size_t ldc, std::vector<std::size_t> C_offsets,
+                           std::size_t batchCount) {
+        CV_Assert(handle);
+
+        const auto opa = trans_a ? CUBLAS_OP_T : CUBLAS_OP_N,
+                   opb = trans_b ? CUBLAS_OP_T : CUBLAS_OP_N;
+        const auto iM = static_cast<int>(M),
+                   iN = static_cast<int>(N),
+                   iK = static_cast<int>(K),
+                   ilda = static_cast<int>(lda),
+                   ildb = static_cast<int>(ldb),
+                   ildc = static_cast<int>(ldc);
+
+        const auto batch_count = static_cast<int>(batchCount);
+
+        AutoBuffer<float> buffer(3 * batch_count);
+        auto A_slices = (float**)(buffer.data());
+        auto B_slices = A_slices + batch_count;
+        auto C_slices = B_slices + batch_count;
+        // collect A, B and C slices
+        for (int i = 0; i < batch_count; i++) {
+            A_slices[i] = (float*)(A.get()) + A_offsets[i];
+            B_slices[i] = (float*)(B.get()) + B_offsets[i];
+            C_slices[i] = (float*)(C.get()) + C_offsets[i];
+        }
+
+        const float **dev_A_slices = 0, **dev_B_slices = 0;
+        float **dev_C_slices = 0;
+        cudaMalloc((void**)&dev_A_slices, batch_count * sizeof(float*));
+        cudaMalloc((void**)&dev_B_slices, batch_count * sizeof(float*));
+        cudaMalloc((void**)&dev_C_slices, batch_count * sizeof(float*));
+        cudaMemcpy(dev_A_slices, A_slices, batch_count * sizeof(float*), cudaMemcpyHostToDevice);
+        cudaMemcpy(dev_B_slices, B_slices, batch_count * sizeof(float*), cudaMemcpyHostToDevice);
+        cudaMemcpy(dev_C_slices, C_slices, batch_count * sizeof(float*), cudaMemcpyHostToDevice);
+
+        // cuBLAS is column-major
+        CUDA4DNN_CHECK_CUBLAS(cublasSgemmBatched(handle.get(), opa, opb, iM, iN, iK, &alpha, dev_A_slices, ilda, dev_B_slices, ildb, &beta, dev_C_slices, ildc, batch_count));
+
+        cudaFree(dev_A_slices);
+        cudaFree(dev_B_slices);
+        cudaFree(dev_C_slices);
+    }
+
 }}}}} /* namespace cv::dnn::cuda4dnn::csl::cublas */
 
 #endif /* OPENCV_DNN_SRC_CUDA4DNN_CSL_CUBLAS_HPP */