[gegl] opencl: opencl kernels are now separated files



commit 0ef6ca2de227fbb128987d8ed63decfffbf7f52f
Author: Victor Oliveira <victormatheus gmail com>
Date:   Fri Jan 11 18:47:50 2013 -0200

    opencl: opencl kernels are now separated files
    
    created a new folder in gegl to contain opencl kernels.
    the point is to have kernels as .cl files that are
    included (with #include).
    
    also, added some perl scripts that take .cl sources and
    generate C strings that are fed to the opencl compiler
    in runtime.
    
    Thanks RPG for help with perl.

 Makefile.am                             |    3 ++-
 configure.ac                            |    1 +
 opencl/brightness-contrast.cl           |   12 ++++++++++++
 opencl/brightness-contrast.cl.h         |   14 ++++++++++++++
 opencl/cltostring.pl                    |   27 +++++++++++++++++++++++++++
 opencl/generate-all.pl                  |    6 ++++++
 operations/common/Makefile.am           |    1 +
 operations/common/brightness-contrast.c |   16 ++--------------
 8 files changed, 65 insertions(+), 15 deletions(-)
---
diff --git a/Makefile.am b/Makefile.am
index 5b78eaa..126d434 100644
--- a/Makefile.am
+++ b/Makefile.am
@@ -11,7 +11,8 @@ SUBDIRS=\
 	tools \
 	examples \
 	tests \
-	po
+	po \
+ 	opencl
 
 if ENABLE_DOCS
 SUBDIRS+= docs
diff --git a/configure.ac b/configure.ac
index 784bae8..7c723e6 100644
--- a/configure.ac
+++ b/configure.ac
@@ -1119,6 +1119,7 @@ tests/xml/Makefile
 tests/xml/data/Makefile
 po/Makefile.in
 gegl-uninstalled.pc
+opencl/Makefile
 ])
 
 # Files with versions in their names
diff --git a/opencl/brightness-contrast.cl b/opencl/brightness-contrast.cl
new file mode 100644
index 0000000..b2b474f
--- /dev/null
+++ b/opencl/brightness-contrast.cl
@@ -0,0 +1,12 @@
+__kernel void gegl_brightness_contrast(__global const float4     *in,
+                                       __global       float4     *out,
+                                       float contrast,
+                                       float brightness)
+{
+  int gid = get_global_id(0);
+  float4 in_v  = in[gid];
+  float4 out_v;
+  out_v.xyz = (in_v.xyz - 0.5f) * contrast + brightness + 0.5f;
+  out_v.w   =  in_v.w;
+  out[gid]  =  out_v;
+}
diff --git a/opencl/brightness-contrast.cl.h b/opencl/brightness-contrast.cl.h
new file mode 100644
index 0000000..c21ef7b
--- /dev/null
+++ b/opencl/brightness-contrast.cl.h
@@ -0,0 +1,14 @@
+static const char* brightnesscontrast_cl_source =
+"__kernel void gegl_brightness_contrast(__global const float4     *in,         \n"
+"                                       __global       float4     *out,        \n"
+"                                       float contrast,                        \n"
+"                                       float brightness)                      \n"
+"{                                                                             \n"
+"  int gid = get_global_id(0);                                                 \n"
+"  float4 in_v  = in[gid];                                                     \n"
+"  float4 out_v;                                                               \n"
+"  out_v.xyz = (in_v.xyz - 0.5f) * contrast + brightness + 0.5f;               \n"
+"  out_v.w   =  in_v.w;                                                        \n"
+"  out[gid]  =  out_v;                                                         \n"
+"}                                                                             \n"
+;
diff --git a/opencl/cltostring.pl b/opencl/cltostring.pl
new file mode 100755
index 0000000..11f1fc0
--- /dev/null
+++ b/opencl/cltostring.pl
@@ -0,0 +1,27 @@
+#!/usr/bin/perl
+
+use File::Basename;
+
+#input file
+$path = $ARGV[0];
+
+open(INPUT_FILE, ${path})
+     or die "Couldn't open ${path} for reading.";
+
+open(OUTPUT_FILE, ">${path}.h");
+
+($name,$dir,$ext) = fileparse($path,'\..*');
+
+$name =~ s/[^a-zA-Z0-9]*//g;
+
+print OUTPUT_FILE "static const char* ${name}_cl_source =\n";
+
+while(<INPUT_FILE>)
+{
+        ~s/\n//;
+        printf( OUTPUT_FILE "\"%-78s\\n\"\n", $_);
+}
+print OUTPUT_FILE ";\n";
+
+close(INPUT_FILE);
+close(OUTPUT_FILE);
diff --git a/opencl/generate-all.pl b/opencl/generate-all.pl
new file mode 100755
index 0000000..1b6d5bd
--- /dev/null
+++ b/opencl/generate-all.pl
@@ -0,0 +1,6 @@
+#!/usr/bin/perl
+
+foreach $f (<*.cl>) {
+  print $f . "\n";
+  system("./cltostring.pl $f");
+}
diff --git a/operations/common/Makefile.am b/operations/common/Makefile.am
index 1f07f66..c106c3b 100644
--- a/operations/common/Makefile.am
+++ b/operations/common/Makefile.am
@@ -3,3 +3,4 @@ include $(top_srcdir)/operations/Makefile-operations.am
 
 AM_CPPFLAGS += \
 	-I$(top_srcdir)/operations/common
+	-I$(top_srcdir)/opencl
diff --git a/operations/common/brightness-contrast.c b/operations/common/brightness-contrast.c
index 3acf092..22ab173 100644
--- a/operations/common/brightness-contrast.c
+++ b/operations/common/brightness-contrast.c
@@ -105,19 +105,7 @@ process (GeglOperation       *op,
   return TRUE;
 }
 
-static const gchar* kernel_source =
-"__kernel void gegl_brightness_contrast(__global const float4     *in,  \n"
-"                                       __global       float4     *out, \n"
-"                                       float contrast,                 \n"
-"                                       float brightness)               \n"
-"{                                                                      \n"
-"  int gid = get_global_id(0);                                          \n"
-"  float4 in_v  = in[gid];                                              \n"
-"  float4 out_v;                                                        \n"
-"  out_v.xyz = (in_v.xyz - 0.5f) * contrast + brightness + 0.5f;        \n"
-"  out_v.w   =  in_v.w;                                                 \n"
-"  out[gid]  =  out_v;                                                  \n"
-"}                                                                      \n";
+#include "opencl/brightness-contrast.cl.h"
 
 /*
  * The class init function sets up information needed for this operations class
@@ -157,7 +145,7 @@ gegl_chant_class_init (GeglChantClass *klass)
       "name",       "gegl:brightness-contrast",
       "categories", "color",
       "description", _("Changes the light level and contrast."),
-      "cl-source"  , kernel_source,
+      "cl-source"  , brightnesscontrast_cl_source,
       "reference-composition", composition,
       NULL);
 }



[Date Prev][Date Next]   [Thread Prev][Thread Next]   [Thread Index] [Date Index] [Author Index]