diff --git a/ACCStream.cpp b/ACCStream.cpp
index bd49663..ccc942a 100644
--- a/ACCStream.cpp
+++ b/ACCStream.cpp
@@ -36,13 +36,19 @@ ACCStream<T>::~ACCStream()
 }
 
 template <class T>
-void ACCStream<T>::write_arrays(const std::vector<T>& h_a, const std::vector<T>& h_b, const std::vector<T>& h_c)
+void ACCStream<T>::init_arrays(T initA, T initB, T initC)
 {
-  T *a = this->a;
-  T *b = this->b;
-  T *c = this->c;
-  #pragma acc update device(a[0:array_size], b[0:array_size], c[0:array_size])
-  {}
+  unsigned int array_size = this->array_size;
+  T * restrict a = this->a;
+  T * restrict b = this->b;
+  T * restrict c = this->c;
+  #pragma acc kernels present(a[0:array_size], b[0:array_size], c[0:array_size]) wait
+  for (int i = 0; i < array_size; i++)
+  {
+    a[i] = initA;
+    b[i] = initB;
+    c[i] = initC;
+  }
 }
 
 template <class T>
diff --git a/ACCStream.h b/ACCStream.h
index 48fea55..54f947b 100644
--- a/ACCStream.h
+++ b/ACCStream.h
@@ -36,7 +36,7 @@ class ACCStream : public Stream<T>
     virtual void mul() override;
     virtual void triad() override;
 
-    virtual void write_arrays(const std::vector<T>& a, const std::vector<T>& b, const std::vector<T>& c) override;
+    virtual void init_arrays(T initA, T initB, T initC) override;
     virtual void read_arrays(std::vector<T>& a, std::vector<T>& b, std::vector<T>& c) override;
 
 
diff --git a/CUDAStream.cu b/CUDAStream.cu
index ff2ec41..08026d9 100644
--- a/CUDAStream.cu
+++ b/CUDAStream.cu
@@ -74,15 +74,22 @@ CUDAStream<T>::~CUDAStream()
   check_error();
 }
 
-template <class T>
-void CUDAStream<T>::write_arrays(const std::vector<T>& a, const std::vector<T>& b, const std::vector<T>& c)
+
+template <typename T>
+__global__ void init_kernel(T * a, T * b, T * c, T initA, T initB, T initC)
 {
-  // Copy host memory to device
-  cudaMemcpy(d_a, a.data(), a.size()*sizeof(T), cudaMemcpyHostToDevice);
+  const int i = blockDim.x * blockIdx.x + threadIdx.x;
+  a[i] = initA;
+  b[i] = initB;
+  c[i] = initC;
+}
+
+template <class T>
+void CUDAStream<T>::init_arrays(T initA, T initB, T initC)
+{
+  init_kernel<<<array_size/TBSIZE, TBSIZE>>>(d_a, d_b, d_c, initA, initB, initC);
   check_error();
-  cudaMemcpy(d_b, b.data(), b.size()*sizeof(T), cudaMemcpyHostToDevice);
-  check_error();
-  cudaMemcpy(d_c, c.data(), c.size()*sizeof(T), cudaMemcpyHostToDevice);
+  cudaDeviceSynchronize();
   check_error();
 }
 
diff --git a/CUDAStream.h b/CUDAStream.h
index 6904a86..912721e 100644
--- a/CUDAStream.h
+++ b/CUDAStream.h
@@ -37,7 +37,7 @@ class CUDAStream : public Stream<T>
     virtual void mul() override;
     virtual void triad() override;
 
-    virtual void write_arrays(const std::vector<T>& a, const std::vector<T>& b, const std::vector<T>& c) override;
+    virtual void init_arrays(T initA, T initB, T initC) override;
     virtual void read_arrays(std::vector<T>& a, std::vector<T>& b, std::vector<T>& c) override;
 
 };
diff --git a/HIPStream.cu b/HIPStream.cu
index 34ecfb6..8c02348 100644
--- a/HIPStream.cu
+++ b/HIPStream.cu
@@ -74,15 +74,21 @@ HIPStream<T>::~HIPStream()
   check_error();
 }
 
-template <class T>
-void HIPStream<T>::write_arrays(const std::vector<T>& a, const std::vector<T>& b, const std::vector<T>& c)
+template <typename T>
+__global__ void init_kernel(hipLaunchParm lp, T * a, T * b, T * c, T initA, T initB, T initC)
 {
-  // Copy host memory to device
-  hipMemcpy(d_a, a.data(), a.size()*sizeof(T), hipMemcpyHostToDevice);
+  const int i = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
+  a[i] = initA;
+  b[i] = initB;
+  c[i] = initC;
+}
+
+template <class T>
+void HIPStream<T>::init_arrays(T initA, T initB, T initC)
+{
+  hipLaunchKernel(HIP_KERNEL_NAME(init_kernel), dim3(array_size/TBSIZE), dim3(TBSIZE), 0, 0, d_a, d_b, d_c, initA, initB, initC);
   check_error();
-  hipMemcpy(d_b, b.data(), b.size()*sizeof(T), hipMemcpyHostToDevice);
-  check_error();
-  hipMemcpy(d_c, c.data(), c.size()*sizeof(T), hipMemcpyHostToDevice);
+  hipDeviceSynchronize();
   check_error();
 }
 
diff --git a/HIPStream.h b/HIPStream.h
index 9015e35..392080a 100644
--- a/HIPStream.h
+++ b/HIPStream.h
@@ -37,7 +37,7 @@ class HIPStream : public Stream<T>
     virtual void mul() override;
     virtual void triad() override;
 
-    virtual void write_arrays(const std::vector<T>& a, const std::vector<T>& b, const std::vector<T>& c) override;
+    virtual void init_arrays(T initA, T initB, T initC) override;
     virtual void read_arrays(std::vector<T>& a, std::vector<T>& b, std::vector<T>& c) override;
 
 };
diff --git a/KOKKOSStream.cpp b/KOKKOSStream.cpp
index 94ac7ee..72b1ee5 100644
--- a/KOKKOSStream.cpp
+++ b/KOKKOSStream.cpp
@@ -34,18 +34,18 @@ KOKKOSStream<T>::~KOKKOSStream()
 }
 
 template <class T>
-void KOKKOSStream<T>::write_arrays(
-        const std::vector<T>& a, const std::vector<T>& b, const std::vector<T>& c)
+void KOKKOSStream<T>::init_arrays(T initA, T initB, T initC)
 {
-  for(int ii = 0; ii < array_size; ++ii)
+  View<double*, DEVICE> a(*d_a);
+  View<double*, DEVICE> b(*d_b);
+  View<double*, DEVICE> c(*d_c);
+  parallel_for(array_size, KOKKOS_LAMBDA (const int index)
   {
-    (*hm_a)(ii) = a[ii];
-    (*hm_b)(ii) = b[ii];
-    (*hm_c)(ii) = c[ii];
-  }
-  deep_copy(*d_a, *hm_a);
-  deep_copy(*d_b, *hm_b);
-  deep_copy(*d_c, *hm_c);
+    a[index] = initA;
+    b[index] - initB;
+    c[index] = initC;
+  });
+  Kokkos::fence();
 }
 
 template <class T>
diff --git a/KOKKOSStream.hpp b/KOKKOSStream.hpp
index d2b9665..ff7cfeb 100644
--- a/KOKKOSStream.hpp
+++ b/KOKKOSStream.hpp
@@ -48,8 +48,7 @@ class KOKKOSStream : public Stream<T>
     virtual void mul() override;
     virtual void triad() override;
 
-    virtual void write_arrays(
-            const std::vector<T>& a, const std::vector<T>& b, const std::vector<T>& c) override;
+    virtual void init_arrays(T initA, T initB, T initC) override;
     virtual void read_arrays(
             std::vector<T>& a, std::vector<T>& b, std::vector<T>& c) override;
 };
diff --git a/OCLStream.cpp b/OCLStream.cpp
index 2a1e5ee..c7e09a8 100644
--- a/OCLStream.cpp
+++ b/OCLStream.cpp
@@ -16,6 +16,18 @@ std::string kernels{R"CLC(
 
   constant TYPE scalar = startScalar;
 
+  kernel void init(
+    global TYPE * restrict a,
+    global TYPE * restrict b,
+    global TYPE * restrict c,
+    TYPE initA, TYPE initB, TYPE initC)
+  {
+    const size_t i = get_global_id(0);
+    a[i] = initA;
+    b[i] = initB;
+    c[i] = initC;
+  }
+
   kernel void copy(
     global const TYPE * restrict a,
     global TYPE * restrict c)
@@ -101,6 +113,7 @@ OCLStream<T>::OCLStream(const unsigned int ARRAY_SIZE, const int device_index)
   }
 
   // Create kernels
+  init_kernel = new cl::KernelFunctor<cl::Buffer, cl::Buffer, cl::Buffer, T, T, T>(program, "init");
   copy_kernel = new cl::KernelFunctor<cl::Buffer, cl::Buffer>(program, "copy");
   mul_kernel = new cl::KernelFunctor<cl::Buffer, cl::Buffer>(program, "mul");
   add_kernel = new cl::KernelFunctor<cl::Buffer, cl::Buffer, cl::Buffer>(program, "add");
@@ -120,12 +133,12 @@ OCLStream<T>::OCLStream(const unsigned int ARRAY_SIZE, const int device_index)
   d_a = cl::Buffer(context, CL_MEM_READ_WRITE, sizeof(T) * ARRAY_SIZE);
   d_b = cl::Buffer(context, CL_MEM_READ_WRITE, sizeof(T) * ARRAY_SIZE);
   d_c = cl::Buffer(context, CL_MEM_READ_WRITE, sizeof(T) * ARRAY_SIZE);
-
 }
 
 template <class T>
 OCLStream<T>::~OCLStream()
 {
+  delete init_kernel;
   delete copy_kernel;
   delete mul_kernel;
   delete add_kernel;
@@ -173,11 +186,13 @@ void OCLStream<T>::triad()
 }
 
 template <class T>
-void OCLStream<T>::write_arrays(const std::vector<T>& a, const std::vector<T>& b, const std::vector<T>& c)
+void OCLStream<T>::init_arrays(T initA, T initB, T initC)
 {
-  cl::copy(queue, a.begin(), a.end(), d_a);
-  cl::copy(queue, b.begin(), b.end(), d_b);
-  cl::copy(queue, c.begin(), c.end(), d_c);
+  (*init_kernel)(
+    cl::EnqueueArgs(queue, cl::NDRange(array_size)),
+    d_a, d_b, d_c, initA, initB, initC
+  );
+  queue.finish();
 }
 
 template <class T>
diff --git a/OCLStream.h b/OCLStream.h
index 54abaa3..845e144 100644
--- a/OCLStream.h
+++ b/OCLStream.h
@@ -38,6 +38,7 @@ class OCLStream : public Stream<T>
     cl::Context context;
     cl::CommandQueue queue;
 
+    cl::KernelFunctor<cl::Buffer, cl::Buffer, cl::Buffer, T, T, T> *init_kernel;
     cl::KernelFunctor<cl::Buffer, cl::Buffer> *copy_kernel;
     cl::KernelFunctor<cl::Buffer, cl::Buffer> * mul_kernel;
     cl::KernelFunctor<cl::Buffer, cl::Buffer, cl::Buffer> *add_kernel;
@@ -53,7 +54,7 @@ class OCLStream : public Stream<T>
     virtual void mul() override;
     virtual void triad() override;
 
-    virtual void write_arrays(const std::vector<T>& a, const std::vector<T>& b, const std::vector<T>& c) override;
+    virtual void init_arrays(T initA, T initB, T initC) override;
     virtual void read_arrays(std::vector<T>& a, std::vector<T>& b, std::vector<T>& c) override;
 
 };
diff --git a/OMP3Stream.cpp b/OMP3Stream.cpp
index f578c7c..b5e1bc2 100644
--- a/OMP3Stream.cpp
+++ b/OMP3Stream.cpp
@@ -26,14 +26,14 @@ OMP3Stream<T>::~OMP3Stream()
 
 
 template <class T>
-void OMP3Stream<T>::write_arrays(const std::vector<T>& h_a, const std::vector<T>& h_b, const std::vector<T>& h_c)
+void OMP3Stream<T>::init_arrays(T initA, T initB, T initC)
 {
   #pragma omp parallel for
   for (int i = 0; i < array_size; i++)
   {
-    a[i] = h_a[i];
-    b[i] = h_b[i];
-    c[i] = h_c[i];
+    a[i] = initA;
+    b[i] = initB;
+    c[i] = initC;
   }
 }
 
diff --git a/OMP3Stream.h b/OMP3Stream.h
index 0f14300..1dadc95 100644
--- a/OMP3Stream.h
+++ b/OMP3Stream.h
@@ -34,7 +34,7 @@ class OMP3Stream : public Stream<T>
     virtual void mul() override;
     virtual void triad() override;
 
-    virtual void write_arrays(const std::vector<T>& a, const std::vector<T>& b, const std::vector<T>& c) override;
+    virtual void init_arrays(T initA, T initB, T initC) override;
     virtual void read_arrays(std::vector<T>& a, std::vector<T>& b, std::vector<T>& c) override;
 
 };
diff --git a/OMP45Stream.cpp b/OMP45Stream.cpp
index 8f684e2..1722662 100644
--- a/OMP45Stream.cpp
+++ b/OMP45Stream.cpp
@@ -35,13 +35,19 @@ OMP45Stream<T>::~OMP45Stream()
 }
 
 template <class T>
-void OMP45Stream<T>::write_arrays(const std::vector<T>& h_a, const std::vector<T>& h_b, const std::vector<T>& h_c)
+void OMP45Stream<T>::init_arrays(T initA, T initB, T initC)
 {
+  unsigned int array_size = this->array_size;
   T *a = this->a;
   T *b = this->b;
   T *c = this->c;
-  #pragma omp target update to(a[0:array_size], b[0:array_size], c[0:array_size])
-  {}
+  #pragma omp target teams distribute parallel for simd map(to: a[0:array_size], b[0:array_size], c[0:array_size])
+  for (int i = 0; i < array_size; i++)
+  {
+    a[i] = initA;
+    b[i] = initB;
+    c[i] = initC;
+  }
 }
 
 template <class T>
diff --git a/OMP45Stream.h b/OMP45Stream.h
index bd812a1..d2a5aaf 100644
--- a/OMP45Stream.h
+++ b/OMP45Stream.h
@@ -37,7 +37,7 @@ class OMP45Stream : public Stream<T>
     virtual void mul() override;
     virtual void triad() override;
 
-    virtual void write_arrays(const std::vector<T>& a, const std::vector<T>& b, const std::vector<T>& c) override;
+    virtual void init_arrays(T initA, T initB, T initC) override;
     virtual void read_arrays(std::vector<T>& a, std::vector<T>& b, std::vector<T>& c) override;
 
 
diff --git a/RAJAStream.cpp b/RAJAStream.cpp
index 33687a1..d872987 100644
--- a/RAJAStream.cpp
+++ b/RAJAStream.cpp
@@ -21,12 +21,6 @@ RAJAStream<T>::RAJAStream(const unsigned int ARRAY_SIZE, const int device_index)
   d_a = new T[ARRAY_SIZE];
   d_b = new T[ARRAY_SIZE];
   d_c = new T[ARRAY_SIZE];
-  forall<policy>(index_set, [=] RAJA_DEVICE (int index)
-  {
-    d_a[index] = 0.0;
-    d_b[index] = 0.0;
-    d_c[index] = 0.0;
-  });
 #else
   cudaMallocManaged((void**)&d_a, sizeof(T)*ARRAY_SIZE, cudaMemAttachGlobal);
   cudaMallocManaged((void**)&d_b, sizeof(T)*ARRAY_SIZE, cudaMemAttachGlobal);
@@ -50,12 +44,17 @@ RAJAStream<T>::~RAJAStream()
 }
 
 template <class T>
-void RAJAStream<T>::write_arrays(
-        const std::vector<T>& a, const std::vector<T>& b, const std::vector<T>& c)
+void RAJAStream<T>::init_arrays(T initA, T initB, T initC)
 {
-  std::copy(a.begin(), a.end(), d_a);
-  std::copy(b.begin(), b.end(), d_b);
-  std::copy(c.begin(), c.end(), d_c);
+  T* a = d_a;
+  T* b = d_b;
+  T* c = d_c;
+  forall<policy>(index_set, [=] RAJA_DEVICE (int index)
+  {
+    a[index] = initA;
+    b[index] = initB;
+    c[index] = initC;
+  });
 }
 
 template <class T>
diff --git a/RAJAStream.hpp b/RAJAStream.hpp
index 454e20e..8ffa5be 100644
--- a/RAJAStream.hpp
+++ b/RAJAStream.hpp
@@ -50,8 +50,7 @@ class RAJAStream : public Stream<T>
     virtual void mul() override;
     virtual void triad() override;
 
-    virtual void write_arrays(
-            const std::vector<T>& a, const std::vector<T>& b, const std::vector<T>& c) override;
+    virtual void init_arrays(T initA, T initB, T initC) override;
     virtual void read_arrays(
             std::vector<T>& a, std::vector<T>& b, std::vector<T>& c) override;
 };
diff --git a/SYCLStream.cpp b/SYCLStream.cpp
index 12e96b4..919a657 100644
--- a/SYCLStream.cpp
+++ b/SYCLStream.cpp
@@ -20,6 +20,7 @@ program * p;
 
 /* Forward declaration of SYCL kernels */
 namespace kernels {
+  class init;
   class copy;
   class mul;
   class add;
@@ -46,6 +47,7 @@ SYCLStream<T>::SYCLStream(const unsigned int ARRAY_SIZE, const int device_index)
 
   /* Pre-build the kernels */
   p = new program(queue->get_context());
+  p->build_from_kernel_name<kernels::init>();
   p->build_from_kernel_name<kernels::copy>();
   p->build_from_kernel_name<kernels::mul>();
   p->build_from_kernel_name<kernels::add>();
@@ -142,17 +144,23 @@ void SYCLStream<T>::triad()
 }
 
 template <class T>
-void SYCLStream<T>::write_arrays(const std::vector<T>& a, const std::vector<T>& b, const std::vector<T>& c)
+void SYCLStream<T>::init_arrays(T initA, T initB, T initC)
 {
-  auto _a = d_a->template get_access<access::mode::write, access::target::host_buffer>();
-  auto _b = d_b->template get_access<access::mode::write, access::target::host_buffer>();
-  auto _c = d_c->template get_access<access::mode::write, access::target::host_buffer>();
-  for (int i = 0; i < array_size; i++)
+  queue->submit([&](handler &cgh)
   {
-    _a[i] = a[i];
-    _b[i] = b[i];
-    _c[i] = c[i];
-  }
+    auto ka = d_a->template get_access<access::mode::write>(cgh);
+    auto kb = d_b->template get_access<access::mode::write>(cgh);
+    auto kc = d_c->template get_access<access::mode::write>(cgh);
+    cgh.parallel_for<kernels::init>(p->get_kernel<kernels::init>(),
+          range<1>{array_size}, [=](item<1> item)
+    {
+      auto id = item.get();
+      ka[id[0]] = initA;
+      kb[id[0]] = initB;
+      kc[id[0]] = initC;
+    });
+  });
+  queue->wait();
 }
 
 template <class T>
diff --git a/SYCLStream.h b/SYCLStream.h
index 8bc515d..4bd21d8 100644
--- a/SYCLStream.h
+++ b/SYCLStream.h
@@ -38,7 +38,7 @@ class SYCLStream : public Stream<T>
     virtual void mul() override;
     virtual void triad() override;
 
-    virtual void write_arrays(const std::vector<T>& a, const std::vector<T>& b, const std::vector<T>& c) override;
+    virtual void init_arrays(T initA, T initB, T initC) override;
     virtual void read_arrays(std::vector<T>& a, std::vector<T>& b, std::vector<T>& c) override;
 
 };
diff --git a/Stream.h b/Stream.h
index 631e305..44b4d8b 100644
--- a/Stream.h
+++ b/Stream.h
@@ -31,7 +31,7 @@ class Stream
     virtual void triad() = 0;
 
     // Copy memory between host and device
-    virtual void write_arrays(const std::vector<T>& a, const std::vector<T>& b, const std::vector<T>& c) = 0;
+    virtual void init_arrays(T initA, T initB, T initC) = 0;
     virtual void read_arrays(std::vector<T>& a, std::vector<T>& b, std::vector<T>& c) = 0;
 
 };
diff --git a/main.cpp b/main.cpp
index 6a15aa7..cb6241f 100644
--- a/main.cpp
+++ b/main.cpp
@@ -84,9 +84,9 @@ void run()
     std::cout << "Precision: double" << std::endl;
 
   // Create host vectors
-  std::vector<T> a(ARRAY_SIZE, startA);
-  std::vector<T> b(ARRAY_SIZE, startB);
-  std::vector<T> c(ARRAY_SIZE, startC);
+  std::vector<T> a(ARRAY_SIZE);
+  std::vector<T> b(ARRAY_SIZE);
+  std::vector<T> c(ARRAY_SIZE);
   std::streamsize ss = std::cout.precision();
   std::cout << std::setprecision(1) << std::fixed
     << "Array size: " << ARRAY_SIZE*sizeof(T)*1.0E-6 << " MB"
@@ -135,7 +135,7 @@ void run()
 
 #endif
 
-  stream->write_arrays(a, b, c);
+  stream->init_arrays(startA, startB, startC);
 
   // List of times
   std::vector<std::vector<double>> timings(4);