/*
* kernel_gauss_lap_pyramid.cl
* input0
* input1
* output
* window, pos_x, pos_y, width, height
*/
#ifndef PYRAMID_UV
#define PYRAMID_UV 0
#endif
#ifndef CL_PYRAMID_ENABLE_DUMP
#define CL_PYRAMID_ENABLE_DUMP 0
#endif
#ifndef ENABLE_MASK_GAUSS_SCALE
#define ENABLE_MASK_GAUSS_SCALE 0
#endif
#define fixed_pixels 8
#define GAUSS_V_R 2
#define GAUSS_H_R 1
#define COEFF_MID 4
#define zero8 (float8)(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)
__constant const float coeffs[9] = {0.0f, 0.0f, 0.152f, 0.222f, 0.252f, 0.222f, 0.152f, 0.0f, 0.0f};
#define ARG_FORMAT4 "(%.1f,%.1f,%.1f,%.1f)"
#define ARGS4(a) a.s0, a.s1, a.s2, a.s3
#define ARG_FORMAT8 "(%.1f,%.1f,%.1f,%.1f,%.1f,%.1f,%.1f,%.1f)"
#define ARGS8(a) a.s0, a.s1, a.s2, a.s3, a.s4, a.s5, a.s6, a.s7
/*
* input: RGBA-CL_UNSIGNED_INT16
* output_gauss: RGBA-CL_UNSIGNED_INT8
* output_lap:RGBA-CL_UNSIGNED_INT16
* each work-item calc 2 lines
*/
__kernel void
kernel_gauss_scale_transform (
__read_only image2d_t input, int in_offset_x,
__write_only image2d_t output_gauss
#if CL_PYRAMID_ENABLE_DUMP
, __write_only image2d_t dump_orig
#endif
)
{
int g_x = get_global_id (0);
int in_x = g_x + in_offset_x;
int g_y = get_global_id (1) * 4;
const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
int g_out_x = get_global_id (0);
int g_out_y = get_global_id (1) * 2;
#if CL_PYRAMID_ENABLE_DUMP
write_imageui (dump_orig, (int2)(g_x, g_y + 0), read_imageui(input, sampler, (int2)(in_x, g_y)));
write_imageui (dump_orig, (int2)(g_x, g_y + 1), read_imageui(input, sampler, (int2)(in_x, g_y + 1)));
write_imageui (dump_orig, (int2)(g_x, g_y + 2), read_imageui(input, sampler, (int2)(in_x, g_y + 2)));
write_imageui (dump_orig, (int2)(g_x, g_y + 3), read_imageui(input, sampler, (int2)(in_x, g_y + 3)));
#endif
float8 result_pre[2] = {zero8, zero8};
float8 result_next[2] = {zero8, zero8};
float8 result_cur[2] = {zero8, zero8};
float4 final_g[2];
float8 tmp_data;
int i_ver;
#pragma unroll
for (i_ver = -GAUSS_V_R; i_ver <= GAUSS_V_R + 2; i_ver++) {
int cur_g_y = g_y + i_ver;
float coeff0 = coeffs[i_ver + COEFF_MID];
float coeff1 = coeffs[i_ver + COEFF_MID - 2];
tmp_data = convert_float8(as_uchar8(convert_ushort4(read_imageui(input, sampler, (int2)(in_x - 1, cur_g_y)))));
result_pre[0] += tmp_data * coeff0;
result_pre[1] += tmp_data * coeff1;
tmp_data = convert_float8(as_uchar8(convert_ushort4(read_imageui(input, sampler, (int2)(in_x, cur_g_y)))));
result_cur[0] += tmp_data * coeff0;
result_cur[1] += tmp_data * coeff1;
tmp_data = convert_float8(as_uchar8(convert_ushort4(read_imageui(input, sampler, (int2)(in_x + 1, cur_g_y)))));
result_next[1] += tmp_data * coeff1;
result_next[0] += tmp_data * coeff0;
}
int i_line;
#pragma unroll
for (i_line = 0; i_line < 2; ++i_line) {
#if !PYRAMID_UV
final_g[i_line] = result_cur[i_line].even * coeffs[COEFF_MID] +
(float4)(result_pre[i_line].s7, result_cur[i_line].s135) * coeffs[COEFF_MID + 1] +
(float4)(result_pre[i_line].s6, result_cur[i_line].s024) * coeffs[COEFF_MID + 2] +
(float4)(result_cur[i_line].s1357) * coeffs[COEFF_MID + 1] +
(float4)(result_cur[i_line].s246, result_next[i_line].s0) * coeffs[COEFF_MID + 2];
#else
final_g[i_line] = result_cur[i_line].s0145 * coeffs[COEFF_MID] +
(float4)(result_pre[i_line].s67, result_cur[i_line].s23) * coeffs[COEFF_MID + 1] +
(float4)(result_pre[i_line].s45, result_cur[i_line].s01) * coeffs[COEFF_MID + 2] +
(float4)(result_cur[i_line].s2367) * coeffs[COEFF_MID + 1] +
(float4)(result_cur[i_line].s45, result_next[i_line].s01) * coeffs[COEFF_MID + 2];
#endif
final_g[i_line] = clamp (final_g[i_line] + 0.5f, 0.0f, 255.0f);
write_imageui (output_gauss, (int2)(g_out_x, g_out_y + i_line), convert_uint4(final_g[i_line]));
}
}
inline float8
read_scale_y (__read_only image2d_t input, const sampler_t sampler, float2 pos_start, float step_x)
{
float8 data;
data.s0 = read_imagef (input, sampler, pos_start).x;
pos_start.x += step_x;
data.s1 = read_imagef (input, sampler, pos_start).x;
pos_start.x += step_x;
data.s2 = read_imagef (input, sampler, pos_start).x;
pos_start.x += step_x;
data.s3 = read_imagef (input, sampler, pos_start).x;
pos_start.x += step_x;
data.s4 = read_imagef (input, sampler, pos_start).x;
pos_start.x += step_x;
data.s5 = read_imagef (input, sampler, pos_start).x;
pos_start.x += step_x;
data.s6 = read_imagef (input, sampler, pos_start).x;
pos_start.x += step_x;
data.s7 = read_imagef (input, sampler, pos_start).x;
return data;
}
inline float8
read_scale_uv (__read_only image2d_t input, const sampler_t sampler, float2 pos_start, float step_x)
{
float8 data;
data.s01 = read_imagef (input, sampler, pos_start).xy;
pos_start.x += step_x;
data.s23 = read_imagef (input, sampler, pos_start).xy;
pos_start.x += step_x;
data.s45 = read_imagef (input, sampler, pos_start).xy;
pos_start.x += step_x;
data.s67 = read_imagef (input, sampler, pos_start).xy;
return data;
}
/*
* input_gauss: RGBA-CL_UNSIGNED_INT18
* input_lap: RGBA-CL_UNSIGNED_INT16
* output: RGBA-CL_UNSIGNED_INT16
* each work-item calc 2 lines
*/
__kernel void
kernel_gauss_lap_reconstruct (
__read_only image2d_t input_gauss,
float in_sampler_offset_x, float in_sampler_offset_y,
__read_only image2d_t input_lap,
__write_only image2d_t output, int out_offset_x, float out_width, float out_height
#if CL_PYRAMID_ENABLE_DUMP
, __write_only image2d_t dump_resize, __write_only image2d_t dump_final
#endif
)
{
int g_x = get_global_id (0);
int g_y = get_global_id (1);
const sampler_t lap_sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
const sampler_t gauss_sampler = CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR;
//if (g_x > out_width + 0.9f || g_y > out_height + 0.5f)
// return;
float8 lap = convert_float8(as_uchar8(convert_ushort4(read_imageui(input_lap, lap_sampler, (int2)(g_x, g_y)))));
lap = (lap - 128.0f) * 2.0f;
float8 data_g;
float2 input_gauss_pos;
float step_x;
input_gauss_pos.x = g_x / out_width + in_sampler_offset_x;
input_gauss_pos.y = g_y / out_height + in_sampler_offset_y;
#if !PYRAMID_UV
step_x = 0.125f / out_width;
data_g = read_scale_y (input_gauss, gauss_sampler, input_gauss_pos, step_x) * 256.0f;
#else
step_x = 0.25f / out_width;
data_g = read_scale_uv (input_gauss, gauss_sampler, input_gauss_pos, step_x) * 256.0f;
#endif
#if CL_PYRAMID_ENABLE_DUMP
write_imageui (dump_resize, (int2)(g_x, g_y), convert_uint4(as_ushort4(convert_uchar8(data_g))));
#endif
data_g += lap + 0.5f;
data_g = clamp (data_g, 0.0f, 255.0f);
write_imageui (output, (int2)(g_x + out_offset_x, g_y), convert_uint4(as_ushort4(convert_uchar8(data_g))));
#if CL_PYRAMID_ENABLE_DUMP
write_imageui (dump_final, (int2)(g_x, g_y), convert_uint4(as_ushort4(convert_uchar8(data_g))));
#endif
}
__kernel void
kernel_pyramid_blend (
__read_only image2d_t input0, __read_only image2d_t input1,
#if !PYRAMID_UV
__global const float8 *input0_mask,
#else
__global const float4 *input0_mask,
#endif
__write_only image2d_t output)
{
sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
const int g_x = get_global_id (0);
const int g_y = get_global_id (1);
int2 pos = (int2) (g_x, g_y);
float8 data0 = convert_float8(as_uchar8(convert_ushort4(read_imageui(input0, sampler, pos))));
float8 data1 = convert_float8(as_uchar8(convert_ushort4(read_imageui(input1, sampler, pos))));
float8 out_data;
#if !PYRAMID_UV
out_data = (data0 - data1) * input0_mask[g_x] + data1;
#else
float8 coeff;
coeff.even = input0_mask[g_x];
coeff.odd = coeff.even;
out_data = (data0 - data1) * coeff + data1;
#endif
out_data = clamp (out_data + 0.5f, 0.0f, 255.0f);
write_imageui(output, pos, convert_uint4(as_ushort4(convert_uchar8(out_data))));
}
__kernel void
kernel_pyramid_scale (
__read_only image2d_t input, __write_only image2d_t output,
int out_offset_x, int output_width, int output_height)
{
const sampler_t sampler = CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR;
int g_x = get_global_id (0);
int g_y = get_global_id (1);
float2 normCoor = (float2)(g_x, g_y) / (float2)(output_width, output_height);
float8 out_data;
float step_x;
#if !PYRAMID_UV
step_x = 0.125f / output_width;
out_data = read_scale_y (input, sampler, normCoor, step_x) * 255.0f;
#else
step_x = 0.25f / output_width;
out_data = read_scale_uv (input, sampler, normCoor, step_x) * 255.0f;
#endif
out_data = clamp (out_data + 0.5f, 0.0f, 255.0f);
write_imageui (output, (int2)(g_x + out_offset_x, g_y), convert_uint4(as_ushort4(convert_uchar8(out_data))));
}
__kernel void
kernel_pyramid_copy (
__read_only image2d_t input, int in_offset_x,
__write_only image2d_t output, int out_offset_x,
int max_g_x, int max_g_y)
{
sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
const int g_x = get_global_id (0);
const int g_y = get_global_id (1);
if (g_x >= max_g_x || g_y >= max_g_y)
return;
uint4 data = read_imageui (input, sampler, (int2)(g_x + in_offset_x, g_y));
write_imageui (output, (int2)(g_x + out_offset_x, g_y), data);
}
/*
* input_gauss: RGBA-CL_UNSIGNED_INT18
* input_lap: RGBA-CL_UNSIGNED_INT16
* output: RGBA-CL_UNSIGNED_INT16
* each work-item calc 2 lines
*/
__kernel void
kernel_lap_transform (
__read_only image2d_t input_gauss0, int gauss0_offset_x,
__read_only image2d_t input_gauss1,
float gauss1_sampler_offset_x, float gauss1_sampler_offset_y,
__write_only image2d_t output, int lap_offset_x, float out_width, float out_height)
{
int g_x = get_global_id (0);
int g_y = get_global_id (1);
const sampler_t gauss0_sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
const sampler_t gauss1_sampler = CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR;
float8 orig = convert_float8(as_uchar8(convert_ushort4(
read_imageui(input_gauss0, gauss0_sampler, (int2)(g_x + gauss0_offset_x, g_y)))));
float8 zoom_in;
float2 gauss1_pos;
float sampler_step;
gauss1_pos.y = (g_y / out_height) + gauss1_sampler_offset_y;
gauss1_pos.x = (g_x / out_width) + gauss1_sampler_offset_x;
#if !PYRAMID_UV
sampler_step = 0.125f / out_width;
zoom_in.s0 = read_imagef (input_gauss1, gauss1_sampler, gauss1_pos).x;
gauss1_pos.x += sampler_step;
zoom_in.s1 = read_imagef (input_gauss1, gauss1_sampler, gauss1_pos).x;
gauss1_pos.x += sampler_step;
zoom_in.s2 = read_imagef (input_gauss1, gauss1_sampler, gauss1_pos).x;
gauss1_pos.x += sampler_step;
zoom_in.s3 = read_imagef (input_gauss1, gauss1_sampler, gauss1_pos).x;
gauss1_pos.x += sampler_step;
zoom_in.s4 = read_imagef (input_gauss1, gauss1_sampler, gauss1_pos).x;
gauss1_pos.x += sampler_step;
zoom_in.s5 = read_imagef (input_gauss1, gauss1_sampler, gauss1_pos).x;
gauss1_pos.x += sampler_step;
zoom_in.s6 = read_imagef (input_gauss1, gauss1_sampler, gauss1_pos).x;
gauss1_pos.x += sampler_step;
zoom_in.s7 = read_imagef (input_gauss1, gauss1_sampler, gauss1_pos).x;
#else
sampler_step = 0.25f / out_width;
zoom_in.s01 = read_imagef (input_gauss1, gauss1_sampler, gauss1_pos).xy;
gauss1_pos.x += sampler_step;
zoom_in.s23 = read_imagef (input_gauss1, gauss1_sampler, gauss1_pos).xy;
gauss1_pos.x += sampler_step;
zoom_in.s45 = read_imagef (input_gauss1, gauss1_sampler, gauss1_pos).xy;
gauss1_pos.x += sampler_step;
zoom_in.s67 = read_imagef (input_gauss1, gauss1_sampler, gauss1_pos).xy;
#endif
float8 lap = (orig - zoom_in * 256.0f) * 0.5f + 128.0f + 0.5f;
lap = clamp (lap, 0.0f, 255.0f);
write_imageui (output, (int2)(g_x + lap_offset_x, g_y), convert_uint4(as_ushort4(convert_uchar8(lap))));
}
/*
* input0: RGBA-CL_UNSIGNED_INT16
* input1: RGBA-CL_UNSIGNED_INT16
* out_diff: RGBA-CL_UNSIGNED_INT16
*/
__kernel void
kernel_image_diff (
__read_only image2d_t input0, int offset0,
__read_only image2d_t input1, int offset1,
__write_only image2d_t out_diff)
{
int g_x = get_global_id (0);
int g_y = get_global_id (1);
const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
int8 data0 = convert_int8(as_uchar8(convert_ushort4(read_imageui(input0, sampler, (int2)(g_x + offset0, g_y)))));
int8 data1 = convert_int8(as_uchar8(convert_ushort4(read_imageui(input1, sampler, (int2)(g_x + offset1, g_y)))));
uint8 diff = abs_diff (data0, data1);
write_imageui (out_diff, (int2)(g_x, g_y), convert_uint4(as_ushort4(convert_uchar8(diff))));
}
/*
* input0: RGBA-CL_UNSIGNED_INT16
*/
#define LEFT_POS (int)(-1)
#define MID_POS (int)(0)
#define RIGHT_POS (int)(1)
__inline int pos_buf_index (int x, int y, int stride)
{
return mad24 (stride, y, x);
}
__kernel void
kernel_seam_dp (
__read_only image2d_t image,
__global short *pos_buf, __global float *sum_buf, int offset_x, int valid_width,
int max_pos, int seam_height, int seam_stride)
{
int l_x = get_local_id (0);
int group_id = get_group_id (0);
if (l_x >= valid_width)
return;
// group0 fill first half slice image curve y = [0, seam_height/2 - 1]
// group1 fill send half slice image curve = [seam_height - 1, seam_height/2]
int first_slice_h = seam_height / 2;
int group_h = (group_id == 0 ? first_slice_h : seam_height - first_slice_h);
__local float slm_sum[4096];
float mid, left, right, cur;
int slm_idx;
int default_pos;
int x = l_x + offset_x;
const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
int y = (group_id == 0 ? 0 : seam_height - 1);
float sum = convert_float(read_imageui(image, sampler, (int2)(x, y)).x);
default_pos = x;
slm_sum[l_x] = sum;
barrier (CLK_LOCAL_MEM_FENCE);
pos_buf[pos_buf_index(x, y, seam_stride)] = convert_short(default_pos);
for (int i = 0; i < group_h; ++i) {
y = (group_id == 0 ? i : seam_height - i - 1);
slm_idx = l_x - 1;
slm_idx = (slm_idx > 0 ? slm_idx : 0);
left = slm_sum[slm_idx];
slm_idx = l_x + 1;
slm_idx = (slm_idx < valid_width ? slm_idx : valid_width - 1);
right = slm_sum[slm_idx];
cur = convert_float(read_imageui(image, sampler, (int2)(x, y)).x);
left = left + cur;
right = right + cur;
mid = sum + cur;
int pos;
pos = (left < mid) ? LEFT_POS : MID_POS;
sum = min (left, mid);
pos = (sum < right) ? pos : RIGHT_POS;
sum = min (sum, right);
slm_sum[l_x] = sum;
barrier (CLK_LOCAL_MEM_FENCE);
pos += default_pos;
pos = clamp (pos, offset_x, max_pos);
//if (l_x == 3)
// printf ("s:%f, pos:%d, mid:%f, offset_x:%d\n", sum.s0, pos.s0, mid.s0, offset_x);
pos_buf[pos_buf_index(x, y, seam_stride)] = convert_short(pos);
}
sum_buf[group_id * seam_stride + x] = sum;
//printf ("sum(x):%f(x:%d)\n", sum_buf[x].s0, x);
}
__kernel void
kernel_seam_mask_blend (
__read_only image2d_t input0, __read_only image2d_t input1,
__read_only image2d_t seam_mask,
__write_only image2d_t output)
{
sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
const int g_x = get_global_id (0);
const int g_y = get_global_id (1);
int2 pos = (int2) (g_x, g_y);
float8 data0 = convert_float8(as_uchar8(convert_ushort4(read_imageui(input0, sampler, pos))));
float8 data1 = convert_float8(as_uchar8(convert_ushort4(read_imageui(input1, sampler, pos))));
float8 coeff0 = convert_float8(as_uchar8(convert_ushort4(read_imageui(seam_mask, sampler, pos)))) / 255.0f;
float8 out_data;
#if !PYRAMID_UV
out_data = (data0 - data1) * coeff0 + data1;
#else
coeff0.even = (coeff0.even + coeff0.odd) * 0.5f;
coeff0.odd = coeff0.even;
out_data = (data0 - data1) * coeff0 + data1;
#endif
out_data = clamp (out_data + 0.5f, 0.0f, 255.0f);
write_imageui(output, pos, convert_uint4(as_ushort4(convert_uchar8(out_data))));
}
#define MASK_GAUSS_R 4
#define MASK_COEFF_MID 7
__constant const float mask_coeffs[] = {0.0f, 0.0f, 0.0f, 0.082f, 0.102f, 0.119f, 0.130f, 0.134f, 0.130f, 0.119f, 0.102f, 0.082f, 0.0f, 0.0f, 0.0f};
/*
* input: RGBA-CL_UNSIGNED_INT16
* output_gauss: RGBA-CL_UNSIGNED_INT8 ?
* output_lap:RGBA-CL_UNSIGNED_INT16
* each work-item calc 2 lines
*/
__kernel void
kernel_mask_gauss_scale_slm (
__read_only image2d_t input,
__write_only image2d_t output_gauss,
int image_width
#if ENABLE_MASK_GAUSS_SCALE
, __write_only image2d_t output_scale
#endif
)
{
#define WI_LINES 2
// input image width MUST < MASK_GAUSS_SLM_WIDTH*4
#define MASK_GAUSS_SLM_WIDTH 256
#define CONV_COEFF 128.0f
int g_x = get_global_id (0);
int g_y = get_global_id (1) * WI_LINES;
const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
__local ushort4 slm_gauss_y[WI_LINES][MASK_GAUSS_SLM_WIDTH];
float8 result_cur[WI_LINES] = {zero8, zero8};
float8 tmp_data;
int i_line;
int cur_g_y;
#pragma unroll
for (i_line = -MASK_GAUSS_R; i_line <= MASK_GAUSS_R + 1; i_line++) {
cur_g_y = g_y + i_line;
tmp_data = convert_float8(as_uchar8(convert_ushort4(read_imageui(input, sampler, (int2)(g_x, cur_g_y)))));
result_cur[0] += tmp_data * mask_coeffs[i_line + MASK_COEFF_MID];
result_cur[1] += tmp_data * mask_coeffs[i_line + MASK_COEFF_MID - 1];
}
((__local ushort8*)(slm_gauss_y[0]))[g_x] = convert_ushort8(result_cur[0] * CONV_COEFF);
((__local ushort8*)(slm_gauss_y[1]))[g_x] = convert_ushort8(result_cur[1] * CONV_COEFF);
barrier (CLK_LOCAL_MEM_FENCE);
float8 final_g[WI_LINES];
float4 result_pre;
float4 result_next;
#pragma unroll
for (i_line = 0; i_line < WI_LINES; ++i_line) {
result_pre = convert_float4(slm_gauss_y[i_line][clamp (g_x * 2 - 1, 0, image_width * 2)]) / CONV_COEFF;
result_next = convert_float4(slm_gauss_y[i_line][clamp (g_x * 2 + 2, 0, image_width * 2)]) / CONV_COEFF;
final_g[i_line] = result_cur[i_line] * mask_coeffs[MASK_COEFF_MID] +
(float8)(result_pre.s3, result_cur[i_line].s0123, result_cur[i_line].s456) *
mask_coeffs[MASK_COEFF_MID + 1] +
(float8)(result_cur[i_line].s1234, result_cur[i_line].s567, result_next.s0) *
mask_coeffs[MASK_COEFF_MID + 1] +
(float8)(result_pre.s23, result_cur[i_line].s0123, result_cur[i_line].s45) *
mask_coeffs[MASK_COEFF_MID + 2] +
(float8)(result_cur[i_line].s2345, result_cur[i_line].s67, result_next.s01) *
mask_coeffs[MASK_COEFF_MID + 2] +
(float8)(result_pre.s123, result_cur[i_line].s0123, result_cur[i_line].s4) *
mask_coeffs[MASK_COEFF_MID + 3] +
(float8)(result_cur[i_line].s3456, result_cur[i_line].s7, result_next.s012) *
mask_coeffs[MASK_COEFF_MID + 3] +
(float8)(result_pre.s0123, result_cur[i_line].s0123) * mask_coeffs[MASK_COEFF_MID + 4] +
(float8)(result_cur[i_line].s4567, result_next.s0123) * mask_coeffs[MASK_COEFF_MID + 4];
final_g[i_line] = clamp (final_g[i_line] + 0.5f, 0.0f, 255.0f);
//if ((g_x == 9 || g_x == 8) && g_y == 0) {
// printf ("(x:%d, y:0), pre:" ARG_FORMAT4 "cur" ARG_FORMAT8 "next" ARG_FORMAT4 "final:" ARG_FORMAT8 "\n",
// g_x, ARGS4(result_pre), ARGS8(result_cur[i_line]), ARGS4(result_next), ARGS8(final_g[i_line]));
//}
write_imageui (output_gauss, (int2)(g_x, g_y + i_line), convert_uint4(as_ushort4(convert_uchar8(final_g[i_line]))));
}
#if ENABLE_MASK_GAUSS_SCALE
write_imageui (output_scale, (int2)(g_x, get_global_id (1)), convert_uint4(final_g[0].even));
#endif
}
__kernel void
kernel_mask_gauss_scale (
__read_only image2d_t input,
__write_only image2d_t output_gauss
#if ENABLE_MASK_GAUSS_SCALE
, __write_only image2d_t output_scale
#endif
)
{
int g_x = get_global_id (0);
int in_x = g_x;
int g_y = get_global_id (1) * 2;
const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
float8 result_pre[2] = {zero8, zero8};
float8 result_next[2] = {zero8, zero8};
float8 result_cur[2] = {zero8, zero8};
float8 final_g[2];
float8 tmp_data;
int i_line;
int cur_g_y;
float coeff0, coeff1;
#pragma unroll
for (i_line = -MASK_GAUSS_R; i_line <= MASK_GAUSS_R + 1; i_line++) {
cur_g_y = g_y + i_line;
coeff0 = mask_coeffs[i_line + MASK_COEFF_MID];
coeff1 = mask_coeffs[i_line + MASK_COEFF_MID - 1];
tmp_data = convert_float8(as_uchar8(convert_ushort4(read_imageui(input, sampler, (int2)(in_x - 1, cur_g_y)))));
result_pre[0] += tmp_data * coeff0;
result_pre[1] += tmp_data * coeff1;
tmp_data = convert_float8(as_uchar8(convert_ushort4(read_imageui(input, sampler, (int2)(in_x, cur_g_y)))));
result_cur[0] += tmp_data * coeff0;
result_cur[1] += tmp_data * coeff1;
tmp_data = convert_float8(as_uchar8(convert_ushort4(read_imageui(input, sampler, (int2)(in_x + 1, cur_g_y)))));
result_next[1] += tmp_data * coeff1;
result_next[0] += tmp_data * coeff0;
}
#pragma unroll
for (i_line = 0; i_line < 2; ++i_line) {
final_g[i_line] = result_cur[i_line] * mask_coeffs[MASK_COEFF_MID] +
(float8)(result_pre[i_line].s7, result_cur[i_line].s0123, result_cur[i_line].s456) *
mask_coeffs[MASK_COEFF_MID + 1] +
(float8)(result_cur[i_line].s1234, result_cur[i_line].s567, result_next[i_line].s0) *
mask_coeffs[MASK_COEFF_MID + 1] +
(float8)(result_pre[i_line].s67, result_cur[i_line].s0123, result_cur[i_line].s45) *
mask_coeffs[MASK_COEFF_MID + 2] +
(float8)(result_cur[i_line].s2345, result_cur[i_line].s67, result_next[i_line].s01) *
mask_coeffs[MASK_COEFF_MID + 2] +
(float8)(result_pre[i_line].s567, result_cur[i_line].s0123, result_cur[i_line].s4) *
mask_coeffs[MASK_COEFF_MID + 3] +
(float8)(result_cur[i_line].s3456,result_cur[i_line].s7, result_next[i_line].s012) *
mask_coeffs[MASK_COEFF_MID + 3] +
(float8)(result_pre[i_line].s4567, result_cur[i_line].s0123) * mask_coeffs[MASK_COEFF_MID + 4] +
(float8)(result_cur[i_line].s4567, result_next[i_line].s0123) * mask_coeffs[MASK_COEFF_MID + 4];
final_g[i_line] = clamp (final_g[i_line] + 0.5f, 0.0f, 255.0f);
write_imageui (output_gauss, (int2)(g_x, g_y + i_line), convert_uint4(as_ushort4(convert_uchar8(final_g[i_line]))));
}
#if ENABLE_MASK_GAUSS_SCALE
write_imageui (output_scale, (int2)(g_x, get_global_id (1)), convert_uint4(final_g[0].even));
#endif
}