diff --git a/SYCLStream.cpp b/SYCLStream.cpp
index 45cf477..995eb1b 100644
--- a/SYCLStream.cpp
+++ b/SYCLStream.cpp
@@ -16,6 +16,9 @@ using namespace cl::sycl;
 bool cached = false;
 std::vector<device> devices;
 void getDeviceList(void);
+#ifdef COMPUTECPP_CE
+program * p;
+#endif
 
 template <class T>
 SYCLStream<T>::SYCLStream(const unsigned int ARRAY_SIZE, const int device_index)
@@ -66,6 +69,17 @@ SYCLStream<T>::SYCLStream(const unsigned int ARRAY_SIZE, const int device_index)
       throw std::runtime_error("SYCL errors detected");
     }
   });
+  
+  #ifdef COMPUTECPP_CE
+  /* Pre-build the kernels */
+  p = new program(queue->get_context());
+  p->build_with_kernel_type<init_kernel>();
+  p->build_with_kernel_type<copy_kernel>();
+  p->build_with_kernel_type<mul_kernel>();
+  p->build_with_kernel_type<add_kernel>();
+  p->build_with_kernel_type<triad_kernel>();
+  p->build_with_kernel_type<dot_kernel>();
+  #endif
 
   // Create buffers
   d_a = new buffer<T>(array_size);
@@ -81,8 +95,9 @@ SYCLStream<T>::~SYCLStream()
   delete d_b;
   delete d_c;
   delete d_sum;
-
-
+  #ifdef COMPUTECPP_CE
+  delete p;
+  #endif
   delete queue;
   devices.clear();
 }
@@ -94,7 +109,12 @@ void SYCLStream<T>::copy()
   {
     auto ka = d_a->template get_access<access::mode::read>(cgh);
     auto kc = d_c->template get_access<access::mode::write>(cgh);
+    #ifdef COMPUTECPP_CE
+    cgh.parallel_for<copy_kernel>(p->get_kernel<copy_kernel>(),
+          range<1>{array_size}, [=](item<1> item)
+    #else
     cgh.parallel_for<copy_kernel>(range<1>{array_size}, [=](item<1> item)
+    #endif
     {
       auto id = item.get_id(0);
       kc[id] = ka[id];
@@ -111,7 +131,12 @@ void SYCLStream<T>::mul()
   {
     auto kb = d_b->template get_access<access::mode::write>(cgh);
     auto kc = d_c->template get_access<access::mode::read>(cgh);
+    #ifdef COMPUTECPP_CE
+    cgh.parallel_for<mul_kernel>(p->get_kernel<mul_kernel>(),
+      range<1>{array_size}, [=](item<1> item)
+    #else
     cgh.parallel_for<mul_kernel>(range<1>{array_size}, [=](item<1> item)
+    #endif
     {
       auto id = item.get_id(0);
       kb[id] = scalar * kc[id];
@@ -128,7 +153,12 @@ void SYCLStream<T>::add()
     auto ka = d_a->template get_access<access::mode::read>(cgh);
     auto kb = d_b->template get_access<access::mode::read>(cgh);
     auto kc = d_c->template get_access<access::mode::write>(cgh);
+    #ifdef COMPUTECPP_CE
+    cgh.parallel_for<add_kernel>(p->get_kernel<add_kernel>(),
+      range<1>{array_size}, [=](item<1> item)
+    #else
     cgh.parallel_for<add_kernel>(range<1>{array_size}, [=](item<1> item)
+    #endif
     {
       auto id = item.get_id(0);
       kc[id] = ka[id] + kb[id];
@@ -146,7 +176,12 @@ void SYCLStream<T>::triad()
     auto ka = d_a->template get_access<access::mode::write>(cgh);
     auto kb = d_b->template get_access<access::mode::read>(cgh);
     auto kc = d_c->template get_access<access::mode::read>(cgh);
+    #ifdef COMPUTECPP_CE
+    cgh.parallel_for<triad_kernel>(p->get_kernel<triad_kernel>(),
+      range<1>{array_size}, [=](item<1> item)
+    #else
     cgh.parallel_for<triad_kernel>(range<1>{array_size}, [=](item<1> item)
+    #endif
     {
       auto id = item.get_id(0);
       ka[id] = kb[id] + scalar * kc[id];
@@ -167,8 +202,12 @@ T SYCLStream<T>::dot()
     auto wg_sum = accessor<T, 1, access::mode::read_write, access::target::local>(range<1>(dot_wgsize), cgh);
 
     size_t N = array_size;
-
+    #ifdef COMPUTECPP_CE
+    cgh.parallel_for<dot_kernel>(p->get_kernel<dot_kernel>(),
+      nd_range<1>(dot_num_groups*dot_wgsize, dot_wgsize), [=](nd_item<1> item)
+    #else
     cgh.parallel_for<dot_kernel>(nd_range<1>(dot_num_groups*dot_wgsize, dot_wgsize), [=](nd_item<1> item)
+    #endif
     {
       size_t i = item.get_global_id(0);
       size_t li = item.get_local_id(0);
@@ -209,7 +248,12 @@ void SYCLStream<T>::init_arrays(T initA, T initB, T initC)
     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);
+    #ifdef COMPUTECPP_CE
+    cgh.parallel_for<init_kernel>(p->get_kernel<init_kernel>(),
+      range<1>{array_size}, [=](item<1> item)
+    #else
     cgh.parallel_for<init_kernel>(range<1>{array_size}, [=](item<1> item)
+    #endif
     {
       auto id = item.get_id(0);
       ka[id] = initA;
diff --git a/SYCLStream.h b/SYCLStream.h
index ab62ecd..9c06f6f 100644
--- a/SYCLStream.h
+++ b/SYCLStream.h
@@ -15,6 +15,9 @@
 
 #define IMPLEMENTATION_STRING "SYCL"
 
+// allows a use of 'parallel_for' currently known to be supported by ComputeCpp
+#define COMPUTECPP_CE
+
 namespace sycl_kernels
 {
   template <class T> class init;