Add get-global-id2d

3 files changed, 73 insertions, 1 deletions

diff --git a/Makefile b/Makefile
index a298383..4530092 100644
--- a/Makefile
+++ b/Makefile
@@ -1,6 +1,6 @@
 CFLAGS=-g
 
-all: hello_world math-int get-global-id
+all: hello_world math-int get-global-id get-global-id-2d
 
 hello_world: hello_world.o util.o
 	gcc -o hello_world $^ -L/usr/local/lib/ -lOpenCL
@@ -10,3 +10,6 @@ math-int: math-int.o util.o
 
 get-global-id: get_global_id.o util.o
 	gcc -o get-global-id $^ -L/usr/local/lib/ -lOpenCL
+
+get-global-id-2d: get_global_id_2d.o util.o
+	gcc -o $@ $^ -L/usr/local/lib/ -lOpenCL
diff --git a/get_global_id_2d.c b/get_global_id_2d.c
new file mode 100644
index 0000000..ae940e7
--- /dev/null
+++ b/get_global_id_2d.c
@@ -0,0 +1,60 @@
+#include <assert.h>
+#include <stdio.h>
+
+#include <CL/cl.h>
+
+#include "util.h"
+
+
+int main (int argc, char ** argv)
+{
+   unsigned i,j;
+   cl_int error;
+
+   cl_mem out_buffer;
+   int out_data[10][10];
+   size_t global_work_size[2] = {10, 10};
+
+   struct clu_context context;
+
+   cluSimpleInit(&context, "global_id2d");
+
+   /* XXX: Delete this to see a missing error path */
+   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,
+                                  2, /* dimensions */
+                                  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 < 10; i++) {
+      for (j = 0; j < 10; j++) {
+         fprintf(stderr, "%2u ", out_data[i][j]);
+      }
+      fprintf(stderr, "\n");
+   }
+}
diff --git a/global_id2d.cl b/global_id2d.cl
new file mode 100644
index 0000000..08f9fd8
--- /dev/null
+++ b/global_id2d.cl
@@ -0,0 +1,9 @@
+__kernel void global_id2d(__global int * out)
+{
+   unsigned x = get_global_id(0);
+   unsigned y = get_global_id(1);
+   unsigned id = (y * 10) + x;
+
+   out[id] = id;
+
+}