1 /* 2 * kernel_geo_map 3 * input_y, input image, CL_R + CL_UNORM_INT8 4 * input_uv, CL_RG + CL_UNORM_INT8 5 * geo_table, CL_RGBA + CL_FLOAT 6 * output_y, CL_RGBA + CL_UNSIGNED_INT16 7 * output_uv, CL_RGBA + CL_UNSIGNED_INT16 8 * 9 * description: 10 * the center of geo_table and output positons are both mapped to (0, 0) 11 */ 12 13 #ifndef ENABLE_LSC 14 #define ENABLE_LSC 0 15 #endif 16 17 #define CONST_DATA_Y 0.0f 18 #define CONST_DATA_UV (float2)(0.5f, 0.5f) 19 20 // 8 bytes for each pixel 21 #define PIXEL_RES_STEP_X 8 22 23 void get_geo_mapped_y ( 24 __read_only image2d_t input, 25 __read_only image2d_t geo_table, float2 table_pos, float step_x, 26 bool *out_of_bound, float2 *input_pos, float8 *out_y) 27 { 28 sampler_t sampler = CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR; 29 float *output_data = (float*)(out_y); 30 int i = 0; 31 32 for (i = 0; i < PIXEL_RES_STEP_X; ++i) { 33 out_of_bound[i] = 34 (min (table_pos.x, table_pos.y) < 0.0f) || 35 (max (table_pos.x, table_pos.y) > 1.0f); 36 input_pos[i] = read_imagef (geo_table, sampler, table_pos).xy; 37 out_of_bound[i] = 38 out_of_bound[i] || 39 (min (input_pos[i].x, input_pos[i].y) < 0.0f) || 40 (max (input_pos[i].x, input_pos[i].y) > 1.0f); 41 //need convert input_pos to (0.0 ~ 1.0)???? 42 output_data[i] = out_of_bound[i] ? CONST_DATA_Y : read_imagef (input, sampler, input_pos[i]).x; 43 table_pos.x += step_x; 44 } 45 } 46 47 void get_lsc_data ( 48 image2d_t lsc_table, int2 g_pos, float step_x, 49 float2 gray_threshold, float8 output, float8 *lsc_data) 50 { 51 sampler_t sampler = CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR; 52 float *lsc_ptr = (float *)(lsc_data); 53 54 float2 pos = convert_float2((int2)(g_pos.x * PIXEL_RES_STEP_X, g_pos.y)) * step_x; 55 for (int i = 0; i < PIXEL_RES_STEP_X; ++i) { 56 lsc_ptr[i] = read_imagef (lsc_table, sampler, pos).x; 57 pos.x += step_x; 58 } 59 60 float8 diff_ratio = (gray_threshold.y - output * 255.0f) / (gray_threshold.y - gray_threshold.x); 61 diff_ratio = clamp (diff_ratio, 0.0f, 1.0f); 62 (*lsc_data) = diff_ratio * diff_ratio * ((*lsc_data) - 1.0f) + 1.0f; 63 } 64 65 __kernel void 66 kernel_geo_map ( 67 __read_only image2d_t input_y, __read_only image2d_t input_uv, 68 __read_only image2d_t geo_table, float2 table_scale_size, 69 #if ENABLE_LSC 70 __read_only image2d_t lsc_table, float2 gray_threshold, 71 #endif 72 __write_only image2d_t output_y, __write_only image2d_t output_uv, float2 out_size) 73 { 74 const int g_x = get_global_id (0); 75 const int g_y_uv = get_global_id (1); 76 const int g_y = get_global_id (1) * 2; 77 float8 output_data; 78 float2 from_pos; 79 bool out_of_bound[8]; 80 float2 input_pos[8]; 81 // map to [-0.5, 0.5) 82 float2 table_scale_step = 1.0f / table_scale_size; 83 float2 out_map_pos; 84 sampler_t sampler = CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR; 85 86 out_map_pos = (convert_float2((int2)(g_x * PIXEL_RES_STEP_X, g_y)) - out_size / 2.0f) * table_scale_step + 0.5f; 87 88 get_geo_mapped_y (input_y, geo_table, out_map_pos, table_scale_step.x, out_of_bound, input_pos, &output_data); 89 90 #if ENABLE_LSC 91 float8 lsc_data; 92 get_lsc_data (lsc_table, (int2)(g_x, g_y), table_scale_step.x, gray_threshold, output_data, &lsc_data); 93 output_data = clamp (output_data * lsc_data, 0.0f, 1.0f); 94 #endif 95 write_imageui (output_y, (int2)(g_x, g_y), convert_uint4(as_ushort4(convert_uchar8(output_data * 255.0f)))); 96 97 output_data.s01 = out_of_bound[0] ? CONST_DATA_UV : read_imagef (input_uv, sampler, input_pos[0]).xy; 98 output_data.s23 = out_of_bound[2] ? CONST_DATA_UV : read_imagef (input_uv, sampler, input_pos[2]).xy; 99 output_data.s45 = out_of_bound[4] ? CONST_DATA_UV : read_imagef (input_uv, sampler, input_pos[4]).xy; 100 output_data.s67 = out_of_bound[6] ? CONST_DATA_UV : read_imagef (input_uv, sampler, input_pos[6]).xy; 101 write_imageui (output_uv, (int2)(g_x, g_y_uv), convert_uint4(as_ushort4(convert_uchar8(output_data * 255.0f)))); 102 103 out_map_pos.y += table_scale_step.y; 104 get_geo_mapped_y (input_y, geo_table, out_map_pos, table_scale_step.x, out_of_bound, input_pos, &output_data); 105 106 #if ENABLE_LSC 107 get_lsc_data (lsc_table, (int2)(g_x, g_y + 1), table_scale_step.x, gray_threshold, output_data, &lsc_data); 108 output_data = clamp (output_data * lsc_data, 0.0f, 1.0f); 109 #endif 110 write_imageui (output_y, (int2)(g_x, g_y + 1), convert_uint4(as_ushort4(convert_uchar8(output_data * 255.0f)))); 111 } 112
