From dec0237353617bb6957bcb8b4a25c6587935354c Mon Sep 17 00:00:00 2001
From: Tom Deakin <tom.deakin@bristol.ac.uk>
Date: Tue, 26 Apr 2016 16:03:28 +0100
Subject: [PATCH] Add mul kernel

---
 src/CUDAStream.cu | 14 +++++++++++++-
 1 file changed, 13 insertions(+), 1 deletion(-)

diff --git a/src/CUDAStream.cu b/src/CUDAStream.cu
index 5629b40..15b1b41 100644
--- a/src/CUDAStream.cu
+++ b/src/CUDAStream.cu
@@ -65,10 +65,21 @@ void CUDAStream<T>::copy()
   check_error();
 }
 
+template <typename T>
+__global__ void mul_kernel(T * b, const T * c)
+{
+  const T scalar = 3.0;
+  const int i = blockDim.x * blockIdx.x + threadIdx.x;
+  b[i] = scalar * c[i];
+}
+
 template <class T>
 void CUDAStream<T>::mul()
 {
-  return;
+  mul_kernel<<<1024, 1024>>>(d_b, d_c);
+  check_error();
+  cudaDeviceSynchronize();
+  check_error();
 }
 
 template <class T>
@@ -85,3 +96,4 @@ void CUDAStream<T>::triad()
 
 template class CUDAStream<float>;
 template class CUDAStream<double>;
+