kat_opencl.c (4400B)
1 /* 2 Copyright 2010-2011, D. E. Shaw Research. 3 All rights reserved. 4 5 Redistribution and use in source and binary forms, with or without 6 modification, are permitted provided that the following conditions are 7 met: 8 9 * Redistributions of source code must retain the above copyright 10 notice, this list of conditions, and the following disclaimer. 11 12 * Redistributions in binary form must reproduce the above copyright 13 notice, this list of conditions, and the following disclaimer in the 14 documentation and/or other materials provided with the distribution. 15 16 * Neither the name of D. E. Shaw Research nor the names of its 17 contributors may be used to endorse or promote products derived from 18 this software without specific prior written permission. 19 20 THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS 21 "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT 22 LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR 23 A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT 24 OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, 25 SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT 26 LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, 27 DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY 28 THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT 29 (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE 30 OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. 31 */ 32 #include "util_opencl.h" 33 #include "kat_main.h" 34 35 // USE_GENCL: gencl.sh is a small shell script 36 // that pre-processes foo.ocl into foo.i, containing 37 // a definition like: 38 // const char opencl_src[] = "preprocessed text of foo.ocl" 39 // Thus, with gencl, this file says 40 // #include <foo.i> 41 // and the binary obtained by compiling it 42 // is fully "baked". Runtime behavior doesn't depend 43 // on the contents of some file (e.g., foo.ocl or some 44 // header that it includes) that might have changed long after this 45 // file was compiled. 46 // 47 // The alternative (USE_GENCL 0) seems to be more along 48 // the lines of what OpenCL designers imagine. It makes the text of the 49 // kernel program the string "#include <foo.c>". This eliminates 50 // the need for the extra machinery in gencl.sh, but runtime 51 // behavior is susceptable to changes in foo.c, or files included 52 // by foo.c long after this file is compiled. It also requires some 53 // hocus pocus to get absolute paths for the -I options needed 54 // to compile the code at runtime. Something like: 55 // override CFLAGS += -DSRCDIR=\"$(dir $(abspath $<)).\" 56 #define USE_GENCL 1 57 58 #if USE_GENCL 59 const char *opencl_src = 60 #include "kat_opencl_kernel.i" 61 ; 62 #else 63 #ifndef SRCDIR 64 #error -DSRCDIR="/absolute/path/to/examples" should have been put on the command-line by GNUmakefile 65 #endif 66 #endif 67 68 void host_execute_tests(kat_instance *tests, unsigned ntests){ 69 UCLInfo *infop; 70 cl_kernel kern; 71 size_t nthreads, tests_sz; 72 cl_mem tests_dev; 73 const char *kernelname = "dev_execute_tests"; 74 cl_int err; 75 76 #if USE_GENCL 77 infop = opencl_init(NULL, opencl_src, ""); 78 #else 79 infop = opencl_init(NULL, "#include <kat_dev_execute.h>", 80 " -I" SRCDIR 81 " -I" SRCDIR "/../include " 82 " -DKAT_KERNEL=__kernel " 83 " -DKAT_GLOBAL=__global "); 84 #endif 85 CHECKERR(kern = clCreateKernel(infop->prog, kernelname, &err)); 86 if (infop->wgsize > 64) infop->wgsize /= 2; 87 nthreads = infop->cores * infop->wgsize; 88 tests_sz = sizeof(*tests) * (ntests+1); // +1 for sentinel test with method==last 89 CHECKERR(tests_dev = clCreateBuffer(infop->ctx, CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR, tests_sz, tests, &err)); 90 CHECK(clEnqueueWriteBuffer(infop->cmdq, tests_dev, CL_TRUE, 0, tests_sz, tests, 0, 0, 0)); 91 CHECK(clSetKernelArg(kern, 0, sizeof(cl_mem), (void*)&tests_dev)); 92 printf("queuing kernel for %lu threads with %lu work group size, %u tests\n", 93 (unsigned long)nthreads, (unsigned long)infop->wgsize, ntests); 94 CHECK(clEnqueueNDRangeKernel(infop->cmdq, kern, 1, 0, &nthreads, &infop->wgsize, 0, 0, 0)); 95 CHECK(clFinish(infop->cmdq)); 96 CHECK(clEnqueueReadBuffer(infop->cmdq, tests_dev, CL_TRUE, 0, tests_sz, tests, 0, 0, 0)); 97 CHECK(clReleaseMemObject(tests_dev)); 98 CHECK(clReleaseKernel(kern)); 99 opencl_done(infop); 100 }