Add loop example

3 files changed, 73 insertions, 1 deletions

diff --git a/Makefile b/Makefile
index 4530092..43862c1 100644
--- a/Makefile
+++ b/Makefile
@@ -1,6 +1,6 @@
 CFLAGS=-g
 
-all: hello_world math-int get-global-id get-global-id-2d
+all: hello_world math-int get-global-id get-global-id-2d loop
 
 hello_world: hello_world.o util.o
 	gcc -o hello_world $^ -L/usr/local/lib/ -lOpenCL
@@ -13,3 +13,6 @@ get-global-id: get_global_id.o util.o
 
 get-global-id-2d: get_global_id_2d.o util.o
 	gcc -o $@ $^ -L/usr/local/lib/ -lOpenCL
+
+loop: loop.o util.o
+	gcc -o $@ $^ -L/usr/local/lib/ -lOpenCL
diff --git a/loop.c b/loop.c
new file mode 100644
index 0000000..adac26d
--- /dev/null
+++ b/loop.c
@@ -0,0 +1,61 @@
+#include <assert.h>
+#include <stdio.h>
+
+#include <CL/cl.h>
+
+#include "util.h"
+
+int main (int argc, char ** argv)
+{
+   int i,j;
+
+   struct clu_context context;
+
+   cl_int error;
+
+   cl_mem out_buffer;
+   int out_data[100];
+   size_t global_work_size = 10;
+
+   if (!cluSimpleInit(&context, "loop")) {
+      return EXIT_FAILURE;
+   }
+
+   out_buffer = clCreateBuffer(context.cl_ctx,
+                              CL_MEM_WRITE_ONLY,
+                              sizeof(out_data),
+                              NULL, &error);
+
+   assert(error == CL_SUCCESS);
+
+   if (!cluKernelSetArg(context.kernel, 0, sizeof(cl_mem), &out_buffer)) {
+      return EXIT_FAILURE;
+   }
+
+   error = clEnqueueNDRangeKernel(context.command_queue,
+                                  context.kernel,
+                                  1,
+                                  NULL,
+                                  &global_work_size,
+                                  &global_work_size,
+                                  0, NULL, NULL);
+
+   assert(error == CL_SUCCESS);
+
+   error = clEnqueueReadBuffer(context.command_queue,
+                               out_buffer,
+                               CL_TRUE,
+                               0,
+                               sizeof(out_data),
+                               out_data,
+                               0, NULL, NULL);
+
+   assert(error == CL_SUCCESS);
+
+   for (i = 0; i < global_work_size; i++) {
+      for (j = 0; j < 10; j++) {
+         fprintf(stderr, "%2u ", out_data[i * global_work_size + j]);
+      }
+      fprintf(stderr, "\n");
+   }
+}
diff --git a/loop.cl b/loop.cl
new file mode 100644
index 0000000..f19c49d
--- /dev/null
+++ b/loop.cl
@@ -0,0 +1,8 @@
+__kernel void loop(__global int * out)
+{
+  unsigned i;
+  for (i = 0; i < 10; i++) {
+    unsigned id = get_global_id(0) + i;
+    out[id] = id;
+  }
+}