diff --git a/src/OCLStream.cpp b/src/OCLStream.cpp
index eac19c0..f60c161 100644
--- a/src/OCLStream.cpp
+++ b/src/OCLStream.cpp
@@ -1,6 +1,47 @@
 
 #include "OCLStream.h"
 
+std::string kernels{R"CLC(
+
+  constant TYPE scalar = 3.0;
+
+  kernel void copy(
+    global const TYPE * restrict a,
+    global TYPE * restrict c)
+  {
+    const size_t i = get_global_id(0);
+    c[i] = a[i];
+  }
+
+  kernel void mul(
+    global TYPE * restrict b,
+    global const TYPE * restrict c)
+  {
+    const size_t i = get_global_id(0);
+    b[i] = scalar * c[i];
+  }
+
+  kernel void add(
+    global const TYPE * restrict a,
+    global const TYPE * restrict b,
+    global TYPE * restrict c)
+  {
+    const size_t i = get_global_id(0);
+    c[i] = a[i] + b[i];
+  }
+
+  kernel void triad(
+    global TYPE * restrict a,
+    global const TYPE * restrict b,
+    global const TYPE * restrict c)
+  {
+    const size_t i = get_global_id(0);
+    a[i] = b[i] + scalar * c[i];
+  }
+
+)CLC"};
+
+
 template <class T>
 OCLStream<T>::OCLStream(const unsigned int ARRAY_SIZE)
 {
@@ -8,31 +49,30 @@ OCLStream<T>::OCLStream(const unsigned int ARRAY_SIZE)
 
   // Setup default OpenCL GPU
   context = cl::Context::getDefault();
-  //queue = cl::CommandQueue::getDefault();
+  queue = cl::CommandQueue::getDefault();
 
   // Create program
+  cl::Program program(kernels);
+  if (sizeof(T) == sizeof(double))
+    program.build("-DTYPE=double");
+  else if (sizeof(T) == sizeof(float))
+    program.build("-DTYPE=float");
 
-  std::string kernels{R"CLC(
-
-    const double scalar = 3.0;
-
-    kernel void copy(
-      global const double * restrict a,
-      global double * restrict c)
-    {
-      const size_t i = get_global_id(0);
-      c[i] = a[i];
-    }
-  )CLC"};
- 
-std::cout << kernels << std::endl;
- 
-  //cl::Program program(kernels);
-  //program.build();
-
-exit(-1);
+  // Create kernels
+  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");
+  triad_kernel = new cl::KernelFunctor<cl::Buffer, cl::Buffer, cl::Buffer>(program, "triad");
 
+}
 
+template <class T>
+OCLStream<T>::~OCLStream()
+{
+  delete[] copy_kernel;
+  delete[] mul_kernel;
+  delete[] add_kernel;
+  delete[] triad_kernel;
 }
 
 template <class T>
diff --git a/src/OCLStream.h b/src/OCLStream.h
index a5cac61..28a9be1 100644
--- a/src/OCLStream.h
+++ b/src/OCLStream.h
@@ -26,6 +26,11 @@ class OCLStream : public Stream<T>
     cl::Context context;
     cl::CommandQueue queue;
 
+    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;
+    cl::KernelFunctor<cl::Buffer, cl::Buffer, cl::Buffer> *triad_kernel;
+
   public:
 
     OCLStream(const unsigned int);