From 4df084d9a85a6a43891104be1b7197a3939e16f8 Mon Sep 17 00:00:00 2001 From: Tom Stellard Date: Fri, 11 May 2012 12:21:07 -0400 Subject: Add use-host-ptr and copy-host-ptr programs --- Makefile | 8 ++++- copy_host_ptr.c | 93 +++++++++++++++++++++++++++++++++++++++++++++++++++++++++ memcpy.cl | 6 ++++ use_host_ptr.c | 89 ++++++++++++++++++++++++++++++++++++++++++++++++++++++ 4 files changed, 195 insertions(+), 1 deletion(-) create mode 100644 copy_host_ptr.c create mode 100644 memcpy.cl create mode 100644 use_host_ptr.c 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 +#include +#include + +#include + +#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 +#include + +#include "cl_simple.h" +#include "cl_util.h" + +#include +#include + + +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; +} -- cgit v1.2.3