diff options
author | Serge Cohen <serge@chocolatnoir.net> | 2017-01-13 11:31:56 +0100 |
---|---|---|
committer | Matthias Vogelgesang <matthias.vogelgesang@kit.edu> | 2017-01-17 09:40:23 +0100 |
commit | a2d16bd6464d877832efaedb07ee66c9b61c0439 (patch) | |
tree | 2f557ff2e147a9c82ec9c81af4ffb9645c6d8f07 /contrib | |
parent | 3a9a7feb9664945ebdc1edf1b60639eee9e432c1 (diff) | |
download | ufo-filters-a2d16bd6464d877832efaedb07ee66c9b61c0439.tar.gz ufo-filters-a2d16bd6464d877832efaedb07ee66c9b61c0439.tar.bz2 ufo-filters-a2d16bd6464d877832efaedb07ee66c9b61c0439.tar.xz ufo-filters-a2d16bd6464d877832efaedb07ee66c9b61c0439.zip |
Use macros pixel access macros in ocl_1liner.cl
Implements macros in ocl_1liner.cl as more direct shortcuts to current
(work-item) pixel of output and input(s).
Diffstat (limited to 'contrib')
-rw-r--r-- | contrib/sxc/src/kernels/ocl-1liner-skel.cl | 10 | ||||
-rw-r--r-- | contrib/sxc/src/ufo-ocl-1liner-task.c | 15 |
2 files changed, 20 insertions, 5 deletions
diff --git a/contrib/sxc/src/kernels/ocl-1liner-skel.cl b/contrib/sxc/src/kernels/ocl-1liner-skel.cl index 182526e..27cf028 100644 --- a/contrib/sxc/src/kernels/ocl-1liner-skel.cl +++ b/contrib/sxc/src/kernels/ocl-1liner-skel.cl @@ -36,12 +36,17 @@ * modifies a pixel value in =out= that is NOT the one indexed by px_index. */ -#define IMG_VAL(hx,hy,image) image[(hx) + (hy)*sizeX] +#define IMG_VAL(hx,hy,image) (image[(hx) + (hy)*sizeX]) + +// Macros to access input pixel without parenthesis or bracket +%s +// Macro to access output pixel without parenthesis or bracket +#define out_px (out[px_index]) __kernel void ocl_1liner ( // The input(s) -%s +%s // The output __global float *out) { @@ -55,3 +60,4 @@ __global float *out) // And here the one line of /variable/ code : %s; } + diff --git a/contrib/sxc/src/ufo-ocl-1liner-task.c b/contrib/sxc/src/ufo-ocl-1liner-task.c index 1463160..d47e9d2 100644 --- a/contrib/sxc/src/ufo-ocl-1liner-task.c +++ b/contrib/sxc/src/ufo-ocl-1liner-task.c @@ -89,18 +89,27 @@ ufo_ocl_1liner_task_setup (UfoTask *task, kernel_skel = ufo_resources_get_kernel_source (resources, skel_filename, error); + char skel_in_macro[1024]; char skel_in[1024]; + int one_in_macro_size; int one_in_size; + char * skel_in_macro_next = skel_in_macro; char * skel_in_next = skel_in; for ( guint i=0; i != priv->num_inputs; ++i ) { - one_in_size = snprintf(skel_in_next, 1024 - (skel_in_next - skel_in), "__global float *in_%d,\n", i); - if ( 0 >= one_in_size ) + one_in_macro_size = snprintf(skel_in_macro_next, 1024 - (skel_in_macro_next - skel_in_macro), + "#define in_%d_px (in_%d[px_index])\n", i, i); + one_in_size = snprintf(skel_in_next, 1024 - (skel_in_next - skel_in), + "__global float *in_%d,\n", i); + + if ( (0 >= one_in_macro_size) || (0 >= one_in_size) ) goto exit; + + skel_in_macro_next += one_in_macro_size; skel_in_next += one_in_size; } - asprintf(&kernel_src, kernel_skel, skel_in, priv->one_line); + asprintf(&kernel_src, kernel_skel, skel_in_macro, skel_in, priv->one_line); if ( ! priv->be_quiet ) { /* Done with the preparation of the OpenCL sources. */ fprintf(stdout, "Current version of the one-liner OpenCL source code :\n%s\n", kernel_src); |