[gegl] opencl: fast bilateral filter produces right results.
- From: Victor Matheus de Araujo Oliveira <vmaolive src gnome org>
- To: commits-list gnome org
- Cc:
- Subject: [gegl] opencl: fast bilateral filter produces right results.
- Date: Wed, 20 Mar 2013 00:57:51 +0000 (UTC)
commit 7fb0c0515f62b60b1b3bc66131924767fe4918e9
Author: Victor Oliveira <victormatheus gmail com>
Date: Tue Mar 19 21:55:38 2013 -0300
opencl: fast bilateral filter produces right results.
gegl/opencl/gegl-cl-init.c | 6 +
opencl/bilateral-filter-fast.cl | 254 +++++++++++++++-------------
opencl/bilateral-filter-fast.cl.h | 254 +++++++++++++++-------------
operations/Makefile-common.am | 3 +-
operations/common/bilateral-filter-fast.c | 88 ++++-------
5 files changed, 309 insertions(+), 296 deletions(-)
---
diff --git a/gegl/opencl/gegl-cl-init.c b/gegl/opencl/gegl-cl-init.c
index 58d04c3..0406b19 100644
--- a/gegl/opencl/gegl-cl-init.c
+++ b/gegl/opencl/gegl-cl-init.c
@@ -455,4 +455,10 @@ gegl_cl_compile_and_build (const char *program_source, const char *kernel_name[]
return cl_data;
}
+#define CL_BUILD(SOURCE, ...) \
+ { \
+ const char *kernel_name[] = {__VA_ARGS__ , NULL}; \
+ return cl_compile_and_build(SOURCE, kernel_name); \
+ } \
+
#undef CL_SAFE_CALL
diff --git a/opencl/bilateral-filter-fast.cl b/opencl/bilateral-filter-fast.cl
index 56eca2e..0205d8f 100644
--- a/opencl/bilateral-filter-fast.cl
+++ b/opencl/bilateral-filter-fast.cl
@@ -18,51 +18,70 @@
#define GRID(x,y,z) grid[x+sw*(y + z * sh)]
-__kernel void bilateral_init(__global float8 *grid,
- int sw,
- int sh,
- int depth)
-{
- const int gid_x = get_global_id(0);
- const int gid_y = get_global_id(1);
+#define LOCAL_W 8
+#define LOCAL_H 8
- for (int d=0; d<depth; d++)
- {
- GRID(gid_x,gid_y,d) = (float8)(0.0f);
- }
-}
+/* found by trial and error on a NVidia GPU */
+#define DEPTH_CHUNK 12
+__attribute__((reqd_work_group_size(8, 8, 1)))
__kernel void bilateral_downsample(__global const float4 *input,
- __global float2 *grid,
- int width,
- int height,
- int sw,
- int sh,
+ __global float8 *grid,
+ int width,
+ int height,
+ int sw,
+ int sh,
+ int depth,
int s_sigma,
float r_sigma)
{
const int gid_x = get_global_id(0);
const int gid_y = get_global_id(1);
- for (int ry=0; ry < s_sigma; ry++)
- for (int rx=0; rx < s_sigma; rx++)
- {
- const int x = clamp(gid_x * s_sigma - s_sigma/2 + rx, 0, width -1);
- const int y = clamp(gid_y * s_sigma - s_sigma/2 + ry, 0, height-1);
+ __local float8 grid_chunk[DEPTH_CHUNK][LOCAL_H][LOCAL_W];
- const float4 val = input[y * width + x];
+ if (gid_x > sw || gid_y > sh) return;
+
+ for (int d = 0; d < depth; d+=DEPTH_CHUNK)
+ {
+ for (int k=0; k < DEPTH_CHUNK; k++)
+ {
+ grid_chunk[k][get_local_id(1)][get_local_id(0)] = (float8)(0.0f);
+ }
- const int4 z = convert_int4(val * (1.0f/r_sigma) + 0.5f);
+ barrier (CLK_LOCAL_MEM_FENCE);
- grid[4*(gid_x+sw*(gid_y + z.x * sh))+0] += (float2)(val.x, 1.0f);
- grid[4*(gid_x+sw*(gid_y + z.y * sh))+1] += (float2)(val.y, 1.0f);
- grid[4*(gid_x+sw*(gid_y + z.z * sh))+2] += (float2)(val.z, 1.0f);
- grid[4*(gid_x+sw*(gid_y + z.w * sh))+3] += (float2)(val.w, 1.0f);
+ for (int ry=0; ry < s_sigma; ry++)
+ for (int rx=0; rx < s_sigma; rx++)
+ {
+ const int x = clamp(gid_x * s_sigma - s_sigma/2 + rx, 0, width -1);
+ const int y = clamp(gid_y * s_sigma - s_sigma/2 + ry, 0, height-1);
- barrier (CLK_GLOBAL_MEM_FENCE);
- }
+ const float4 val = input[y * width + x];
+
+ const int4 z = convert_int4(val * (1.0f/r_sigma) + 0.5f);
+
+ // z >= d && z < d+DEPTH_CHUNK
+ int4 inbounds = (z >= d & z < d+DEPTH_CHUNK);
+
+ if (inbounds.x) grid_chunk[z.x-d][get_local_id(1)][get_local_id(0)].s01 += (float2)(val.x, 1.0f);
+ if (inbounds.y) grid_chunk[z.y-d][get_local_id(1)][get_local_id(0)].s23 += (float2)(val.y, 1.0f);
+ if (inbounds.z) grid_chunk[z.z-d][get_local_id(1)][get_local_id(0)].s45 += (float2)(val.z, 1.0f);
+ if (inbounds.w) grid_chunk[z.w-d][get_local_id(1)][get_local_id(0)].s67 += (float2)(val.w, 1.0f);
+
+ barrier (CLK_LOCAL_MEM_FENCE);
+ }
+
+ for (int s=d, e=d+min(DEPTH_CHUNK, depth-d); s < e; s++)
+ {
+ grid[gid_x+sw*(gid_y + s * sh)] = grid_chunk[s-d][get_local_id(1)][get_local_id(0)];
+ }
+ }
}
+#undef LOCAL_W
+#undef LOCAL_H
+
#define LOCAL_W 16
#define LOCAL_H 16
@@ -76,9 +95,6 @@ __kernel void bilateral_blur(__global const float8 *grid,
int sh,
int depth)
{
- __local float8 img1[LOCAL_H+2][LOCAL_W+2];
- __local float8 img2[LOCAL_H+2][LOCAL_W+2];
-
const int gid_x = get_global_id(0);
const int gid_y = get_global_id(1);
@@ -89,114 +105,82 @@ __kernel void bilateral_blur(__global const float8 *grid,
float8 vp = (float8)(0.0f);
float8 v = (float8)(0.0f);
- float8 k;
-
- int x = clamp(gid_x, 0, sw-1);
- int y = clamp(gid_y, 0, sw-1);
+ __local float8 data[LOCAL_H+2][LOCAL_W+2];
for (int d=0; d<depth; d++)
{
- int xp = max(x-1, 0);
- int xn = min(x+1, sw-1);
-
- int yp = max(y-1, 0);
- int yn = min(y+1, sh-1);
-
- int zp = max(d-1, 0);
- int zn = min(d+1, depth-1);
-
- /* the corners are not going to be used, don't bother to load them */
-
- if (ly == 0) {
- k = GRID(x, yp, d);
- img1[0][lx+1] = k;
- img2[0][lx+1] = k;
- }
-
- if (ly == LOCAL_H-1) {
- k = GRID(x, yn, d);
- img1[LOCAL_H+1][lx+1] = k;
- img2[LOCAL_H+1][lx+1] = k;
- }
-
- if (lx == 0) {
- k = GRID(xp, y, d);
- img1[ly+1][0] = k;
- img2[ly+1][0] = k;
- }
-
- if (lx == LOCAL_W-1) {
- k = GRID(xn, y, d);
- img1[ly+1][LOCAL_W+1] = k;
- img2[ly+1][LOCAL_W+1] = k;
- }
+ for (int ky=get_local_id(1)-1; ky<LOCAL_H+1; ky+=get_local_size(1))
+ for (int kx=get_local_id(0)-1; kx<LOCAL_W+1; kx+=get_local_size(0))
+ {
+ int xx = clamp((int)get_group_id(0)*LOCAL_W+kx, 0, sw-1);
+ int yy = clamp((int)get_group_id(1)*LOCAL_H+ky, 0, sh-1);
- img1[ly+1][lx+1] = GRID(x, y, d);
+ data[ky+1][kx+1] = GRID(xx, yy, d);
+ }
- barrier (CLK_LOCAL_MEM_FENCE);
+ barrier (CLK_LOCAL_MEM_FENCE);
/* blur x */
- img2[ly+1][lx+1] = (img1[ly+1][lx] + 2.0f * img1[ly+1][lx+1] + img1[ly+1][lx+2]) / 4.0f;
+ data[ly ][lx+1] = (data[ly ][lx] + 2.0f * data[ly ][lx+1] + data[ly ][lx+2]) / 4.0f;
+ data[ly+1][lx+1] = (data[ly+1][lx] + 2.0f * data[ly+1][lx+1] + data[ly+1][lx+2]) / 4.0f;
+ data[ly+2][lx+1] = (data[ly+2][lx] + 2.0f * data[ly+2][lx+1] + data[ly+2][lx+2]) / 4.0f;
- barrier (CLK_LOCAL_MEM_FENCE);
+ barrier (CLK_LOCAL_MEM_FENCE);
/* blur y */
- v = (img2[ly][lx+1] + 2.0f * img2[ly+1][lx+1] + img2[ly+2][lx+1]) / 4.0f;
-
- /* last three v values */
-
- if (d == 0)
- {
- /* this is like CLAMP */
- vpp = img1[ly+1][lx+1];
- vp = img1[ly+1][lx+1];
- }
- else
- {
- vpp = vp;
- vp = v;
+ if (d==0) {
+ v = (data[ly][lx+1] + 2.0f * data[ly+1][lx+1] + data[ly+2][lx+1]) / 4.0f;
+ vpp = v;
+ vp = v;
+ }
+ else {
+ vpp = vp;
+ vp = v;
- float8 blurred = (vpp + 2.0f * vp + v) / 4.0f;
+ v = (data[ly][lx+1] + 2.0f * data[ly+1][lx+1] + data[ly+2][lx+1]) / 4.0f;
- /* XXX: OpenCL 1.1 doesn't support writes to 3D textures */
+ float8 blurred = (vpp + 2.0f * vp + v) / 4.0f;
- if (gid_x < sw && gid_y < sh)
- {
- blurz_r[x+sw*(y+sh*(d-1))] = blurred.s01;
- blurz_g[x+sw*(y+sh*(d-1))] = blurred.s23;
- blurz_b[x+sw*(y+sh*(d-1))] = blurred.s45;
- blurz_a[x+sw*(y+sh*(d-1))] = blurred.s67;
- }
+ if (gid_x < sw && gid_y < sh) {
+ blurz_r[gid_x+sw*(gid_y+sh*(d-1))] = blurred.s01;
+ blurz_g[gid_x+sw*(gid_y+sh*(d-1))] = blurred.s23;
+ blurz_b[gid_x+sw*(gid_y+sh*(d-1))] = blurred.s45;
+ blurz_a[gid_x+sw*(gid_y+sh*(d-1))] = blurred.s67;
}
+ }
}
- /* last z */
-
vpp = vp;
vp = v;
float8 blurred = (vpp + 2.0f * vp + v) / 4.0f;
- if (gid_x < sw && gid_y < sh)
- {
- blurz_r[x+sw*(y+sh*(depth-1))] = blurred.s01;
- blurz_g[x+sw*(y+sh*(depth-1))] = blurred.s23;
- blurz_b[x+sw*(y+sh*(depth-1))] = blurred.s45;
- blurz_a[x+sw*(y+sh*(depth-1))] = blurred.s67;
- }
+ if (gid_x < sw && gid_y < sh) {
+ blurz_r[gid_x+sw*(gid_y+sh*(depth-1))] = blurred.s01;
+ blurz_g[gid_x+sw*(gid_y+sh*(depth-1))] = blurred.s23;
+ blurz_b[gid_x+sw*(gid_y+sh*(depth-1))] = blurred.s45;
+ blurz_a[gid_x+sw*(gid_y+sh*(depth-1))] = blurred.s67;
+ }
}
-const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR;
+#undef LOCAL_W
+#undef LOCAL_H
-__kernel void bilateral_interpolate(__global const float4 *input,
- __read_only image3d_t blurz_r,
- __read_only image3d_t blurz_g,
- __read_only image3d_t blurz_b,
- __read_only image3d_t blurz_a,
- __global float4 *smoothed,
+/* I could use texture memory here, but it's not worth the cost of
+ converting blurz_[r,g,b,a] from buffers to textures */
+
+__kernel void bilateral_interpolate(__global const float4 *input,
+ __global const float2 *blurz_r,
+ __global const float2 *blurz_g,
+ __global const float2 *blurz_b,
+ __global const float2 *blurz_a,
+ __global float4 *smoothed,
int width,
+ int sw,
+ int sh,
+ int depth,
int s_sigma,
float r_sigma)
{
@@ -209,10 +193,44 @@ __kernel void bilateral_interpolate(__global const float4 *input,
float8 val;
- val.s04 = (read_imagef (blurz_r, sampler, (float4)(xf, yf, zf.x, 0.0f))).xy;
- val.s15 = (read_imagef (blurz_g, sampler, (float4)(xf, yf, zf.y, 0.0f))).xy;
- val.s26 = (read_imagef (blurz_b, sampler, (float4)(xf, yf, zf.z, 0.0f))).xy;
- val.s37 = (read_imagef (blurz_a, sampler, (float4)(xf, yf, zf.w, 0.0f))).xy;
+ int x1 = (int)xf;
+ int y1 = (int)yf;
+ int4 z1 = convert_int4(zf);
+
+ int x2 = min(x1+1, sw-1);
+ int y2 = min(y1+1, sh-1);
+ int4 z2 = min(z1+1, depth-1);
+
+ float x_alpha = xf - x1;
+ float y_alpha = yf - y1;
+ float4 z_alpha = zf - floor(zf); /* it's weird that z1 doesn't work here */
+
+ #define BLURZ_R(x,y,z) blurz_r[x+sw*(y+z*sh)]
+ #define BLURZ_G(x,y,z) blurz_g[x+sw*(y+z*sh)]
+ #define BLURZ_B(x,y,z) blurz_b[x+sw*(y+z*sh)]
+ #define BLURZ_A(x,y,z) blurz_a[x+sw*(y+z*sh)]
+
+ /* trilinear interpolation */
+
+ val.s04 = mix(mix(mix(BLURZ_R(x1, y1, z1.x), BLURZ_R(x2, y1, z1.x), x_alpha),
+ mix(BLURZ_R(x1, y2, z1.x), BLURZ_R(x2, y2, z1.x), x_alpha), y_alpha),
+ mix(mix(BLURZ_R(x1, y1, z2.x), BLURZ_R(x2, y1, z2.x), x_alpha),
+ mix(BLURZ_R(x1, y2, z2.x), BLURZ_R(x2, y2, z2.x), x_alpha), y_alpha), z_alpha.x);
+
+ val.s15 = mix(mix(mix(BLURZ_G(x1, y1, z1.y), BLURZ_G(x2, y1, z1.y), x_alpha),
+ mix(BLURZ_G(x1, y2, z1.y), BLURZ_G(x2, y2, z1.y), x_alpha), y_alpha),
+ mix(mix(BLURZ_G(x1, y1, z2.y), BLURZ_G(x2, y1, z2.y), x_alpha),
+ mix(BLURZ_G(x1, y2, z2.y), BLURZ_G(x2, y2, z2.y), x_alpha), y_alpha), z_alpha.y);
+
+ val.s26 = mix(mix(mix(BLURZ_B(x1, y1, z1.z), BLURZ_B(x2, y1, z1.z), x_alpha),
+ mix(BLURZ_B(x1, y2, z1.z), BLURZ_B(x2, y2, z1.z), x_alpha), y_alpha),
+ mix(mix(BLURZ_B(x1, y1, z2.z), BLURZ_B(x2, y1, z2.z), x_alpha),
+ mix(BLURZ_B(x1, y2, z2.z), BLURZ_B(x2, y2, z2.z), x_alpha), y_alpha), z_alpha.z);
+
+ val.s37 = mix(mix(mix(BLURZ_A(x1, y1, z1.w), BLURZ_A(x2, y1, z1.w), x_alpha),
+ mix(BLURZ_A(x1, y2, z1.w), BLURZ_A(x2, y2, z1.w), x_alpha), y_alpha),
+ mix(mix(BLURZ_A(x1, y1, z2.w), BLURZ_A(x2, y1, z2.w), x_alpha),
+ mix(BLURZ_A(x1, y2, z2.w), BLURZ_A(x2, y2, z2.w), x_alpha), y_alpha), z_alpha.w);
smoothed[y * width + x] = val.s0123/val.s4567;
}
diff --git a/opencl/bilateral-filter-fast.cl.h b/opencl/bilateral-filter-fast.cl.h
index 7648245..2d3072a 100644
--- a/opencl/bilateral-filter-fast.cl.h
+++ b/opencl/bilateral-filter-fast.cl.h
@@ -19,51 +19,70 @@ static const char* bilateral_filter_fast_cl_source =
" \n"
"#define GRID(x,y,z) grid[x+sw*(y + z * sh)] \n"
" \n"
-"__kernel void bilateral_init(__global float8 *grid, \n"
-" int sw, \n"
-" int sh, \n"
-" int depth) \n"
-"{ \n"
-" const int gid_x = get_global_id(0); \n"
-" const int gid_y = get_global_id(1); \n"
+"#define LOCAL_W 8 \n"
+"#define LOCAL_H 8 \n"
" \n"
-" for (int d=0; d<depth; d++) \n"
-" { \n"
-" GRID(gid_x,gid_y,d) = (float8)(0.0f); \n"
-" } \n"
-"} \n"
+"/* found by trial and error on a NVidia GPU */ \n"
+"#define DEPTH_CHUNK 12 \n"
" \n"
+"__attribute__((reqd_work_group_size(8, 8, 1))) \n"
"__kernel void bilateral_downsample(__global const float4 *input, \n"
-" __global float2 *grid, \n"
-" int width, \n"
-" int height, \n"
-" int sw, \n"
-" int sh, \n"
+" __global float8 *grid, \n"
+" int width, \n"
+" int height, \n"
+" int sw, \n"
+" int sh, \n"
+" int depth, \n"
" int s_sigma, \n"
" float r_sigma) \n"
"{ \n"
" const int gid_x = get_global_id(0); \n"
" const int gid_y = get_global_id(1); \n"
" \n"
-" for (int ry=0; ry < s_sigma; ry++) \n"
-" for (int rx=0; rx < s_sigma; rx++) \n"
-" { \n"
-" const int x = clamp(gid_x * s_sigma - s_sigma/2 + rx, 0, width -1); \n"
-" const int y = clamp(gid_y * s_sigma - s_sigma/2 + ry, 0, height-1); \n"
+" __local float8 grid_chunk[DEPTH_CHUNK][LOCAL_H][LOCAL_W]; \n"
" \n"
-" const float4 val = input[y * width + x]; \n"
+" if (gid_x > sw || gid_y > sh) return; \n"
+" \n"
+" for (int d = 0; d < depth; d+=DEPTH_CHUNK) \n"
+" { \n"
+" for (int k=0; k < DEPTH_CHUNK; k++) \n"
+" { \n"
+" grid_chunk[k][get_local_id(1)][get_local_id(0)] = (float8)(0.0f); \n"
+" } \n"
" \n"
-" const int4 z = convert_int4(val * (1.0f/r_sigma) + 0.5f); \n"
+" barrier (CLK_LOCAL_MEM_FENCE); \n"
" \n"
-" grid[4*(gid_x+sw*(gid_y + z.x * sh))+0] += (float2)(val.x, 1.0f); \n"
-" grid[4*(gid_x+sw*(gid_y + z.y * sh))+1] += (float2)(val.y, 1.0f); \n"
-" grid[4*(gid_x+sw*(gid_y + z.z * sh))+2] += (float2)(val.z, 1.0f); \n"
-" grid[4*(gid_x+sw*(gid_y + z.w * sh))+3] += (float2)(val.w, 1.0f); \n"
+" for (int ry=0; ry < s_sigma; ry++) \n"
+" for (int rx=0; rx < s_sigma; rx++) \n"
+" { \n"
+" const int x = clamp(gid_x * s_sigma - s_sigma/2 + rx, 0, width -1);\n"
+" const int y = clamp(gid_y * s_sigma - s_sigma/2 + ry, 0, height-1);\n"
" \n"
-" barrier (CLK_GLOBAL_MEM_FENCE); \n"
-" } \n"
+" const float4 val = input[y * width + x]; \n"
+" \n"
+" const int4 z = convert_int4(val * (1.0f/r_sigma) + 0.5f); \n"
+" \n"
+" // z >= d && z < d+DEPTH_CHUNK \n"
+" int4 inbounds = (z >= d & z < d+DEPTH_CHUNK); \n"
+" \n"
+" if (inbounds.x) grid_chunk[z.x-d][get_local_id(1)][get_local_id(0)].s01 += (float2)(val.x,
1.0f);\n"
+" if (inbounds.y) grid_chunk[z.y-d][get_local_id(1)][get_local_id(0)].s23 += (float2)(val.y,
1.0f);\n"
+" if (inbounds.z) grid_chunk[z.z-d][get_local_id(1)][get_local_id(0)].s45 += (float2)(val.z,
1.0f);\n"
+" if (inbounds.w) grid_chunk[z.w-d][get_local_id(1)][get_local_id(0)].s67 += (float2)(val.w,
1.0f);\n"
+" \n"
+" barrier (CLK_LOCAL_MEM_FENCE); \n"
+" } \n"
+" \n"
+" for (int s=d, e=d+min(DEPTH_CHUNK, depth-d); s < e; s++) \n"
+" { \n"
+" grid[gid_x+sw*(gid_y + s * sh)] = grid_chunk[s-d][get_local_id(1)][get_local_id(0)];\n"
+" } \n"
+" } \n"
"} \n"
" \n"
+"#undef LOCAL_W \n"
+"#undef LOCAL_H \n"
+" \n"
"#define LOCAL_W 16 \n"
"#define LOCAL_H 16 \n"
" \n"
@@ -77,9 +96,6 @@ static const char* bilateral_filter_fast_cl_source =
" int sh, \n"
" int depth) \n"
"{ \n"
-" __local float8 img1[LOCAL_H+2][LOCAL_W+2]; \n"
-" __local float8 img2[LOCAL_H+2][LOCAL_W+2]; \n"
-" \n"
" const int gid_x = get_global_id(0); \n"
" const int gid_y = get_global_id(1); \n"
" \n"
@@ -90,114 +106,82 @@ static const char* bilateral_filter_fast_cl_source =
" float8 vp = (float8)(0.0f); \n"
" float8 v = (float8)(0.0f); \n"
" \n"
-" float8 k; \n"
-" \n"
-" int x = clamp(gid_x, 0, sw-1); \n"
-" int y = clamp(gid_y, 0, sw-1); \n"
+" __local float8 data[LOCAL_H+2][LOCAL_W+2]; \n"
" \n"
" for (int d=0; d<depth; d++) \n"
" { \n"
-" int xp = max(x-1, 0); \n"
-" int xn = min(x+1, sw-1); \n"
-" \n"
-" int yp = max(y-1, 0); \n"
-" int yn = min(y+1, sh-1); \n"
-" \n"
-" int zp = max(d-1, 0); \n"
-" int zn = min(d+1, depth-1); \n"
-" \n"
-" /* the corners are not going to be used, don't bother to load them */ \n"
-" \n"
-" if (ly == 0) { \n"
-" k = GRID(x, yp, d); \n"
-" img1[0][lx+1] = k; \n"
-" img2[0][lx+1] = k; \n"
-" } \n"
-" \n"
-" if (ly == LOCAL_H-1) { \n"
-" k = GRID(x, yn, d); \n"
-" img1[LOCAL_H+1][lx+1] = k; \n"
-" img2[LOCAL_H+1][lx+1] = k; \n"
-" } \n"
-" \n"
-" if (lx == 0) { \n"
-" k = GRID(xp, y, d); \n"
-" img1[ly+1][0] = k; \n"
-" img2[ly+1][0] = k; \n"
-" } \n"
-" \n"
-" if (lx == LOCAL_W-1) { \n"
-" k = GRID(xn, y, d); \n"
-" img1[ly+1][LOCAL_W+1] = k; \n"
-" img2[ly+1][LOCAL_W+1] = k; \n"
-" } \n"
+" for (int ky=get_local_id(1)-1; ky<LOCAL_H+1; ky+=get_local_size(1)) \n"
+" for (int kx=get_local_id(0)-1; kx<LOCAL_W+1; kx+=get_local_size(0)) \n"
+" { \n"
+" int xx = clamp((int)get_group_id(0)*LOCAL_W+kx, 0, sw-1); \n"
+" int yy = clamp((int)get_group_id(1)*LOCAL_H+ky, 0, sh-1); \n"
" \n"
-" img1[ly+1][lx+1] = GRID(x, y, d); \n"
+" data[ky+1][kx+1] = GRID(xx, yy, d); \n"
+" } \n"
" \n"
-" barrier (CLK_LOCAL_MEM_FENCE); \n"
+" barrier (CLK_LOCAL_MEM_FENCE); \n"
" \n"
" /* blur x */ \n"
" \n"
-" img2[ly+1][lx+1] = (img1[ly+1][lx] + 2.0f * img1[ly+1][lx+1] + img1[ly+1][lx+2]) / 4.0f;\n"
+" data[ly ][lx+1] = (data[ly ][lx] + 2.0f * data[ly ][lx+1] + data[ly ][lx+2]) / 4.0f;\n"
+" data[ly+1][lx+1] = (data[ly+1][lx] + 2.0f * data[ly+1][lx+1] + data[ly+1][lx+2]) / 4.0f;\n"
+" data[ly+2][lx+1] = (data[ly+2][lx] + 2.0f * data[ly+2][lx+1] + data[ly+2][lx+2]) / 4.0f;\n"
" \n"
-" barrier (CLK_LOCAL_MEM_FENCE); \n"
+" barrier (CLK_LOCAL_MEM_FENCE); \n"
" \n"
" /* blur y */ \n"
" \n"
-" v = (img2[ly][lx+1] + 2.0f * img2[ly+1][lx+1] + img2[ly+2][lx+1]) / 4.0f;\n"
-" \n"
-" /* last three v values */ \n"
-" \n"
-" if (d == 0) \n"
-" { \n"
-" /* this is like CLAMP */ \n"
-" vpp = img1[ly+1][lx+1]; \n"
-" vp = img1[ly+1][lx+1]; \n"
-" } \n"
-" else \n"
-" { \n"
-" vpp = vp; \n"
-" vp = v; \n"
+" if (d==0) { \n"
+" v = (data[ly][lx+1] + 2.0f * data[ly+1][lx+1] + data[ly+2][lx+1]) / 4.0f;\n"
+" vpp = v; \n"
+" vp = v; \n"
+" } \n"
+" else { \n"
+" vpp = vp; \n"
+" vp = v; \n"
" \n"
-" float8 blurred = (vpp + 2.0f * vp + v) / 4.0f; \n"
+" v = (data[ly][lx+1] + 2.0f * data[ly+1][lx+1] + data[ly+2][lx+1]) / 4.0f;\n"
" \n"
-" /* XXX: OpenCL 1.1 doesn't support writes to 3D textures */ \n"
+" float8 blurred = (vpp + 2.0f * vp + v) / 4.0f; \n"
" \n"
-" if (gid_x < sw && gid_y < sh) \n"
-" { \n"
-" blurz_r[x+sw*(y+sh*(d-1))] = blurred.s01; \n"
-" blurz_g[x+sw*(y+sh*(d-1))] = blurred.s23; \n"
-" blurz_b[x+sw*(y+sh*(d-1))] = blurred.s45; \n"
-" blurz_a[x+sw*(y+sh*(d-1))] = blurred.s67; \n"
-" } \n"
+" if (gid_x < sw && gid_y < sh) { \n"
+" blurz_r[gid_x+sw*(gid_y+sh*(d-1))] = blurred.s01; \n"
+" blurz_g[gid_x+sw*(gid_y+sh*(d-1))] = blurred.s23; \n"
+" blurz_b[gid_x+sw*(gid_y+sh*(d-1))] = blurred.s45; \n"
+" blurz_a[gid_x+sw*(gid_y+sh*(d-1))] = blurred.s67; \n"
" } \n"
+" } \n"
" } \n"
" \n"
-" /* last z */ \n"
-" \n"
" vpp = vp; \n"
" vp = v; \n"
" \n"
" float8 blurred = (vpp + 2.0f * vp + v) / 4.0f; \n"
" \n"
-" if (gid_x < sw && gid_y < sh) \n"
-" { \n"
-" blurz_r[x+sw*(y+sh*(depth-1))] = blurred.s01; \n"
-" blurz_g[x+sw*(y+sh*(depth-1))] = blurred.s23; \n"
-" blurz_b[x+sw*(y+sh*(depth-1))] = blurred.s45; \n"
-" blurz_a[x+sw*(y+sh*(depth-1))] = blurred.s67; \n"
-" } \n"
+" if (gid_x < sw && gid_y < sh) { \n"
+" blurz_r[gid_x+sw*(gid_y+sh*(depth-1))] = blurred.s01; \n"
+" blurz_g[gid_x+sw*(gid_y+sh*(depth-1))] = blurred.s23; \n"
+" blurz_b[gid_x+sw*(gid_y+sh*(depth-1))] = blurred.s45; \n"
+" blurz_a[gid_x+sw*(gid_y+sh*(depth-1))] = blurred.s67; \n"
+" } \n"
"} \n"
" \n"
-"const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR;\n"
+"#undef LOCAL_W \n"
+"#undef LOCAL_H \n"
" \n"
-"__kernel void bilateral_interpolate(__global const float4 *input, \n"
-" __read_only image3d_t blurz_r, \n"
-" __read_only image3d_t blurz_g, \n"
-" __read_only image3d_t blurz_b, \n"
-" __read_only image3d_t blurz_a, \n"
-" __global float4 *smoothed, \n"
+"/* I could use texture memory here, but it's not worth the cost of \n"
+" converting blurz_[r,g,b,a] from buffers to textures */ \n"
+" \n"
+"__kernel void bilateral_interpolate(__global const float4 *input, \n"
+" __global const float2 *blurz_r, \n"
+" __global const float2 *blurz_g, \n"
+" __global const float2 *blurz_b, \n"
+" __global const float2 *blurz_a, \n"
+" __global float4 *smoothed, \n"
" int width, \n"
+" int sw, \n"
+" int sh, \n"
+" int depth, \n"
" int s_sigma, \n"
" float r_sigma) \n"
"{ \n"
@@ -210,10 +194,44 @@ static const char* bilateral_filter_fast_cl_source =
" \n"
" float8 val; \n"
" \n"
-" val.s04 = (read_imagef (blurz_r, sampler, (float4)(xf, yf, zf.x, 0.0f))).xy;\n"
-" val.s15 = (read_imagef (blurz_g, sampler, (float4)(xf, yf, zf.y, 0.0f))).xy;\n"
-" val.s26 = (read_imagef (blurz_b, sampler, (float4)(xf, yf, zf.z, 0.0f))).xy;\n"
-" val.s37 = (read_imagef (blurz_a, sampler, (float4)(xf, yf, zf.w, 0.0f))).xy;\n"
+" int x1 = (int)xf; \n"
+" int y1 = (int)yf; \n"
+" int4 z1 = convert_int4(zf); \n"
+" \n"
+" int x2 = min(x1+1, sw-1); \n"
+" int y2 = min(y1+1, sh-1); \n"
+" int4 z2 = min(z1+1, depth-1); \n"
+" \n"
+" float x_alpha = xf - x1; \n"
+" float y_alpha = yf - y1; \n"
+" float4 z_alpha = zf - floor(zf); /* it's weird that z1 doesn't work here */ \n"
+" \n"
+" #define BLURZ_R(x,y,z) blurz_r[x+sw*(y+z*sh)] \n"
+" #define BLURZ_G(x,y,z) blurz_g[x+sw*(y+z*sh)] \n"
+" #define BLURZ_B(x,y,z) blurz_b[x+sw*(y+z*sh)] \n"
+" #define BLURZ_A(x,y,z) blurz_a[x+sw*(y+z*sh)] \n"
+" \n"
+" /* trilinear interpolation */ \n"
+" \n"
+" val.s04 = mix(mix(mix(BLURZ_R(x1, y1, z1.x), BLURZ_R(x2, y1, z1.x), x_alpha),\n"
+" mix(BLURZ_R(x1, y2, z1.x), BLURZ_R(x2, y2, z1.x), x_alpha), y_alpha),\n"
+" mix(mix(BLURZ_R(x1, y1, z2.x), BLURZ_R(x2, y1, z2.x), x_alpha),\n"
+" mix(BLURZ_R(x1, y2, z2.x), BLURZ_R(x2, y2, z2.x), x_alpha), y_alpha), z_alpha.x);\n"
+" \n"
+" val.s15 = mix(mix(mix(BLURZ_G(x1, y1, z1.y), BLURZ_G(x2, y1, z1.y), x_alpha),\n"
+" mix(BLURZ_G(x1, y2, z1.y), BLURZ_G(x2, y2, z1.y), x_alpha), y_alpha),\n"
+" mix(mix(BLURZ_G(x1, y1, z2.y), BLURZ_G(x2, y1, z2.y), x_alpha),\n"
+" mix(BLURZ_G(x1, y2, z2.y), BLURZ_G(x2, y2, z2.y), x_alpha), y_alpha), z_alpha.y);\n"
+" \n"
+" val.s26 = mix(mix(mix(BLURZ_B(x1, y1, z1.z), BLURZ_B(x2, y1, z1.z), x_alpha),\n"
+" mix(BLURZ_B(x1, y2, z1.z), BLURZ_B(x2, y2, z1.z), x_alpha), y_alpha),\n"
+" mix(mix(BLURZ_B(x1, y1, z2.z), BLURZ_B(x2, y1, z2.z), x_alpha),\n"
+" mix(BLURZ_B(x1, y2, z2.z), BLURZ_B(x2, y2, z2.z), x_alpha), y_alpha), z_alpha.z);\n"
+" \n"
+" val.s37 = mix(mix(mix(BLURZ_A(x1, y1, z1.w), BLURZ_A(x2, y1, z1.w), x_alpha),\n"
+" mix(BLURZ_A(x1, y2, z1.w), BLURZ_A(x2, y2, z1.w), x_alpha), y_alpha),\n"
+" mix(mix(BLURZ_A(x1, y1, z2.w), BLURZ_A(x2, y1, z2.w), x_alpha),\n"
+" mix(BLURZ_A(x1, y2, z2.w), BLURZ_A(x2, y2, z2.w), x_alpha), y_alpha), z_alpha.w);\n"
" \n"
" smoothed[y * width + x] = val.s0123/val.s4567; \n"
"} \n"
diff --git a/operations/Makefile-common.am b/operations/Makefile-common.am
index 184bed6..c641f58 100644
--- a/operations/Makefile-common.am
+++ b/operations/Makefile-common.am
@@ -7,7 +7,8 @@ libgegl = $(top_builddir)/gegl/libgegl-$(GEGL_API_VERSION).la $(BABL_LIBS)
op_libs = $(DEP_LIBS) $(libgegl) $(OPENCL_LIBS)
GEGLHEADERS = $(wildcard $(top_srcdir)/gegl/*.h)\
- $(wildcard $(top_srcdir)/gegl/buffer/*.h)
+ $(wildcard $(top_srcdir)/gegl/buffer/*.h)\
+ $(wildcard $(top_srcdir)/gegl/opencl/*.h)
AM_CPPFLAGS = \
-I$(top_srcdir) \
diff --git a/operations/common/bilateral-filter-fast.c b/operations/common/bilateral-filter-fast.c
index 28edc0c..d2fea54 100644
--- a/operations/common/bilateral-filter-fast.c
+++ b/operations/common/bilateral-filter-fast.c
@@ -297,12 +297,8 @@ cl_bilateral (cl_mem in_tex,
cl_mem grid = NULL;
cl_mem blur[4] = {NULL, NULL, NULL, NULL};
- cl_mem blur_tex[4] = {NULL, NULL, NULL, NULL};
-
- cl_image_format format = {CL_RG, CL_FLOAT};
GEGL_CL_BUILD(bilateral_filter_fast,
- "bilateral_init",
"bilateral_downsample",
"bilateral_blur",
"bilateral_interpolate")
@@ -317,56 +313,34 @@ cl_bilateral (cl_mem in_tex,
for(c = 0; c < 4; c++)
{
blur[c] = gegl_clCreateBuffer (gegl_cl_get_context (),
- CL_MEM_WRITE_ONLY,
+ CL_MEM_READ_WRITE,
sw * sh * depth * sizeof(cl_float2),
NULL, &cl_err);
CL_CHECK;
-
- blur_tex[c] = gegl_clCreateImage3D (gegl_cl_get_context (),
- CL_MEM_READ_ONLY,
- &format,
- sw, sh, depth,
- 0, 0, NULL, &cl_err);
- CL_CHECK;
}
{
- global_ws[0] = sw;
- global_ws[1] = sh;
+ local_ws[0] = 8;
+ local_ws[1] = 8;
- GEGL_CL_ARG_START(cl_data->kernel[0])
- GEGL_CL_ARG(cl_mem, grid)
- GEGL_CL_ARG(cl_int, sw)
- GEGL_CL_ARG(cl_int, sh)
- GEGL_CL_ARG(cl_int, depth)
- GEGL_CL_ARG_END
+ global_ws[0] = ((sw + local_ws[0] - 1)/local_ws[0])*local_ws[0];
+ global_ws[1] = ((sh + local_ws[1] - 1)/local_ws[1])*local_ws[1];
- cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
- cl_data->kernel[0], 2,
- NULL, global_ws, NULL,
- 0, NULL, NULL);
- CL_CHECK;
- }
-
-
- {
- global_ws[0] = sw;
- global_ws[1] = sh;
-
- GEGL_CL_ARG_START(cl_data->kernel[1])
+ GEGL_CL_ARG_START(cl_data->kernel[0])
GEGL_CL_ARG(cl_mem, in_tex)
GEGL_CL_ARG(cl_mem, grid)
GEGL_CL_ARG(cl_int, width)
GEGL_CL_ARG(cl_int, height)
GEGL_CL_ARG(cl_int, sw)
GEGL_CL_ARG(cl_int, sh)
+ GEGL_CL_ARG(cl_int, depth)
GEGL_CL_ARG(cl_int, s_sigma)
GEGL_CL_ARG(cl_float, r_sigma)
GEGL_CL_ARG_END
cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
- cl_data->kernel[1], 2,
- NULL, global_ws, NULL,
+ cl_data->kernel[0], 2,
+ NULL, global_ws, local_ws,
0, NULL, NULL);
CL_CHECK;
}
@@ -375,10 +349,10 @@ cl_bilateral (cl_mem in_tex,
local_ws[0] = 16;
local_ws[1] = 16;
- global_ws[0] = ((sw + 15)/16)*16;
- global_ws[1] = ((sh + 15)/16)*16;
+ global_ws[0] = ((sw + local_ws[0] - 1)/local_ws[0])*local_ws[0];
+ global_ws[1] = ((sh + local_ws[1] - 1)/local_ws[1])*local_ws[1];
- GEGL_CL_ARG_START(cl_data->kernel[2])
+ GEGL_CL_ARG_START(cl_data->kernel[1])
GEGL_CL_ARG(cl_mem, grid)
GEGL_CL_ARG(cl_mem, blur[0])
GEGL_CL_ARG(cl_mem, blur[1])
@@ -390,43 +364,33 @@ cl_bilateral (cl_mem in_tex,
GEGL_CL_ARG_END
cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
- cl_data->kernel[2], 2,
+ cl_data->kernel[1], 2,
NULL, global_ws, local_ws,
0, NULL, NULL);
CL_CHECK;
}
- for(c = 0; c < 4; c++)
- {
- const size_t dst_origin[3] = {0, 0, 0};
- const size_t dst_region[3] = {sw, sh, depth};
-
- cl_err = gegl_clEnqueueCopyBufferToImage (gegl_cl_get_command_queue (),
- blur[c],
- blur_tex[c],
- 0, dst_origin, dst_region,
- 0, NULL, NULL);
- CL_CHECK;
- }
-
{
global_ws[0] = width;
global_ws[1] = height;
- GEGL_CL_ARG_START(cl_data->kernel[3])
+ GEGL_CL_ARG_START(cl_data->kernel[2])
GEGL_CL_ARG(cl_mem, in_tex)
- GEGL_CL_ARG(cl_mem, blur_tex[0])
- GEGL_CL_ARG(cl_mem, blur_tex[1])
- GEGL_CL_ARG(cl_mem, blur_tex[2])
- GEGL_CL_ARG(cl_mem, blur_tex[3])
+ GEGL_CL_ARG(cl_mem, blur[0])
+ GEGL_CL_ARG(cl_mem, blur[1])
+ GEGL_CL_ARG(cl_mem, blur[2])
+ GEGL_CL_ARG(cl_mem, blur[3])
GEGL_CL_ARG(cl_mem, out_tex)
GEGL_CL_ARG(cl_int, width)
+ GEGL_CL_ARG(cl_int, sw)
+ GEGL_CL_ARG(cl_int, sh)
+ GEGL_CL_ARG(cl_int, depth)
GEGL_CL_ARG(cl_int, s_sigma)
GEGL_CL_ARG(cl_float, r_sigma)
GEGL_CL_ARG_END
cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
- cl_data->kernel[3], 2,
+ cl_data->kernel[2], 2,
NULL, global_ws, NULL,
0, NULL, NULL);
CL_CHECK;
@@ -440,12 +404,18 @@ cl_bilateral (cl_mem in_tex,
for(c = 0; c < 4; c++)
{
GEGL_CL_RELEASE(blur[c]);
- GEGL_CL_RELEASE(blur_tex[c]);
}
return FALSE;
error:
+ if (grid) GEGL_CL_RELEASE(grid);
+
+ for(c = 0; c < 4; c++)
+ {
+ if (blur[c]) GEGL_CL_RELEASE(blur[c]);
+ }
+
return TRUE;
}
[
Date Prev][
Date Next] [
Thread Prev][
Thread Next]
[
Thread Index]
[
Date Index]
[
Author Index]
