commit 34f048852a9077015c67b9513d6d67a4517a9071
Author: Victor Oliveira <victormatheus gmail com>
Date: Tue Jun 4 22:57:39 2013 -0300
missing files in previous commit
opencl/edge-laplace.cl.h | 182 ++++++++++++++++++++++++++++++++++++++++++++++
opencl/edge-sobel.cl.h | 73 ++++++++++++++++++
2 files changed, 255 insertions(+), 0 deletions(-)
---
diff --git a/opencl/edge-laplace.cl.h b/opencl/edge-laplace.cl.h
new file mode 100644
index 0000000..bd12a34
--- /dev/null
+++ b/opencl/edge-laplace.cl.h
@@ -0,0 +1,182 @@
+static const char* edge_laplace_cl_source =
+"#define LAPLACE_RADIUS 1 \n"
+"void minmax(float x1, float x2, float x3, \n"
+" float x4, float x5, \n"
+" float *min_result, \n"
+" float *max_result) \n"
+"{ \n"
+" float min1, min2, max1, max2; \n"
+" \n"
+" if (x1 > x2) \n"
+" { \n"
+" max1 = x1; \n"
+" min1 = x2; \n"
+" } \n"
+" else \n"
+" { \n"
+" max1 = x2; \n"
+" min1 = x1; \n"
+" } \n"
+" \n"
+" if (x3 > x4) \n"
+" { \n"
+" max2 = x3; \n"
+" min2 = x4; \n"
+" } \n"
+" else \n"
+" { \n"
+" max2 = x4; \n"
+" min2 = x3; \n"
+" } \n"
+" \n"
+" if (min1 < min2) \n"
+" *min_result = fmin(min1, x5); \n"
+" else \n"
+" *min_result = fmin(min2, x5); \n"
+" if (max1 > max2) \n"
+" *max_result = fmax(max1, x5); \n"
+" else \n"
+" *max_result = fmax(max2, x5); \n"
+"} \n"
+" \n"
+"kernel void pre_edgelaplace (global float4 *in, \n"
+" global float4 *out) \n"
+"{ \n"
+" int gidx = get_global_id(0); \n"
+" int gidy = get_global_id(1); \n"
+" \n"
+" int src_width = get_global_size(0) + LAPLACE_RADIUS * 2; \n"
+" int src_height = get_global_size(1); \n"
+" \n"
+" int i = gidx + LAPLACE_RADIUS, j = gidy + LAPLACE_RADIUS; \n"
+" int gid1d = i + j * src_width; \n"
+" \n"
+" float pix_fl[4] = { \n"
+" in[gid1d - 1 - src_width].x, in[gid1d - 1 - src_width].y, \n"
+" in[gid1d - 1 - src_width].z, in[gid1d - 1 - src_width].w \n"
+" }; \n"
+" float pix_fm[4] = { \n"
+" in[gid1d - src_width].x, in[gid1d - src_width].y, \n"
+" in[gid1d - src_width].z, in[gid1d - src_width].w \n"
+" }; \n"
+" float pix_fr[4] = { \n"
+" in[gid1d + 1 - src_width].x, in[gid1d + 1 - src_width].y, \n"
+" in[gid1d + 1 - src_width].z, in[gid1d + 1 - src_width].w \n"
+" }; \n"
+" float pix_ml[4] = { \n"
+" in[gid1d - 1 ].x, in[gid1d - 1 ].y, \n"
+" in[gid1d - 1 ].z, in[gid1d - 1 ].w \n"
+" }; \n"
+" float pix_mm[4] = { \n"
+" in[gid1d ].x, in[gid1d ].y, \n"
+" in[gid1d ].z, in[gid1d ].w \n"
+" }; \n"
+" float pix_mr[4] = { \n"
+" in[gid1d + 1 ].x, in[gid1d + 1 ].y, \n"
+" in[gid1d + 1 ].z, in[gid1d + 1 ].w \n"
+" }; \n"
+" float pix_bl[4] = { \n"
+" in[gid1d - 1 + src_width].x, in[gid1d - 1 + src_width].y, \n"
+" in[gid1d - 1 + src_width].z, in[gid1d - 1 + src_width].w \n"
+" }; \n"
+" float pix_bm[4] = { \n"
+" in[gid1d + src_width].x, in[gid1d + src_width].y, \n"
+" in[gid1d + src_width].z, in[gid1d + src_width].w \n"
+" }; \n"
+" float pix_br[4] = { \n"
+" in[gid1d + 1 + src_width].x, in[gid1d + 1 + src_width].y, \n"
+" in[gid1d + 1 + src_width].z, in[gid1d + 1 + src_width].w \n"
+" }; \n"
+" \n"
+" int c; \n"
+" float minval, maxval; \n"
+" float gradient[4]; \n"
+" \n"
+" for (c = 0;c < 3; ++c) \n"
+" { \n"
+" minmax(pix_fm[c], pix_bm[c], pix_ml[c], pix_mr[c], \n"
+" pix_mm[c], &minval, &maxval); \n"
+" gradient[c] = 0.5f * \n"
+" fmax((maxval - pix_mm[c]),(pix_mm[c] - minval)); \n"
+" gradient[c] = \n"
+" (pix_fl[c] + pix_fm[c] + pix_fr[c] + \n"
+" pix_ml[c] + pix_mr[c] + pix_bl[c] + \n"
+" pix_bm[c] + pix_br[c] - 8.0f * pix_mm[c]) > \n"
+" 0.0f ? gradient[c] : -1.0f * gradient[c]; \n"
+" } \n"
+" gradient[3] = pix_mm[3]; \n"
+" \n"
+" out[gid1d] = (float4) \n"
+" (gradient[0], gradient[1], gradient[2], gradient[3]); \n"
+"} \n"
+" \n"
+"kernel void knl_edgelaplace (global float4 *in, \n"
+" global float4 *out) \n"
+"{ \n"
+" int gidx = get_global_id(0); \n"
+" int gidy = get_global_id(1); \n"
+" \n"
+" int src_width = get_global_size(0) + LAPLACE_RADIUS * 2; \n"
+" int src_height = get_global_size(1); \n"
+" \n"
+" int i = gidx + LAPLACE_RADIUS, j = gidy + LAPLACE_RADIUS; \n"
+" int gid1d = i + j * src_width; \n"
+" \n"
+" float pix_fl[4] = { \n"
+" in[gid1d - 1 - src_width].x, in[gid1d - 1 - src_width].y, \n"
+" in[gid1d - 1 - src_width].z, in[gid1d - 1 - src_width].w \n"
+" }; \n"
+" float pix_fm[4] = { \n"
+" in[gid1d - src_width].x, in[gid1d - src_width].y, \n"
+" in[gid1d - src_width].z, in[gid1d - src_width].w \n"
+" }; \n"
+" float pix_fr[4] = { \n"
+" in[gid1d + 1 - src_width].x, in[gid1d + 1 - src_width].y, \n"
+" in[gid1d + 1 - src_width].z, in[gid1d + 1 - src_width].w \n"
+" }; \n"
+" float pix_ml[4] = { \n"
+" in[gid1d - 1 ].x, in[gid1d - 1 ].y, \n"
+" in[gid1d - 1 ].z, in[gid1d - 1 ].w \n"
+" }; \n"
+" float pix_mm[4] = { \n"
+" in[gid1d ].x, in[gid1d ].y, \n"
+" in[gid1d ].z, in[gid1d ].w \n"
+" }; \n"
+" float pix_mr[4] = { \n"
+" in[gid1d + 1 ].x, in[gid1d + 1 ].y, \n"
+" in[gid1d + 1 ].z, in[gid1d + 1 ].w \n"
+" }; \n"
+" float pix_bl[4] = { \n"
+" in[gid1d - 1 + src_width].x, in[gid1d - 1 + src_width].y, \n"
+" in[gid1d - 1 + src_width].z, in[gid1d - 1 + src_width].w \n"
+" }; \n"
+" float pix_bm[4] = { \n"
+" in[gid1d + src_width].x, in[gid1d + src_width].y, \n"
+" in[gid1d + src_width].z, in[gid1d + src_width].w \n"
+" }; \n"
+" float pix_br[4] = { \n"
+" in[gid1d + 1 + src_width].x, in[gid1d + 1 + src_width].y, \n"
+" in[gid1d + 1 + src_width].z, in[gid1d + 1 + src_width].w \n"
+" }; \n"
+" \n"
+" int c; \n"
+" float value[4]; \n"
+" \n"
+" for (c = 0;c < 3; ++c) \n"
+" { \n"
+" float current = pix_mm[c]; \n"
+" current = \n"
+" ((current > 0.0f) && \n"
+" (pix_fl[c] < 0.0f || pix_fm[c] < 0.0f || \n"
+" pix_fr[c] < 0.0f || pix_ml[c] < 0.0f || \n"
+" pix_mr[c] < 0.0f || pix_bl[c] < 0.0f || \n"
+" pix_bm[c] < 0.0f || pix_br[c] < 0.0f ) \n"
+" ) ? current : 0.0f; \n"
+" value[c] = current; \n"
+" } \n"
+" value[3] = pix_mm[3]; \n"
+" \n"
+" out[gidx + gidy * get_global_size(0)] = (float4) \n"
+" (value[0], value[1], value[2], value[3]); \n"
+"} \n"
+;
diff --git a/opencl/edge-sobel.cl.h b/opencl/edge-sobel.cl.h
new file mode 100644
index 0000000..6f1816b
--- /dev/null
+++ b/opencl/edge-sobel.cl.h
@@ -0,0 +1,73 @@
+static const char* edge_sobel_cl_source =
+"#define SOBEL_RADIUS 1 \n"
+"kernel void kernel_edgesobel(global float4 *in, \n"
+" global float4 *out, \n"
+" const int horizontal, \n"
+" const int vertical, \n"
+" const int keep_signal, \n"
+" const int has_alpha) \n"
+"{ \n"
+" int gidx = get_global_id(0); \n"
+" int gidy = get_global_id(1); \n"
+" \n"
+" float4 hor_grad = 0.0f; \n"
+" float4 ver_grad = 0.0f; \n"
+" float4 gradient = 0.0f; \n"
+" \n"
+" int dst_width = get_global_size(0); \n"
+" int src_width = dst_width + SOBEL_RADIUS * 2; \n"
+" \n"
+" int i = gidx + SOBEL_RADIUS, j = gidy + SOBEL_RADIUS; \n"
+" int gid1d = i + j * src_width; \n"
+" \n"
+" float4 pix_fl = in[gid1d - 1 - src_width]; \n"
+" float4 pix_fm = in[gid1d - src_width]; \n"
+" float4 pix_fr = in[gid1d + 1 - src_width]; \n"
+" float4 pix_ml = in[gid1d - 1 ]; \n"
+" float4 pix_mm = in[gid1d ]; \n"
+" float4 pix_mr = in[gid1d + 1 ]; \n"
+" float4 pix_bl = in[gid1d - 1 + src_width]; \n"
+" float4 pix_bm = in[gid1d + src_width]; \n"
+" float4 pix_br = in[gid1d + 1 + src_width]; \n"
+" \n"
+" if (horizontal) \n"
+" { \n"
+" hor_grad += \n"
+" - 1.0f * pix_fl + 1.0f * pix_fr \n"
+" - 2.0f * pix_ml + 2.0f * pix_mr \n"
+" - 1.0f * pix_bl + 1.0f * pix_br; \n"
+" } \n"
+" if (vertical) \n"
+" { \n"
+" ver_grad += \n"
+" - 1.0f * pix_fl - 2.0f * pix_fm \n"
+" - 1.0f * pix_fr + 1.0f * pix_bl \n"
+" + 2.0f * pix_bm + 1.0f * pix_br; \n"
+" } \n"
+" \n"
+" if (horizontal && vertical) \n"
+" { \n"
+" gradient = sqrt( \n"
+" hor_grad * hor_grad + \n"
+" ver_grad * ver_grad) / 1.41f; \n"
+" } \n"
+" else \n"
+" { \n"
+" if (keep_signal) \n"
+" gradient = hor_grad + ver_grad; \n"
+" else \n"
+" gradient = fabs(hor_grad + ver_grad); \n"
+" } \n"
+" \n"
+" if (has_alpha) \n"
+" { \n"
+" gradient.w = pix_mm.w; \n"
+" } \n"
+" else \n"
+" { \n"
+" gradient.w = 1.0f; \n"
+" } \n"
+" \n"
+" out[gidx + gidy * dst_width] = gradient; \n"
+"} \n"
+;