Add use-host-ptr and copy-host-ptr programs

4 files changed, 195 insertions, 1 deletions

diff --git a/Makefile b/Makefile
index 543bf66..bc7bf2a 100644
--- a/Makefile
+++ b/Makefile
@@ -4,7 +4,7 @@ COMMON_OBJECTS = cl_simple.o cl_util.o
 LDFLAGS = -L/usr/local/lib -lOpenCL
 
 
-all: hello_world math-int get-global-id get-global-id-2d loop get-global-id-3d mat-mul
+all: hello_world math-int get-global-id get-global-id-2d loop get-global-id-3d mat-mul use-host-ptr copy-host-ptr
 
 hello_world: hello_world.o $(COMMON_OBJECTS)
 	gcc -o $@ $^ $(LDFLAGS)
@@ -29,3 +29,9 @@ mat-mul: mat_mul.o $(COMMON_OBJECTS)
 
 memset: memset.o $(COMMON_OBJECTS)
 	gcc -o $@ $^ $(LDFLAGS)
+
+use-host-ptr: use_host_ptr.o $(COMMON_OBJECTS)
+	gcc -o $@ $^ $(LDFLAGS)
+
+copy-host-ptr: copy_host_ptr.o $(COMMON_OBJECTS)
+	gcc -o $@ $^ $(LDFLAGS)
diff --git a/copy_host_ptr.c b/copy_host_ptr.c
new file mode 100644
index 0000000..f585fb7
--- /dev/null
+++ b/copy_host_ptr.c
@@ -0,0 +1,93 @@
+#include <assert.h>
+#include <stdlib.h>
+#include <stdio.h>
+
+#include <CL/cl.h>
+
+#include "cl_simple.h"
+#include "cl_util.h"
+
+int main (int argc, char ** argv)
+{
+
+   struct cl_simple_context context;
+   cl_mem src_buffer;
+   cl_int error = CL_SUCCESS;
+   unsigned num_items;
+   int * src_data;
+   int * out;
+   unsigned src_size;
+   unsigned i;
+   int copy_host_ptr;
+   cl_mem_flags mem_flags;
+   size_t global_work_size = 1;
+
+   if (argc != 3) {
+      fprintf(stderr, "Usage: copy-host-ptr num_items copy_host_ptr(T/F)\n");
+      return EXIT_FAILURE;
+   }
+
+   num_items = atoi(argv[1]);
+   copy_host_ptr = atoi(argv[2]);
+   src_size = num_items * sizeof(int);
+
+   src_data = malloc(src_size * sizeof(int));
+   out = malloc(src_size * sizeof(int));
+
+   for (i = 0; i < num_items; i++) {
+      src_data[i] = i;
+   }
+
+   if (!clSimpleSimpleInit(&context, "memcpy")) {
+      return EXIT_FAILURE;
+   }
+
+   if (!clSimpleSetOutputBuffer(&context, src_size)) {
+      return EXIT_FAILURE;
+   }
+
+   if (copy_host_ptr) {
+      src_buffer = clCreateBuffer(context.cl_ctx,
+                                  CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
+                                  src_size, src_data, &error);
+      if (error != CL_SUCCESS) {
+         fprintf(stderr, "clCreateBuffer() failed.\n");
+         return EXIT_FAILURE;
+      }
+   } else {
+      if (!clSimpleCreateBuffer(&src_buffer, context.cl_ctx,
+                                CL_MEM_READ_ONLY, src_size)) {
+         return EXIT_FAILURE;
+      }
+      if (!clSimpleEnqueueWriteBuffer(context.command_queue, src_buffer,
+                                      src_size, src_data)) {
+         return EXIT_FAILURE;
+      }
+   }
+
+   if (!clSimpleKernelSetArg(context.kernel, 1, sizeof(cl_mem), &src_buffer)
+      || !clSimpleKernelSetArg(context.kernel, 2, sizeof(int), &num_items)) {
+      return EXIT_FAILURE;
+   }
+
+   if (!clSimpleEnqueueNDRangeKernel(context.command_queue,
+                                     context.kernel, 1, &global_work_size,
+                                     &global_work_size)) {
+      return EXIT_FAILURE;
+   }
+
+   if (!clSimpleReadOutput(&context, out, src_size)) {
+      return EXIT_FAILURE;
+   }
+
+   for (i = 0; i < num_items; i++) {
+      fprintf(stderr, "%4d ", out[i]);
+      if (num_items % 10 == 9)
+         fprintf(stderr, "\n");
+   }
+   fprintf(stderr, "\n");
+
+   return EXIT_SUCCESS;
+
+}
+
diff --git a/memcpy.cl b/memcpy.cl
new file mode 100644
index 0000000..5451cad
--- /dev/null
+++ b/memcpy.cl
@@ -0,0 +1,6 @@
+__kernel void memcpy(__global int * dst, __global int * src, int items)
+{
+   for (int i = 0; i < items; i++) {
+      dst[i] = src[i];
+   }
+}
diff --git a/use_host_ptr.c b/use_host_ptr.c
new file mode 100644
index 0000000..ad4d345
--- /dev/null
+++ b/use_host_ptr.c
@@ -0,0 +1,89 @@
+#include <CL/cl.h>
+#include <CL/cl_platform.h>
+
+#include "cl_simple.h"
+#include "cl_util.h"
+
+#include <stdlib.h>
+#include <stdio.h>
+
+
+int main(int argc, char ** argv)
+{
+   cl_int error;
+   struct cl_simple_context context;
+   int * out;
+   int * new_out;
+   size_t global_work_size = 10;
+   unsigned out_size;
+   int iterations, i, j;
+
+   if (argc !=2 ) {
+      fprintf(stderr, "Usage: use-host-ptr iterations\n");
+      return EXIT_FAILURE;
+   }
+
+   iterations = atoi(argv[1]);
+   out_size = (global_work_size * iterations + 1) * sizeof(int);
+   out = malloc(out_size);
+   out[global_work_size * iterations] = 12345;
+
+   if (!clSimpleSimpleInit(&context, "loop_lt")) {
+      return EXIT_FAILURE;
+   }
+
+   context.out_buffer = clCreateBuffer(context.cl_ctx,
+                                       CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR,
+                                       out_size,
+                                       out,
+                                       &error);
+   if (error != CL_SUCCESS) {
+      fprintf(stderr, "clCreateBuffer() failed: %s\n",
+                       clUtilErrorString(error));
+      return EXIT_FAILURE;
+   }
+
+   if (   !clSimpleKernelSetArg(context.kernel, 0, sizeof(cl_mem),
+                                &context.out_buffer)
+       || !clSimpleKernelSetArg(context.kernel, 1, sizeof(int),
+                                &iterations)) {
+      return EXIT_FAILURE;
+   }
+
+   if (!clSimpleEnqueueNDRangeKernel(context.command_queue,
+                                     context.kernel,
+                                     1, &global_work_size, &global_work_size)) {
+      return EXIT_FAILURE;
+   }
+
+   new_out = clEnqueueMapBuffer(context.command_queue,
+                            context.out_buffer,
+                            CL_TRUE, /* Blocking */
+                            CL_MAP_READ,
+                            0, /* Offset */
+                            out_size, /* cb */
+                            0, NULL, NULL, &error);
+   if (error != CL_SUCCESS) {
+      fprintf(stderr, "clEnqueueMapBuffer() failed: %s\n",
+                      clUtilErrorString(error));
+      return EXIT_FAILURE;
+   }
+
+
+   error = clEnqueueUnmapMemObject(context.command_queue, context.out_buffer,
+                                   new_out, 0, NULL, NULL);
+   if (error != CL_SUCCESS) {
+      fprintf(stderr, "clEnqueueUnmapMemObject() failed: %s\n",
+                      clUtilErrorString(error));
+   }
+
+   for (i = 0; i < global_work_size; i++) {
+      for (j = 0; j < iterations; j++) {
+         unsigned id = (i * iterations) + j;
+         fprintf(stderr, "%2d ", new_out[id]);
+      }
+      fprintf(stderr, "\n");
+   }
+   fprintf(stderr, "Prestore = %d\n", out[global_work_size * iterations]);
+   return EXIT_SUCCESS;
+}