From ef9df7f67c9158f696367b1b4a48f16d74391eb8 Mon Sep 17 00:00:00 2001
From: "Dakota St. Laurent" <d.h.stlaurent@gmail.com>
Date: Tue, 9 Jun 2015 16:07:18 -0400
Subject: [PATCH] add code to 01 that repeatedly writes buffers to show
 overhead; change vector names in kernel code to v1,v2,v3 to show that they're
 not the same as A,B,C in the C++ code

---
 README.md     | 10 +++++++--
 example01.cpp | 60 ++++++++++++++++++++++++++++++++++++++++++---------
 2 files changed, 58 insertions(+), 12 deletions(-)

diff --git a/README.md b/README.md
index 1fb19e0..730074e 100644
--- a/README.md
+++ b/README.md
@@ -9,8 +9,14 @@ timing results of vectors added together on a CPU vs GPU. the size of the vector
 
 - CPU code
 - GPU code equivalent
+- GPU code where the buffers are created and destroyed each iteration (this is guaranteed to be slower, of course, but it would be nice to see the amount of overhead that occurs with inefficient copying...)
 
 todo:
 
-- GPU code where a work-group barrier is initiated after each thread is done its work
-- GPU code where the buffers are created and destroyed each iteration (this is guaranteed to be slower, of course, but it would be nice to see the amount of overhead that occurs with inefficient copying...)
+- GPU code where a work-group barrier is initiated after each thread is done its work (equivalent to the regular GPU code, but with one extra line...)
+
+## TODO
+
+- figure out how OpenCL manages memory (when are buffers cleared on the GPU?)
+- figure out how to view the OpenCL assembly code if possible (is warp divergence happening?)
+
diff --git a/example01.cpp b/example01.cpp
index 365576a..f8ea162 100644
--- a/example01.cpp
+++ b/example01.cpp
@@ -3,7 +3,7 @@
 #include <ctime>
 
 
-double time_add_vectors(int n, int k) {
+double timeAddVectorsCPU(int n, int k) {
     // adds two vectors of size n, k times, returns total duration
     std::clock_t start;
     double duration;
@@ -26,7 +26,6 @@ double time_add_vectors(int n, int k) {
     return duration;
 }
 
-
 int main() {
     // get all platforms (drivers), e.g. NVIDIA
     std::vector<cl::Platform> all_platforms;
@@ -58,7 +57,7 @@ int main() {
     std::string kernel_code=
         //  is equivalent to the host's "time_add_vectors" function, except the
         //  timing will be done on the host.
-        "   void kernel looped_add(global const int* A, global const int* B, global int* C, "
+        "   void kernel looped_add(global const int* v1, global const int* v2, global int* v3, "
         "                          global const int* constants) {"
         "       int ID, Nthreads, n, k, ratio, start, stop;"
         "       ID = get_global_id(0);"
@@ -73,9 +72,23 @@ int main() {
         "       int i, j;" // will the compiler optimize this anyway? probably.
         "       for (i=0; i<k; i++) {"
         "           for (j=start; j<stop; j++)"
-        "               C[j] = A[j] + B[j];"
-        "           barrier(CLK_GLOBAL_MEM_FENCE);"
+        "               v3[j] = v1[j] + v2[j];"
         "       }"
+        "   }"
+        ""
+        "   void kernel add(global const int* v1, global const int* v2, global int* v3, "
+        "                   global const int* constants) {"
+        "       int ID, Nthreads, n, ratio, start, stop;"
+        "       ID = get_global_id(0);"
+        "       Nthreads = get_global_size(0);"
+        "       n = constants[0];"
+        ""
+        "       ratio = (n / Nthreads);"
+        "       start = ratio * ID;"
+        "       stop  = ratio * (ID+1);"
+        ""
+        "       for (int i=start; i<stop; i++)"
+        "           v3[i] = v1[i] + v2[i];"
         "   }";
     sources.push_back({kernel_code.c_str(), kernel_code.length()});
 
@@ -92,13 +105,14 @@ int main() {
     int constants[2] = {n, k};
 
     // run the CPU code
-    float CPUtime = time_add_vectors(n, k);
+    float CPUtime = timeAddVectorsCPU(n, k);
 
     // run some GPU code; this block allocates space, writes buffers, and then
     // adds the same two vectors multiple times -- i.e. it's equivalent to the
     // host (CPU) code above, but with some necessary overhead.
     cl::CommandQueue queue(context, default_device);
     cl::KernelFunctor looped_add(cl::Kernel(program, "looped_add"), queue, cl::NullRange, cl::NDRange(Nthreads), cl::NullRange);
+    cl::KernelFunctor add(cl::Kernel(program, "add"), queue, cl::NullRange, cl::NDRange(Nthreads), cl::NullRange);
 
     // construct vectors
     int A[n], B[n], C[n];
@@ -125,20 +139,46 @@ int main() {
     queue.enqueueWriteBuffer(buffer_constants, CL_TRUE, 0, sizeof(int)*2, constants);
 
     // RUN ZE KERNEL
-    std::cout << "Running...";
     looped_add(buffer_A, buffer_B, buffer_C, buffer_constants);
-    std::cout << "Cool." << std::endl;
 
-    // read result from GPU to here
+    // read result from GPU to here; including for the sake of timing
     queue.enqueueReadBuffer(buffer_C, CL_TRUE, 0, sizeof(int)*n, C);
-    
     GPUtime1 = (std::clock() - start_time) / (double) CLOCKS_PER_SEC;
 
+    // do the same thing, except copy the arrays over every iteration
+    double GPUtime2;
+    start_time = std::clock();
+
+    cl::Buffer buffer_A2(context, CL_MEM_READ_WRITE, sizeof(int)*n);
+    cl::Buffer buffer_B2(context, CL_MEM_READ_WRITE, sizeof(int)*n);
+    cl::Buffer buffer_C2(context, CL_MEM_READ_WRITE, sizeof(int)*n);
+    cl::Buffer buffer_constants2(context, CL_MEM_READ_ONLY, sizeof(int)*2);
+    for (int i=0; i<k; i++) {
+        queue.enqueueWriteBuffer(buffer_A2, CL_TRUE, 0, sizeof(int)*n, A);
+        queue.enqueueWriteBuffer(buffer_B2, CL_TRUE, 0, sizeof(int)*n, B);
+        queue.enqueueWriteBuffer(buffer_constants2, CL_TRUE, 0, sizeof(int)*2, constants);
+
+        add(buffer_A2, buffer_B2, buffer_C2, buffer_constants2);
+    }
+    queue.enqueueReadBuffer(buffer_C2, CL_TRUE, 0, sizeof(int)*n, C);
+    GPUtime2 = (std::clock() - start_time) / (double) CLOCKS_PER_SEC;
+
     // let's compare!
     double time_ratio = (CPUtime / GPUtime1);
+    std::cout << "VERSION 1 -----------" << std::endl;
     std::cout << "CPU time: " << CPUtime  << std::endl;
     std::cout << "GPU time: " << GPUtime1 << std::endl;
     std::cout << "GPU is ";
+    if (time_ratio > 1)
+        std::cout << time_ratio << " times faster!" << std::endl;
+    else
+        std::cout << time_ratio << " times slower :(" << std::endl;
+
+    time_ratio = (CPUtime / GPUtime2);
+    std::cout << "\nVERSION 2 -----------" << std::endl;
+    std::cout << "CPU time: " << CPUtime  << std::endl;
+    std::cout << "GPU time: " << GPUtime2 << std::endl;
+    std::cout << "GPU is ";
     if (time_ratio > 1)
         std::cout << time_ratio << " times faster!" << std::endl;
     else