From 59bc019bea96f7e7b69720cd796947798a710beb Mon Sep 17 00:00:00 2001 From: Tom Stellard Date: Mon, 30 Apr 2012 10:16:30 -0400 Subject: Turn loop into a test and add some loop test cases to run_tests.sh --- loop.c | 77 ++++++++++++++++++++++++++++++++---------------------------- loop.cl | 9 ------- loop_ge.cl | 9 +++++++ loop_gt.cl | 9 +++++++ loop_le.cl | 9 +++++++ loop_lt.cl | 9 +++++++ run_tests.sh | 19 +++++++++++++++ 7 files changed, 96 insertions(+), 45 deletions(-) delete mode 100644 loop.cl create mode 100644 loop_ge.cl create mode 100644 loop_gt.cl create mode 100644 loop_le.cl create mode 100644 loop_lt.cl diff --git a/loop.c b/loop.c index 151dfde..9a806ef 100644 --- a/loop.c +++ b/loop.c @@ -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.cl b/loop.cl deleted file mode 100644 index 3379e9e..0000000 --- a/loop.cl +++ /dev/null @@ -1,9 +0,0 @@ -__kernel void loop(__global int * out, int iterations) -{ - unsigned i; - unsigned base = get_global_id(0) * iterations; - for (i = 0; i < iterations; i++) { - unsigned id = (base + i); - out[id] = id; - } -} 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; + } +} diff --git a/loop_lt.cl b/loop_lt.cl new file mode 100644 index 0000000..d52c5a8 --- /dev/null +++ b/loop_lt.cl @@ -0,0 +1,9 @@ +__kernel void loop_lt(__global int * out, int iterations) +{ + unsigned i; + unsigned base = get_global_id(0) * iterations; + for (i = 0; i < iterations; i++) { + unsigned id = (base + i); + out[id] = id; + } +} 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" -- cgit v1.2.3