From da1526aa49530df7faa588f0335642e1a2fbf90f Mon Sep 17 00:00:00 2001
From: Andrey Morozov <no@email>
Date: Sat, 17 Jul 2010 13:50:30 +0000
Subject: [PATCH] Added implementation SetTo() without mask

---
 modules/gpu/cuda/mat_operators.cu | 24 ++++++++++++++++++++++--
 1 file changed, 22 insertions(+), 2 deletions(-)

diff --git a/modules/gpu/cuda/mat_operators.cu b/modules/gpu/cuda/mat_operators.cu
index 57eb9bb..5d69759 100644
--- a/modules/gpu/cuda/mat_operators.cu
+++ b/modules/gpu/cuda/mat_operators.cu
@@ -40,6 +40,7 @@
 //
 //M*/
 
+#include <stddef.h>
 #include "cuda_shared.hpp"
 #include "cuda_runtime.h"
 
@@ -47,11 +48,30 @@ __constant__ float scalar_d[4];
 
 namespace mat_operators
 {
+
+    template <typename T, int channels, int count = channels>
+    struct unroll
+    {
+        __device__ static void unroll_set(T * mat, size_t i)
+        {
+            mat[i] = static_cast<T>(scalar_d[i % channels]);
+            unroll<T, channels, count - 1>::unroll_set(mat, i+1);
+        }
+    };
+
+    template <typename T, int channels>
+    struct unroll<T,channels,0>
+    {
+        __device__ static void unroll_set(T * , size_t){}
+    };
+
+
+
     template <typename T, int channels>
     __global__ void kernel_set_to_without_mask(T * mat)
     {
-        int i = blockIdx.x * blockDim.x + threadIdx.x;
-        mat[i * sizeof(T)] = static_cast<T>(scalar_d[i % channels]);
+        size_t i = (blockIdx.x * blockDim.x + threadIdx.x) * sizeof(T);
+        unroll<T, channels>::unroll_set(mat, i);
     }
 }
 
-- 
2.7.4