Turn loop into a test and add some loop test cases to run_tests.sh

6 files changed, 88 insertions, 37 deletions

diff --git a/loop.c b/loop.c
index 151dfde..9a806ef 100644
--- a/loop.c
+++ b/loop.c
@@ -15,61 +15,66 @@ int main (int argc, char ** argv)
 
    cl_int error;
 
-   cl_mem out_buffer;
-   int * out_data;
-   size_t global_work_size = 5;
+   unsigned pass = 1;
+   int * out;
+   char * kernel_name;
+   size_t global_work_size = 10;
    int iterations;
-   unsigned out_data_size;
+   unsigned out_size;
 
-   if (argc != 2) {
-      fprintf(stderr, "Usage: loop iterations\n");
+   if (argc != 3) {
+      fprintf(stderr, "Usage: loop kernel_name iterations\n");
       return EXIT_FAILURE;
    }
 
-   iterations = atoi(argv[1]);
-   out_data_size = global_work_size * iterations * sizeof(int);
-   out_data = malloc(out_data_size);
+   kernel_name = argv[1];
+   iterations = atoi(argv[2]);
+   out_size = global_work_size * iterations * sizeof(int);
+   out = malloc(out_size);
 
-   if (!clSimpleSimpleInit(&context, "loop")) {
+   if (!clSimpleSimpleInit(&context, kernel_name)) {
       return EXIT_FAILURE;
    }
 
-   out_buffer = clCreateBuffer(context.cl_ctx,
-                              CL_MEM_WRITE_ONLY,
-                              out_data_size,
-                              NULL, &error);
-
-   assert(error == CL_SUCCESS);
-
-   if (!clSimpleKernelSetArg(context.kernel, 0, sizeof(cl_mem), &out_buffer)
-       || !clSimpleKernelSetArg(context.kernel, 1, sizeof(int), &iterations)) {
+   if (!clSimpleSetOutputBuffer(&context, out_size)) {
       return EXIT_FAILURE;
    }
 
-   error = clEnqueueNDRangeKernel(context.command_queue,
-                                  context.kernel,
-                                  1,
-                                  NULL,
-                                  &global_work_size,
-                                  &global_work_size,
-                                  0, NULL, NULL);
+   if (error != CL_SUCCESS) {
+      fprintf(stderr, "clCreateBuffer() failed\n");
+      return EXIT_FAILURE;
+   }
 
-   assert(error == CL_SUCCESS);
+   if (!clSimpleKernelSetArg(context.kernel, 1, sizeof(int), &iterations)) {
+      return EXIT_FAILURE;
+   }
 
-   error = clEnqueueReadBuffer(context.command_queue,
-                               out_buffer,
-                               CL_TRUE,
-                               0,
-                               out_data_size,
-                               out_data,
-                               0, NULL, NULL);
+   if( !clSimpleEnqueueNDRangeKernel(context.command_queue,
+                              context.kernel,
+                              1, &global_work_size, &global_work_size)) {
+      return EXIT_FAILURE;
+   }
 
-   assert(error == CL_SUCCESS);
+   if (!clSimpleReadOutput(&context, out, out_size)) {
+      return EXIT_FAILURE;
+   }
 
    for (i = 0; i < global_work_size; i++) {
       for (j = 0; j < iterations; j++) {
-         fprintf(stderr, "%2u ", out_data[(i * iterations) + j]);
+            unsigned id = (i * iterations) + j;
+            if (out[id] != id) {
+               fprintf(stderr, "\n\nExpected %d, but got %d\n\n", id, out[id]);
+               pass = 0;
+            }
+         fprintf(stderr, "%2u ", out[id]);
       }
       fprintf(stderr, "\n");
    }
+   if (pass) {
+      fprintf(stderr, "Pass\n");
+      return EXIT_SUCCESS;
+   } else {
+      fprintf(stderr, "Fail\n");
+      return EXIT_FAILURE;
+   }
 }
diff --git a/loop_ge.cl b/loop_ge.cl
new file mode 100644
index 0000000..8ef8c7c
--- /dev/null
+++ b/loop_ge.cl
@@ -0,0 +1,9 @@
+__kernel void loop_ge(__global int * out, int iterations)
+{
+  unsigned i;
+  unsigned base = get_global_id(0) * iterations;
+  for (i = iterations; i >= 1; i--) {
+    unsigned id = (base + i);
+    out[id] = id;
+  }
+}
diff --git a/loop_gt.cl b/loop_gt.cl
new file mode 100644
index 0000000..6e95fcc
--- /dev/null
+++ b/loop_gt.cl
@@ -0,0 +1,9 @@
+__kernel void loop_gt(__global int * out, int iterations)
+{
+  unsigned i;
+  unsigned base = get_global_id(0) * iterations;
+  for (i = iterations; i > 0; i--) {
+    unsigned id = (base + i);
+    out[id] = id;
+  }
+}
diff --git a/loop_le.cl b/loop_le.cl
new file mode 100644
index 0000000..a798fd4
--- /dev/null
+++ b/loop_le.cl
@@ -0,0 +1,9 @@
+__kernel void loop_le(__global int * out, int iterations)
+{
+  unsigned i;
+  unsigned base = get_global_id(0) * iterations;
+  for (i = 0; i <= iterations - 1; i++) {
+    unsigned id = (base + i);
+    out[id] = id;
+  }
+}
diff --git a/loop.cl b/loop_lt.cl
index 3379e9e..d52c5a8 100644
--- a/loop.cl
+++ b/loop_lt.cl
@@ -1,4 +1,4 @@
-__kernel void loop(__global int * out, int iterations)
+__kernel void loop_lt(__global int * out, int iterations)
 {
   unsigned i;
   unsigned base = get_global_id(0) * iterations;
diff --git a/run_tests.sh b/run_tests.sh
index 36c157c..e5be8ea 100644
--- a/run_tests.sh
+++ b/run_tests.sh
@@ -17,6 +17,10 @@ run_test () {
   fi
 }
 
+clear_memory() {
+  `./memset 0 $1 1 1 &>/dev/null`
+}
+
 ################################################################################
 #Integer Add                                                                   #
 ################################################################################
@@ -107,8 +111,10 @@ run_test "./math-int mod_four 5 4 1"
 ###############################################################################
 
 #One work group
+clear_memory 100
 run_test "./get-global-id 100 100"
 #One work item per work group
+clear_memory 100
 run_test "./get-global-id 250 1"
 #Large number of threads in one work group
 #run_test "./get-global-id 1000000 1000000"
@@ -203,4 +209,17 @@ run_test "./math-int if_ne 31 -31 1"
 #False
 run_test "./math-int if_ne 3 3 0"
 
+################################################################################
+#loop tests                                                                    #
+################################################################################
+
+clear_memory 100
+run_test "./loop loop_lt 10"
+clear_memory 100
+run_test "./loop loop_le 10"
+clear_memory 100
+run_test "./loop loop_gt 10"
+clear_memory 100
+run_test "./loop loop_ge 10"
+
 echo "$PASS passes, $FAIL fails"