/*
|
* kernel_geo_map
|
* input_y, input image, CL_R + CL_UNORM_INT8
|
* input_uv, CL_RG + CL_UNORM_INT8
|
* geo_table, CL_RGBA + CL_FLOAT
|
* output_y, CL_RGBA + CL_UNSIGNED_INT16
|
* output_uv, CL_RGBA + CL_UNSIGNED_INT16
|
*
|
* description:
|
* the center of geo_table and output positons are both mapped to (0, 0)
|
*/
|
|
#ifndef ENABLE_LSC
|
#define ENABLE_LSC 0
|
#endif
|
|
#define CONST_DATA_Y 0.0f
|
#define CONST_DATA_UV (float2)(0.5f, 0.5f)
|
|
// 8 bytes for each pixel
|
#define PIXEL_RES_STEP_X 8
|
|
void get_geo_mapped_y (
|
__read_only image2d_t input,
|
__read_only image2d_t geo_table, float2 table_pos, float step_x,
|
bool *out_of_bound, float2 *input_pos, float8 *out_y)
|
{
|
sampler_t sampler = CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR;
|
float *output_data = (float*)(out_y);
|
int i = 0;
|
|
for (i = 0; i < PIXEL_RES_STEP_X; ++i) {
|
out_of_bound[i] =
|
(min (table_pos.x, table_pos.y) < 0.0f) ||
|
(max (table_pos.x, table_pos.y) > 1.0f);
|
input_pos[i] = read_imagef (geo_table, sampler, table_pos).xy;
|
out_of_bound[i] =
|
out_of_bound[i] ||
|
(min (input_pos[i].x, input_pos[i].y) < 0.0f) ||
|
(max (input_pos[i].x, input_pos[i].y) > 1.0f);
|
//need convert input_pos to (0.0 ~ 1.0)????
|
output_data[i] = out_of_bound[i] ? CONST_DATA_Y : read_imagef (input, sampler, input_pos[i]).x;
|
table_pos.x += step_x;
|
}
|
}
|
|
void get_lsc_data (
|
image2d_t lsc_table, int2 g_pos, float step_x,
|
float2 gray_threshold, float8 output, float8 *lsc_data)
|
{
|
sampler_t sampler = CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR;
|
float *lsc_ptr = (float *)(lsc_data);
|
|
float2 pos = convert_float2((int2)(g_pos.x * PIXEL_RES_STEP_X, g_pos.y)) * step_x;
|
for (int i = 0; i < PIXEL_RES_STEP_X; ++i) {
|
lsc_ptr[i] = read_imagef (lsc_table, sampler, pos).x;
|
pos.x += step_x;
|
}
|
|
float8 diff_ratio = (gray_threshold.y - output * 255.0f) / (gray_threshold.y - gray_threshold.x);
|
diff_ratio = clamp (diff_ratio, 0.0f, 1.0f);
|
(*lsc_data) = diff_ratio * diff_ratio * ((*lsc_data) - 1.0f) + 1.0f;
|
}
|
|
__kernel void
|
kernel_geo_map (
|
__read_only image2d_t input_y, __read_only image2d_t input_uv,
|
__read_only image2d_t geo_table, float2 table_scale_size,
|
#if ENABLE_LSC
|
__read_only image2d_t lsc_table, float2 gray_threshold,
|
#endif
|
__write_only image2d_t output_y, __write_only image2d_t output_uv, float2 out_size)
|
{
|
const int g_x = get_global_id (0);
|
const int g_y_uv = get_global_id (1);
|
const int g_y = get_global_id (1) * 2;
|
float8 output_data;
|
float2 from_pos;
|
bool out_of_bound[8];
|
float2 input_pos[8];
|
// map to [-0.5, 0.5)
|
float2 table_scale_step = 1.0f / table_scale_size;
|
float2 out_map_pos;
|
sampler_t sampler = CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR;
|
|
out_map_pos = (convert_float2((int2)(g_x * PIXEL_RES_STEP_X, g_y)) - out_size / 2.0f) * table_scale_step + 0.5f;
|
|
get_geo_mapped_y (input_y, geo_table, out_map_pos, table_scale_step.x, out_of_bound, input_pos, &output_data);
|
|
#if ENABLE_LSC
|
float8 lsc_data;
|
get_lsc_data (lsc_table, (int2)(g_x, g_y), table_scale_step.x, gray_threshold, output_data, &lsc_data);
|
output_data = clamp (output_data * lsc_data, 0.0f, 1.0f);
|
#endif
|
write_imageui (output_y, (int2)(g_x, g_y), convert_uint4(as_ushort4(convert_uchar8(output_data * 255.0f))));
|
|
output_data.s01 = out_of_bound[0] ? CONST_DATA_UV : read_imagef (input_uv, sampler, input_pos[0]).xy;
|
output_data.s23 = out_of_bound[2] ? CONST_DATA_UV : read_imagef (input_uv, sampler, input_pos[2]).xy;
|
output_data.s45 = out_of_bound[4] ? CONST_DATA_UV : read_imagef (input_uv, sampler, input_pos[4]).xy;
|
output_data.s67 = out_of_bound[6] ? CONST_DATA_UV : read_imagef (input_uv, sampler, input_pos[6]).xy;
|
write_imageui (output_uv, (int2)(g_x, g_y_uv), convert_uint4(as_ushort4(convert_uchar8(output_data * 255.0f))));
|
|
out_map_pos.y += table_scale_step.y;
|
get_geo_mapped_y (input_y, geo_table, out_map_pos, table_scale_step.x, out_of_bound, input_pos, &output_data);
|
|
#if ENABLE_LSC
|
get_lsc_data (lsc_table, (int2)(g_x, g_y + 1), table_scale_step.x, gray_threshold, output_data, &lsc_data);
|
output_data = clamp (output_data * lsc_data, 0.0f, 1.0f);
|
#endif
|
write_imageui (output_y, (int2)(g_x, g_y + 1), convert_uint4(as_ushort4(convert_uchar8(output_data * 255.0f))));
|
}
|
