diff --git a/.gitignore b/.gitignore
index 1f283e2..6b4a81a 100644
--- a/.gitignore
+++ b/.gitignore
@@ -1,4 +1,6 @@
 /target
 .idea
 primes.txt
-timings.csv
\ No newline at end of file
+timings.csv
+*.csv
+*.~lock*
\ No newline at end of file
diff --git a/src/kernel_controller/bench.rs b/src/kernel_controller/bench.rs
new file mode 100644
index 0000000..43ed9cf
--- /dev/null
+++ b/src/kernel_controller/bench.rs
@@ -0,0 +1,55 @@
+/*
+ * opencl demos with rust
+ * Copyright (C) 2020 trivernis
+ * See LICENSE for more information
+ */
+
+use crate::kernel_controller::KernelController;
+use std::time::{Duration, Instant};
+
+pub struct BenchStatistics {
+    pub calc_count: u32,
+    pub num_tasks: usize,
+    pub write_duration: Duration,
+    pub calc_duration: Duration,
+    pub read_duration: Duration,
+}
+
+impl KernelController {
+    /// Benches an integer
+    pub fn bench_int(&self, calc_count: u32, num_tasks: usize) -> ocl::Result<BenchStatistics> {
+        let write_start = Instant::now();
+        let input_buffer = self
+            .pro_que
+            .buffer_builder()
+            .len(num_tasks)
+            .fill_val(0u32)
+            .build()?;
+        let write_duration = write_start.elapsed();
+
+        let kernel = self
+            .pro_que
+            .kernel_builder("bench_int")
+            .arg(calc_count)
+            .arg(&input_buffer)
+            .build()?;
+        let calc_start = Instant::now();
+        unsafe {
+            kernel.enq()?;
+        }
+        self.pro_que.finish()?;
+        let calc_duration = calc_start.elapsed();
+        let mut output = vec![0u32; num_tasks];
+        let read_start = Instant::now();
+        input_buffer.read(&mut output).enq()?;
+        let read_duration = read_start.elapsed();
+
+        Ok(BenchStatistics {
+            num_tasks,
+            calc_count,
+            read_duration,
+            calc_duration,
+            write_duration,
+        })
+    }
+}
diff --git a/src/kernel_controller/kernel.cl b/src/kernel_controller/kernel.cl
index 79c8bd6..9a82e66 100644
--- a/src/kernel_controller/kernel.cl
+++ b/src/kernel_controller/kernel.cl
@@ -58,4 +58,13 @@ __kernel void check_prime(__global const ulong *IN, __global bool *OUT) {
     }
 
     OUT[id] = true;
+}
+
+__kernel void bench_int(const uint limit, __global int *NUMBERS) {
+    uint id = get_global_id(0);
+    int num = NUMBERS[id];
+    for (int i = 0; i < limit; i++) {
+        num += i;
+    }
+    NUMBERS[id] = num;
 }
\ No newline at end of file
diff --git a/src/kernel_controller/mod.rs b/src/kernel_controller/mod.rs
index 8003837..eec8cf6 100644
--- a/src/kernel_controller/mod.rs
+++ b/src/kernel_controller/mod.rs
@@ -8,6 +8,7 @@ use ocl::core::DeviceInfo;
 use ocl::enums::DeviceInfoResult;
 use ocl::ProQue;
 
+pub mod bench;
 pub mod primes;
 
 pub struct KernelController {
@@ -31,6 +32,14 @@ impl KernelController {
             "Max Mem Alloc: {} bytes",
             device.info(DeviceInfo::MaxMemAllocSize)?
         );
+        println!(
+            "Max Compute Units: {}",
+            device.info(DeviceInfo::MaxComputeUnits)?
+        );
+        println!(
+            "Max Work Group Size: {}",
+            device.info(DeviceInfo::MaxWorkGroupSize)?
+        );
         println!();
         Ok(Self { pro_que })
     }
diff --git a/src/main.rs b/src/main.rs
index 3783564..d616fce 100644
--- a/src/main.rs
+++ b/src/main.rs
@@ -6,17 +6,18 @@
 
 use crate::kernel_controller::primes::is_prime;
 use crate::kernel_controller::KernelController;
+use crate::output::csv::CSVWriter;
+use crate::output::{create_csv_write_thread, create_prime_write_thread};
 use rayon::prelude::*;
-use std::fs::{File, OpenOptions};
-use std::io::{BufWriter, Write};
+use std::fs::OpenOptions;
+use std::io::BufWriter;
 use std::mem;
 use std::path::PathBuf;
-use std::sync::mpsc::{channel, Sender};
-use std::thread::{self, JoinHandle};
-use std::time::Instant;
+use std::time::{Duration, Instant};
 use structopt::StructOpt;
 
 mod kernel_controller;
+mod output;
 
 #[derive(StructOpt, Clone, Debug)]
 #[structopt()]
@@ -77,7 +78,7 @@ fn calculate_primes(prime_opts: CalculatePrimes, controller: KernelController) -
             .open(prime_opts.output_file)
             .unwrap(),
     );
-    let mut timings = BufWriter::new(
+    let timings = BufWriter::new(
         OpenOptions::new()
             .create(true)
             .truncate(true)
@@ -85,17 +86,26 @@ fn calculate_primes(prime_opts: CalculatePrimes, controller: KernelController) -
             .open(prime_opts.timings_file)
             .unwrap(),
     );
-    timings
-        .write_all("offset,count,gpu_duration,filter_duration,duration\n".as_bytes())
-        .unwrap();
-    let (sender, handle) = create_write_thread(output);
+    let timings = CSVWriter::new(
+        timings,
+        &[
+            "offset",
+            "count",
+            "gpu_duration",
+            "filter_duration",
+            "total_duration",
+        ],
+    );
+
+    let (prime_sender, prime_handle) = create_prime_write_thread(output);
+    let (csv_sender, csv_handle) = create_csv_write_thread(timings);
 
     let mut offset = prime_opts.start_offset;
     if offset % 2 == 0 {
         offset += 1;
     }
     if offset < 2 {
-        sender.send(vec![2]).unwrap();
+        prime_sender.send(vec![2]).unwrap();
     }
     loop {
         let start = Instant::now();
@@ -121,26 +131,21 @@ fn calculate_primes(prime_opts: CalculatePrimes, controller: KernelController) -
             elapsed_ms,
             prime_opts.numbers_per_step as f64 / start.elapsed().as_secs_f64()
         );
-        timings
-            .write_all(
-                format!(
-                    "{},{},{},{},{}\n",
-                    offset,
-                    primes.len(),
-                    prime_result.gpu_duration.as_secs_f64() * 1000f64,
-                    prime_result.filter_duration.as_secs_f64() * 1000f64,
-                    elapsed_ms
-                )
-                .as_bytes(),
-            )
+        csv_sender
+            .send(vec![
+                offset.to_string(),
+                primes.len().to_string(),
+                duration_to_ms_string(&prime_result.gpu_duration),
+                duration_to_ms_string(&prime_result.filter_duration),
+                elapsed_ms.to_string(),
+            ])
             .unwrap();
-        timings.flush().unwrap();
 
         if prime_opts.cpu_validate {
             validate_primes_on_cpu(&primes)
         }
         println!();
-        sender.send(primes).unwrap();
+        prime_sender.send(primes).unwrap();
 
         if (prime_opts.numbers_per_step as u128 * 2 + offset as u128)
             > prime_opts.max_number as u128
@@ -150,8 +155,9 @@ fn calculate_primes(prime_opts: CalculatePrimes, controller: KernelController) -
         offset += prime_opts.numbers_per_step as u64 * 2;
     }
 
-    mem::drop(sender);
-    handle.join().unwrap();
+    mem::drop(prime_sender);
+    prime_handle.join().unwrap();
+    csv_handle.join().unwrap();
 
     Ok(())
 }
@@ -173,16 +179,6 @@ fn validate_primes_on_cpu(primes: &Vec<u64>) {
     }
 }
 
-fn create_write_thread(mut writer: BufWriter<File>) -> (Sender<Vec<u64>>, JoinHandle<()>) {
-    let (tx, rx) = channel();
-    let handle = thread::spawn(move || {
-        for primes in rx {
-            for prime in primes {
-                writer.write_all(format!("{}\n", prime).as_bytes()).unwrap();
-            }
-            writer.flush().unwrap();
-        }
-    });
-
-    (tx, handle)
+fn duration_to_ms_string(duration: &Duration) -> String {
+    format!("{}", duration.as_secs_f64() * 1000f64)
 }
diff --git a/src/output/csv.rs b/src/output/csv.rs
new file mode 100644
index 0000000..dfaf08e
--- /dev/null
+++ b/src/output/csv.rs
@@ -0,0 +1,52 @@
+/*
+ * opencl demos with rust
+ * Copyright (C) 2020 trivernis
+ * See LICENSE for more information
+ */
+
+use std::collections::HashMap;
+use std::io::{Result, Write};
+
+pub struct CSVWriter<W: Write> {
+    inner: W,
+    columns: Vec<String>,
+}
+
+impl<W> CSVWriter<W>
+where
+    W: Write,
+{
+    /// Creates a new CSVWriter with a defined list of columns
+    pub fn new(writer: W, columns: &[&str]) -> Self {
+        Self {
+            inner: writer,
+            columns: columns.iter().map(|column| column.to_string()).collect(),
+        }
+    }
+
+    /// Adds a new row of values to the file
+    pub fn add_row(&mut self, items: Vec<String>) -> Result<()> {
+        self.inner.write_all(
+            items
+                .iter()
+                .fold("".to_string(), |a, b| format!("{},{}", a, b))
+                .as_bytes(),
+        )?;
+        self.inner.write_all("\n".as_bytes())
+    }
+
+    /// Adds a new row of values stored in a map to the file
+    #[allow(dead_code)]
+    pub fn add_row_map(&mut self, item_map: &HashMap<String, String>) -> Result<()> {
+        let mut items = Vec::new();
+        for key in &self.columns {
+            items.push(item_map.get(key).cloned().unwrap_or("".to_string()));
+        }
+
+        self.add_row(items)
+    }
+
+    pub fn flush(&mut self) -> Result<()> {
+        self.inner.flush()
+    }
+}
diff --git a/src/output/mod.rs b/src/output/mod.rs
new file mode 100644
index 0000000..96f48c9
--- /dev/null
+++ b/src/output/mod.rs
@@ -0,0 +1,43 @@
+/*
+ * opencl demos with rust
+ * Copyright (C) 2020 trivernis
+ * See LICENSE for more information
+ */
+use crate::output::csv::CSVWriter;
+use std::fmt::Display;
+use std::fs::File;
+use std::io::{BufWriter, Write};
+use std::sync::mpsc::{channel, Sender};
+use std::thread::{self, JoinHandle};
+
+pub mod csv;
+
+pub fn create_prime_write_thread(
+    mut writer: BufWriter<File>,
+) -> (Sender<Vec<u64>>, JoinHandle<()>) {
+    let (tx, rx) = channel();
+    let handle = thread::spawn(move || {
+        for primes in rx {
+            for prime in primes {
+                writer.write_all(format!("{}\n", prime).as_bytes()).unwrap();
+            }
+            writer.flush().unwrap();
+        }
+    });
+
+    (tx, handle)
+}
+
+pub fn create_csv_write_thread(
+    mut writer: CSVWriter<BufWriter<File>>,
+) -> (Sender<Vec<String>>, JoinHandle<()>) {
+    let (tx, rx) = channel();
+    let handle = thread::spawn(move || {
+        for row in rx {
+            writer.add_row(row).unwrap();
+        }
+        writer.flush().unwrap();
+    });
+
+    (tx, handle)
+}