diff options
| -rw-r--r-- | loop.c | 77 | ||||
| -rw-r--r-- | loop_ge.cl | 9 | ||||
| -rw-r--r-- | loop_gt.cl | 9 | ||||
| -rw-r--r-- | loop_le.cl | 9 | ||||
| -rw-r--r-- | loop_lt.cl (renamed from loop.cl) | 2 | ||||
| -rw-r--r-- | run_tests.sh | 19 |
6 files changed, 88 insertions, 37 deletions
@@ -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; + } +} @@ -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" |