opencv
diff --git a/‎modules/dnn/src/cuda/mvn.cu
Copy file name to clipboard
+62-7Lines changed: 62 additions & 7 deletions b/‎modules/dnn/src/cuda/mvn.cu
Copy file name to clipboard
+62-7Lines changed: 62 additions & 7 deletions
diff --git a/‎modules/dnn/src/cuda4dnn/kernels/mvn.hpp
Copy file name to clipboard
+7-1Lines changed: 7 additions & 1 deletion b/‎modules/dnn/src/cuda4dnn/kernels/mvn.hpp
Copy file name to clipboard
+7-1Lines changed: 7 additions & 1 deletion
diff --git a/‎modules/dnn/src/cuda4dnn/primitives/layer_norm.hpp
Copy file name to clipboard
+93Lines changed: 93 additions & 0 deletions b/‎modules/dnn/src/cuda4dnn/primitives/layer_norm.hpp
Copy file name to clipboard
+93Lines changed: 93 additions & 0 deletions
@@ -68,15 +68,36 @@ namespace raw {
     }
 
     template <class T>
-    __global__ void normalize_mean_variance_channelwise(Span<T> output, View<T> input, View<T> scale, View<T> bias, View<float> means, View<float> stdev, size_type inner_size, size_type C) {
+    __global__ void normalize_mean_variance_channelwise(Span<T> output, View<T> input, View<T> scale, View<T> bias, View<float> means, View<float> inv_stddev, size_type inner_size, size_type C) {
         for (auto idx : grid_stride_range(output.size())) {
             const index_type outer_idx = idx / inner_size;
             const index_type c = outer_idx % C;
-            auto s = static_cast<float>(scale[c]) * stdev[outer_idx];
+            auto s = static_cast<float>(scale[c]) * inv_stddev[outer_idx];
             auto b = static_cast<float>(bias[c]);
             output[idx] = (static_cast<float>(input[idx]) - means[outer_idx]) * s + b;
         }
     }
+
+    template <class T>
+    __global__ void normalize_mean_variance_layernorm(Span<T> output, View<T> input, View<T> scale, View<float> means, View<float> inv_stddev, size_type inner_size) {
+        for (auto idx : grid_stride_range(output.size())) {
+            const index_type outer_idx = idx / inner_size;
+            const index_type inner_idx = idx % inner_size;
+            auto s = static_cast<float>(scale[inner_idx]) * inv_stddev[outer_idx];
+            output[idx] = (static_cast<float>(input[idx]) - means[outer_idx]) * s;
+        }
+    }
+
+    template <class T>
+    __global__ void normalize_mean_variance_layernorm_with_bias(Span<T> output, View<T> input, View<T> scale, View<T> bias, View<float> means, View<float> inv_stddev, size_type inner_size) {
+        for (auto idx : grid_stride_range(output.size())) {
+            const index_type outer_idx = idx / inner_size;
+            const index_type inner_idx = idx % inner_size;
+            auto s = static_cast<float>(scale[inner_idx]) * inv_stddev[outer_idx];
+            auto b = static_cast<float>(bias[inner_idx]);
+            output[idx] = (static_cast<float>(input[idx]) - means[outer_idx]) * s + b;
+        }
+    }
 }
 
 template <class T>
@@ -154,20 +175,54 @@ template void normalize_mean_variance(const Stream&, Span<__half>, View<__half>,
 template void normalize_mean_variance(const Stream&, Span<float>, View<float>, View<float>, View<float>, std::size_t);
 
 template <class T>
-void normalize_mean_variance_channelwise(const Stream& stream, Span<T> output, View<T> input, View<T> scale, View<T> bias, View<float> means, View<float> stdev, std::size_t inner_size, std::size_t C)
+void normalize_mean_variance_channelwise(const Stream& stream, Span<T> output, View<T> input, View<T> scale, View<T> bias, View<float> means, View<float> inv_stddev, std::size_t inner_size, std::size_t C)
 {
     CV_Assert(input.size() == output.size());
     CV_Assert(input.size() / inner_size == means.size());
-    CV_Assert(means.size() == stdev.size());
+    CV_Assert(means.size() == inv_stddev.size());
 
     auto kernel = raw::normalize_mean_variance_channelwise<T>;
     auto policy = make_policy(kernel, output.size(), 0, stream);
-    launch_kernel(kernel, policy, output, input, scale, bias, means, stdev, inner_size, C);
+    launch_kernel(kernel, policy, output, input, scale, bias, means, inv_stddev, inner_size, C);
+}
+
+#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 530)
+template void normalize_mean_variance_channelwise(const Stream&, Span<__half> /*output*/, View<__half> /*input*/, View<__half> /*scale*/, View<__half> /*bias*/, View<float> /*means*/, View<float> /*inv_stddev*/, std::size_t, std::size_t);
+#endif
+template void normalize_mean_variance_channelwise(const Stream&, Span<float> /*output*/, View<float> /*input*/, View<float> /*scale*/, View<float> /*bias*/, View<float> /*means*/, View<float> /*inv_stddev*/, std::size_t, std::size_t);
+
+template <class T>
+void normalize_mean_variance_layernorm(const Stream& stream, Span<T> output, View<T> input, View<T> scale, View<float> means, View<float> inv_stddev, std::size_t inner_size)
+{
+    CV_Assert(input.size() == output.size());
+    CV_Assert(input.size() / inner_size == means.size());
+    CV_Assert(means.size() == inv_stddev.size());
+
+    auto kernel = raw::normalize_mean_variance_layernorm<T>;
+    auto policy = make_policy(kernel, output.size(), 0, stream);
+    launch_kernel(kernel, policy, output, input, scale, means, inv_stddev, inner_size);
+}
+
+#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 530)
+template void normalize_mean_variance_layernorm(const Stream&, Span<__half> /*output*/, View<__half> /*input*/, View<__half> /*scale*/, View<float> /*means*/, View<float> /*inv_stddev*/, std::size_t);
+#endif
+template void normalize_mean_variance_layernorm(const Stream&, Span<float> /*output*/, View<float> /*input*/, View<float> /*scale*/, View<float> /*means*/, View<float> /*inv_stddev*/, std::size_t);
+
+template <class T>
+void normalize_mean_variance_layernorm(const Stream& stream, Span<T> output, View<T> input, View<T> scale, View<T> bias, View<float> means, View<float> inv_stddev, std::size_t inner_size)
+{
+    CV_Assert(input.size() == output.size());
+    CV_Assert(input.size() / inner_size == means.size());
+    CV_Assert(means.size() == inv_stddev.size());
+
+    auto kernel = raw::normalize_mean_variance_layernorm_with_bias<T>;
+    auto policy = make_policy(kernel, output.size(), 0, stream);
+    launch_kernel(kernel, policy, output, input, scale, bias, means, inv_stddev, inner_size);
 }
 
 #if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 530)
-template void normalize_mean_variance_channelwise(const Stream&, Span<__half> /*output*/, View<__half> /*input*/, View<__half> /*scale*/, View<__half> /*bias*/, View<float> /*means*/, View<float> /*stdev*/, std::size_t, std::size_t);
+template void normalize_mean_variance_layernorm(const Stream&, Span<__half> /*output*/, View<__half> /*input*/, View<__half> /*scale*/, View<__half> /*bias*/, View<float> /*means*/, View<float> /*inv_stddev*/, std::size_t);
 #endif
-template void normalize_mean_variance_channelwise(const Stream&, Span<float> /*output*/, View<float> /*input*/, View<float> /*scale*/, View<float> /*bias*/, View<float> /*means*/, View<float> /*stdev*/, std::size_t, std::size_t);
+template void normalize_mean_variance_layernorm(const Stream&, Span<float> /*output*/, View<float> /*input*/, View<float> /*scale*/, View<float> /*bias*/, View<float> /*means*/, View<float> /*inv_stddev*/, std::size_t);
 
 }}}} /* namespace cv::dnn::cuda4dnn::kernels */
@@ -27,7 +27,13 @@ template <class T>
 void normalize_mean_variance(const csl::Stream& stream, csl::Span<T> output, csl::View<T> input, csl::View<float> means, csl::View<float> scale, std::size_t inner_size);
 
 template <class T>
-void normalize_mean_variance_channelwise(const csl::Stream &stream, csl::Span<T> output, csl::View<T> input, csl::View<T> scale, csl::View<T> bias, csl::View<float> means, csl::View<float> stdev, std::size_t inner_size, std::size_t C);
+void normalize_mean_variance_channelwise(const csl::Stream &stream, csl::Span<T> output, csl::View<T> input, csl::View<T> scale, csl::View<T> bias, csl::View<float> means, csl::View<float> inv_stddev, std::size_t inner_size, std::size_t C);
+
+template <class T>
+void normalize_mean_variance_layernorm(const csl::Stream &stream, csl::Span<T> output, csl::View<T> input, csl::View<T> scale, csl::View<float> means, csl::View<float> inv_stddev, std::size_t inner_size);
+
+template <class T>
+void normalize_mean_variance_layernorm(const csl::Stream &stream, csl::Span<T> output, csl::View<T> input, csl::View<T> scale, csl::View<T> bias, csl::View<float> means, csl::View<float> inv_stddev, std::size_t inner_size);
 
 }}}} /* namespace cv::dnn::cuda4dnn::kernels */
 
 
@@ -0,0 +1,93 @@
+// This file is part of OpenCV project.
+// It is subject to the license terms in the LICENSE file found in the top-level directory
+// of this distribution and at https://opencv.org/license.html.
+
+#ifndef OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_LAYER_NORM_HPP
+#define OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_LAYER_NORM_HPP
+
+#include "../../op_cuda.hpp"
+
+#include "../csl/stream.hpp"
+#include "../csl/span.hpp"
+#include "../csl/tensor.hpp"
+#include "../csl/workspace.hpp"
+
+#include "../kernels/fill_copy.hpp"
+#include "../kernels/mvn.hpp"
+
+#include <opencv2/core.hpp>
+
+#include <cstddef>
+#include <vector>
+#include <utility>
+
+namespace cv { namespace dnn { namespace cuda4dnn {
+
+    template <class T>
+    class LayerNormOp final : public CUDABackendNode {
+     public:
+        using wrapper_type = GetCUDABackendWrapperType<T>;
+
+        LayerNormOp(csl::Stream stream_, int normalized_axis, float epsilon_, size_t loops)
+            : stream(std::move(stream_)), epsilon(epsilon_) {
+            CV_CheckGE(normalized_axis, 0, "LayerNorm/CUDA: axis needs to be normalized");
+            axis = static_cast<size_t>(normalized_axis);
+
+            csl::WorkspaceBuilder builder;
+            builder.require<float>(loops);
+            builder.require<float>(loops);
+            scratch_mem_in_bytes = builder.required_workspace_size();
+        }
+
+        void forward(const std::vector<cv::Ptr<BackendWrapper>>& inputs,
+                     const std::vector<cv::Ptr<BackendWrapper>>& outputs,
+                     csl::Workspace& workspace) override {
+            auto input_wrapper = inputs[0].dynamicCast<wrapper_type>();
+            auto scale_wrapper = inputs[1].dynamicCast<wrapper_type>();
+
+            auto input = input_wrapper->getView();
+            auto scale = scale_wrapper->getView();
+
+            auto output_wrapper = outputs[0].dynamicCast<wrapper_type>();
+            auto output = output_wrapper->getSpan();
+
+            auto loops = input.size_range(0, axis);
+            auto norm_size = input.size_range(axis, input.rank());
+            if (norm_size == 1) {
+                kernels::fill<T>(stream, output, 0.f);
+                return;
+            } else {
+                auto ws_allocator = csl::WorkspaceAllocator(workspace);
+
+                auto mean = ws_allocator.get_span<float>(loops);
+                kernels::fill<float>(stream, mean, 0.f);
+
+                auto inv_stddev = ws_allocator.get_span<float>(loops);
+                kernels::fill<float>(stream, inv_stddev, 0.f);
+
+                kernels::reduce_mean_sqr_sum<T>(stream, mean, inv_stddev, input, norm_size);
+                kernels::compute_normalization_scale(stream, inv_stddev, mean, inv_stddev, norm_size, epsilon);
+                if (inputs.size() == 3) {
+                    auto bias_wrapper = inputs[2].dynamicCast<wrapper_type>();
+                    auto bias = bias_wrapper->getView();
+                    kernels::normalize_mean_variance_layernorm<T>(stream, output, input, scale, bias, mean, inv_stddev, norm_size);
+                } else {
+                    kernels::normalize_mean_variance_layernorm<T>(stream, output, input, scale, mean, inv_stddev, norm_size);
+                }
+            }
+        }
+
+        std::size_t get_workspace_memory_in_bytes() const noexcept override { return scratch_mem_in_bytes; }
+
+     private:
+        csl::Stream stream;
+
+        float epsilon;
+        size_t axis;
+
+        std::size_t scratch_mem_in_bytes;
+    };
+
+}}} // cv::dnn::cuda4dnn
+
+#endif // OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_LAYER_NORM_HPP