cl_kernel/kernel_fisheye.cl

/*
 * kernel_fisheye_2_gps
 * input_y,      input image, CL_R + CL_UNORM_INT8 //sampler
 * input_uv, CL_RG + CL_UNORM_INT8  //sampler
 * output_y,  CL_RGBA + CL_UNSIGNED_INT8, // 4-pixel
 * output_uv,  CL_RGBA + CL_UNSIGNED_INT8, // 4-pixel
 *
 * all angles are in radian
 */

#define PI 3.1415926f
#define PIXEL_PER_WI 4

typedef struct {
    float    center_x;
    float    center_y;
    float    wide_angle;
    float    radius;
    float    rotate_angle;
} FisheyeInfo;

__inline float2 calculate_fisheye_pos (float2 gps_pos, const FisheyeInfo *info)
{
    float z = cos (gps_pos.y);
    float x = sin (gps_pos.y) * cos (gps_pos.x);
    float y = sin (gps_pos.y) * sin (gps_pos.x);
    float r_angle = acos (y);
    float r = r_angle * (info->radius * 2.0f) / info->wide_angle;
    float xz_size = sqrt(x * x + z * z);

    float2 dst;
    dst.x = -r * x / xz_size;
    dst.y = -r * z / xz_size;

    float2 ret;
    ret.x = cos(info->rotate_angle) * dst.x - sin(info->rotate_angle) * dst.y;
    ret.y = sin(info->rotate_angle) * dst.x + cos (info->rotate_angle) * dst.y;

    return ret + (float2)(info->center_x, info->center_y);
}

__kernel void
kernel_fisheye_table (
    const FisheyeInfo info, const float2 fisheye_image_size,
    __write_only image2d_t table, const float2 radian_per_pixel, const float2 table_center)
{
    int2 out_pos = (int2)(get_global_id (0), get_global_id (1));
    float2 gps_pos = (convert_float2 (out_pos) - table_center) * radian_per_pixel + PI / 2.0f;
    float2 pos = calculate_fisheye_pos (gps_pos, &info);
    float2 min_pos = (float2)(info.center_x - info.radius, info.center_y - info.radius);
    float2 max_pos = (float2)(info.center_x + info.radius, info.center_y + info.radius);
    pos = clamp (pos, min_pos, max_pos);
    pos /= fisheye_image_size;
    write_imagef (table, out_pos, (float4)(pos, 0.0f, 0.0f));
}

__kernel void
kernel_lsc_table (
    __read_only image2d_t geo_table, __write_only image2d_t lsc_table,
    __global float *lsc_array, int array_size, const FisheyeInfo info, const float2 fisheye_image_size)
{
    sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
    int2 pos = (int2) (get_global_id (0), get_global_id (1));

    float2 geo_data = read_imagef (geo_table, sampler, pos).xy * fisheye_image_size;
    float2 dist = geo_data - (float2)(info.center_x, info.center_y);
    float r = sqrt (dist.x * dist.x + dist.y * dist.y);
    r /= (1.0f * info.radius / array_size);

    int min_idx = r;
    int max_idx = r + 1.0f;
    float lsc_data = max_idx > (array_size - 1) ? lsc_array[array_size - 1] :
                     (r - min_idx) * (lsc_array[max_idx] - lsc_array[min_idx]) + lsc_array[min_idx];

    write_imagef (lsc_table, pos, (float4)(lsc_data, 0.0f, 0.0f, 1.0f));
}

__kernel void
kernel_fisheye_2_gps (
    __read_only image2d_t input_y, __read_only image2d_t input_uv,
    const float2 input_y_size, const FisheyeInfo info,
    __write_only image2d_t output_y, __write_only image2d_t output_uv,
    const float2 dst_center, const float2 radian_per_pixel)
{
    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;
    float2 src_pos[4];
    float4 src_data;
    float *src_ptr = (float*)(&src_data);
    sampler_t sampler = CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR;

    float2 gps_start_pos =
        (convert_float2((int2)(g_x * PIXEL_PER_WI, g_y)) - dst_center) * radian_per_pixel + PI / 2.0f;
    float2 gps_pos = gps_start_pos;

#pragma unroll
    for (int i = 0; i < PIXEL_PER_WI; ++i) {
        float2 pos = calculate_fisheye_pos (gps_pos, &info);
        src_pos[i] = pos / input_y_size;
        src_ptr[i] = read_imagef (input_y, sampler, src_pos[i]).x;
        gps_pos.x += radian_per_pixel.x;
    }
    write_imageui (output_y, (int2)(g_x, g_y), convert_uint4(convert_uchar4(src_data * 255.0f)));

    src_data.s01 = read_imagef (input_uv, sampler, src_pos[0]).xy;
    src_data.s23 = read_imagef (input_uv, sampler, src_pos[2]).xy;
    write_imageui (output_uv, (int2)(g_x, g_y_uv), convert_uint4(convert_uchar4(src_data * 255.0f)));

    gps_pos = gps_start_pos;
    gps_pos.y += radian_per_pixel.y;
#pragma unroll
    for (int i = 0; i < PIXEL_PER_WI; ++i) {
        float2 pos = calculate_fisheye_pos (gps_pos, &info);
        pos /= input_y_size;
        src_ptr[i] = read_imagef (input_y, sampler, pos).x;
        gps_pos.x += radian_per_pixel.x;
    }
    write_imageui (output_y, (int2)(g_x, g_y + 1), convert_uint4(convert_uchar4(src_data * 255.0f)));

}