summaryrefslogtreecommitdiffstats
path: root/contrib
diff options
context:
space:
mode:
authorSerge Cohen <serge@chocolatnoir.net>2017-01-13 11:31:56 +0100
committerMatthias Vogelgesang <matthias.vogelgesang@kit.edu>2017-01-17 09:40:23 +0100
commita2d16bd6464d877832efaedb07ee66c9b61c0439 (patch)
tree2f557ff2e147a9c82ec9c81af4ffb9645c6d8f07 /contrib
parent3a9a7feb9664945ebdc1edf1b60639eee9e432c1 (diff)
downloadufo-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.cl10
-rw-r--r--contrib/sxc/src/ufo-ocl-1liner-task.c15
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);