add functions addSquare and elementWiseSqrt

Author: joeljonsson

src/numerics/miscCuda.cu

@@ -43,6 +43,32 @@ __global__ void elementWiseDiv(const T* numerator,
 }
 
 
+/**
+ * Add the square of one vector to another. For 0 <= idx < size:
+ *      in1[idx] += in2[idx] * in2[idx];
+ */
+__global__ void addSquare(float* in1,
+                          const float* in2,
+                          const size_t size) {
+
+    for(size_t idx = blockIdx.x * blockDim.x + threadIdx.x; idx < size; idx += blockDim.x * gridDim.x) {
+        in1[idx] += in2[idx] * in2[idx];
+    }
+}
+
+
+/**
+ * Take the square root of each element in the input vector
+ */
+__global__ void elementWiseSqrt(float* __restrict__ input,
+                                const size_t size) {
+
+    for(size_t idx = blockIdx.x * blockDim.x + threadIdx.x; idx < size; idx += blockDim.x * gridDim.x) {
+        input[idx] = sqrtf(input[idx]);
+    }
+}
+
+
 template<typename T>
 __global__ void copyKernel(const T* in, T* out, const size_t size){
 

src/numerics/miscCuda.hpp

@@ -85,6 +85,12 @@ __global__ void elementWiseMult(T* in1, const T* in2, const size_t size);
 template<typename T, typename S>
 __global__ void elementWiseDiv(const T* numerator, const S* denominator, S* out, const size_t size);
 
+
+__global__ void addSquare(float* in1, const float* in2, size_t size);
+
+__global__ void elementWiseSqrt(float* input, size_t size);
+
+
 template<typename T>
 __global__ void copyKernel(const T* in, T* out, const size_t size);