[gegl] opencl: opencl kernels are now separated files
- From: Victor Matheus de Araujo Oliveira <vmaolive src gnome org>
- To: commits-list gnome org
- Cc:
- Subject: [gegl] opencl: opencl kernels are now separated files
- Date: Fri, 11 Jan 2013 20:54:12 +0000 (UTC)
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]