From 5c5a02b49917e40594835919941a4e2b45a80792 Mon Sep 17 00:00:00 2001
From: "Dakota St. Laurent" <d.h.stlaurent@gmail.com>
Date: Fri, 19 Jun 2015 09:50:49 -0400
Subject: [PATCH] added fifth version, which gives very strange results...

---
 example01/main.cpp | 70 ++++++++++++++++++++++++++++++++--------------
 1 file changed, 49 insertions(+), 21 deletions(-)

diff --git a/example01/main.cpp b/example01/main.cpp
index 5284a40..6be5adf 100644
--- a/example01/main.cpp
+++ b/example01/main.cpp
@@ -1,6 +1,5 @@
 #include <iostream>
 #include <ctime>
-#include <string.h>
 #include "../CL/cl.hpp"
 
 
@@ -42,9 +41,6 @@ double timeAddVectorsCPU(int n, int k) {
 
 int main(int argc, char* argv[]) {
 
-    for (int j=0; j<argc; j++)
-        std::cout << argv[j] << std::endl;
-
     bool verbose;
     if (argc == 1 || std::strcmp(argv[1], "0") == 0)
         verbose = true;
@@ -120,23 +116,28 @@ int main(int argc, char* argv[]) {
         "           v3[i] = v1[i] + v2[i];"
         "   }"
         ""
-        "   void kernel add_single(global const int* v1, global const int* v2, global int* v3) { "
+        "   void kernel add_single(global const int* v1, global const int* v2, global int* v3, "
+        "                          global const int* constants) { "
+        "       int k  = constants[1];"
         "       int ID = get_global_id(0);"
-        "       v3[ID] = v1[ID] + v2[ID];"
+        "       for (int i=0; i<k; i++)"
+        "           v3[ID] = v1[ID] + v2[ID];"
         "   }"
-        ""  // same as add_single, but with the overhead of the looped functions
+        ""  // same as add_single, but with the overhead (indexing, determining ratio) of versions 01 and 02
         "   void kernel add_single_overhead(global const int* v1, global const int* v2, global int* v3,"
         "                                   global const int* constants) {"
-        "       int ID, NUM_GLOBAL_WITEMS, n, ratio, start, stop;"
+        "       int ID, NUM_GLOBAL_WITEMS, n, k, ratio, start, stop;"
         "       ID = get_global_id(0);"
         "       NUM_GLOBAL_WITEMS = get_global_size(0);"
         "       n = constants[0];"
+        "       k = constants[1];"
         ""
         "       ratio = (n / NUM_GLOBAL_WITEMS);"
         "       start = ratio * ID;"
         "       stop  = ratio * (ID+1);"
         ""
-        "       v3[ID] = v1[ID] + v2[ID];"
+        "       for (int i=0; i<k; i++)"
+        "           v3[ID] = v1[ID] + v2[ID];"
         "   }";
     sources.push_back({kernel_code.c_str(), kernel_code.length()});
 
@@ -163,11 +164,12 @@ int main(int argc, char* argv[]) {
         B[i] = n - i - 1;
         C[i] = 0;
     }
+    std::clock_t start_time;
+
 
     // VERSION 1 ==========================================
     // start timer
     double GPUtime1;
-    std::clock_t start_time;
     start_time = std::clock();
 
     // allocate space
@@ -192,8 +194,8 @@ int main(int argc, char* argv[]) {
 
     // 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;
     queue.enqueueBarrier();
+    GPUtime1 = (std::clock() - start_time) / (double) CLOCKS_PER_SEC;
 
     // VERSION 2 ==========================================
     double GPUtime2;
@@ -208,15 +210,15 @@ int main(int argc, char* argv[]) {
         queue.enqueueWriteBuffer(buffer_B2, CL_TRUE, 0, sizeof(int)*n, B);
         queue.enqueueWriteBuffer(buffer_constants2, CL_TRUE, 0, sizeof(int)*2, constants);
 
-        add_looped.setArg(0, buffer_A);
-        add_looped.setArg(1, buffer_B);
-        add_looped.setArg(2, buffer_C);
+        add_looped.setArg(0, buffer_A2);
+        add_looped.setArg(1, buffer_B2);
+        add_looped.setArg(2, buffer_C2);
         add_looped.setArg(3, buffer_constants);
         queue.enqueueNDRangeKernel(add, cl::NullRange, cl::NDRange(NUM_GLOBAL_WITEMS), cl::NDRange(32));
     }
     queue.enqueueReadBuffer(buffer_C2, CL_TRUE, 0, sizeof(int)*n, C);
-    GPUtime2 = (std::clock() - start_time) / (double) CLOCKS_PER_SEC;
     queue.enqueueBarrier();
+    GPUtime2 = (std::clock() - start_time) / (double) CLOCKS_PER_SEC;
 
     // VERSION 3 ==========================================
     double GPUtime3;
@@ -225,16 +227,19 @@ int main(int argc, char* argv[]) {
     cl::Buffer buffer_A3(context, CL_MEM_READ_WRITE, sizeof(int)*n);
     cl::Buffer buffer_B3(context, CL_MEM_READ_WRITE, sizeof(int)*n);
     cl::Buffer buffer_C3(context, CL_MEM_READ_WRITE, sizeof(int)*n);
+    cl::Buffer buffer_constants3(context, CL_MEM_READ_ONLY, sizeof(int) * 2);
     queue.enqueueWriteBuffer(buffer_A3, CL_TRUE, 0, sizeof(int)*n, A);
     queue.enqueueWriteBuffer(buffer_B3, CL_TRUE, 0, sizeof(int)*n, B);
+    queue.enqueueWriteBuffer(buffer_constants3, CL_TRUE, 0, sizeof(int)*2, constants);
 
     add_single.setArg(0, buffer_A3);
     add_single.setArg(1, buffer_B3);
     add_single.setArg(2, buffer_C3);
+    add_single.setArg(3, buffer_constants3);
     queue.enqueueNDRangeKernel(add_single, cl::NullRange, cl::NDRange(n), cl::NDRange(32));
 
-    // end timer
     queue.enqueueReadBuffer(buffer_C3, CL_TRUE, 0, sizeof(int)*n, C);
+    queue.enqueueBarrier();
     GPUtime3 = (std::clock() - start_time) / (double) CLOCKS_PER_SEC;
 
     // VERSION 4 ==========================================
@@ -255,20 +260,43 @@ int main(int argc, char* argv[]) {
     add_single_overhead.setArg(3, buffer_constants4);
     queue.enqueueNDRangeKernel(add_single_overhead, cl::NullRange, cl::NDRange(n), cl::NDRange(32));
 
-    // end timer
     queue.enqueueReadBuffer(buffer_C4, CL_TRUE, 0, sizeof(int)*n, C);
+    queue.enqueueBarrier();
     GPUtime4 = (std::clock() - start_time) / (double) CLOCKS_PER_SEC;
+
+    // VERSION 5 ==========================================
+    double GPUtime5;
+    start_time = std::clock();
+
+    cl::Buffer buffer_A5(context, CL_MEM_READ_WRITE, sizeof(int) * n);
+    cl::Buffer buffer_B5(context, CL_MEM_READ_WRITE, sizeof(int) * n);
+    cl::Buffer buffer_C5(context, CL_MEM_READ_WRITE, sizeof(int) * n);
+    cl::Buffer buffer_constants5(context, CL_MEM_READ_ONLY, sizeof(int) * 2);
+    queue.enqueueWriteBuffer(buffer_A5, CL_TRUE, 0, sizeof(int)*n, A);
+    queue.enqueueWriteBuffer(buffer_B5, CL_TRUE, 0, sizeof(int)*n, B);
+    queue.enqueueWriteBuffer(buffer_constants5, CL_TRUE, 0, sizeof(int)*2, constants);
+
+    add_looped.setArg(0, buffer_A5);
+    add_looped.setArg(1, buffer_B5);
+    add_looped.setArg(2, buffer_C5);
+    add_looped.setArg(3, buffer_constants5);
+    queue.enqueueNDRangeKernel(add_looped, cl::NullRange, cl::NDRange(n), cl::NDRange(32));
+
+    queue.enqueueReadBuffer(buffer_C5, CL_TRUE, 0, sizeof(int)*n, C);
+    queue.enqueueBarrier();
+    GPUtime5 = (std::clock() - start_time) / (double) CLOCKS_PER_SEC;
     
     // let's compare!
-    double GPUtimes[4] = {GPUtime1, GPUtime2, GPUtime3, GPUtime4};
+    const int NUM_VERSIONS = 5;
+    double GPUtimes[NUM_VERSIONS] = {GPUtime1, GPUtime2, GPUtime3, GPUtime4, GPUtime5};
     if (verbose) {
-        for (int i=0; i<4; i++)
+        for (int i=0; i<NUM_VERSIONS; i++)
             compareResults(CPUtime, GPUtimes[i], i+1);
     } else {
         std::cout << CPUtime << ",";
-        for (int i=0; i<3; i++)
+        for (int i=0; i<NUM_VERSIONS-1; i++)
             std::cout << GPUtimes[i] << ",";
-        std::cout << GPUtimes[4] << std::endl;
+        std::cout << GPUtimes[NUM_VERSIONS-1] << std::endl;
     }
     return 0;
 }