cl_kernel/kernel_image_warp.cl

/**
* \brief Image warping kernel function.
* \param[in] input Input image object.
* \param[out] output scaled output image object.
* \param[in] warp_config: image warping parameters
*/

#ifndef WARP_Y
#define WARP_Y 1
#endif

// 8 bytes for each Y pixel
#define PIXEL_X_STEP   8

typedef struct {
    int frame_id;
    int width;
    int height;
    float trim_ratio;
    float proj_mat[9];
} CLWarpConfig;

__kernel void
kernel_image_warp_8_pixel (
    __read_only image2d_t input,
    __write_only image2d_t output,
    CLWarpConfig warp_config)
{
    // dest coordinate
    int d_x = get_global_id(0);
    int d_y = get_global_id(1);

    int out_width = get_image_width (output);
    int out_height = get_image_height (output);

    const sampler_t sampler = CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR;

    // source coordinate
    float s_x = 0.0f;
    float s_y = 0.0f;
    float warp_x = 0.0f;
    float warp_y = 0.0f;
    float w = 0.0f;

    float t_x = 0.0f;
    float t_y = 0.0f;

    float16 pixel = 0.0f;
    float* output_pixel = (float*)(&pixel);
    int i = 0;

    t_y = d_y;
#pragma unroll
    for (i = 0; i < PIXEL_X_STEP; i++) {
        t_x = (float)(PIXEL_X_STEP * d_x + i);

        s_x = warp_config.proj_mat[0] * t_x +
              warp_config.proj_mat[1] * t_y +
              warp_config.proj_mat[2];
        s_y = warp_config.proj_mat[3] * t_x +
              warp_config.proj_mat[4] * t_y +
              warp_config.proj_mat[5];
        w = warp_config.proj_mat[6] * t_x +
            warp_config.proj_mat[7] * t_y +
            warp_config.proj_mat[8];
        w = w != 0.0f ? 1.0f / w : 0.0f;

        warp_x = (s_x * w) / (float)(PIXEL_X_STEP * out_width);
        warp_y = (s_y * w) / (float)out_height;

#if WARP_Y
        output_pixel[i] = read_imagef(input, sampler, (float2)(warp_x, warp_y)).x;
#else
        float2 temp = read_imagef(input, sampler, (float2)(warp_x, warp_y)).xy;
        output_pixel[2 * i] = temp.x;
        output_pixel[2 * i + 1] = temp.y;
#endif
    }

#if WARP_Y
    write_imageui(output, (int2)(d_x, d_y), convert_uint4(as_ushort4(convert_uchar8(pixel.lo * 255.0f))));
#else
    write_imageui(output, (int2)(d_x, d_y), as_uint4(convert_uchar16(pixel * 255.0f)));
#endif

}

__kernel void
kernel_image_warp_1_pixel (
    __read_only image2d_t input,
    __write_only image2d_t output,
    CLWarpConfig warp_config)
{
    // dest coordinate
    int d_x = get_global_id(0);
    int d_y = get_global_id(1);

    int out_width = get_image_width (output);
    int out_height = get_image_height (output);

    const sampler_t sampler = CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR;

    // source coordinate
    float s_x = warp_config.proj_mat[0] * d_x +
                warp_config.proj_mat[1] * d_y +
                warp_config.proj_mat[2];
    float s_y = warp_config.proj_mat[3] * d_x +
                warp_config.proj_mat[4] * d_y +
                warp_config.proj_mat[5];
    float w = warp_config.proj_mat[6] * d_x +
              warp_config.proj_mat[7] * d_y +
              warp_config.proj_mat[8];
    w = w != 0.0f ? 1.0f / w : 0.0f;

    float warp_x = (s_x * w) / (float)out_width;
    float warp_y = (s_y * w) / (float)out_height;

    float4 pixel = read_imagef(input, sampler, (float2)(warp_x, warp_y));

    write_imagef(output, (int2)(d_x, d_y), pixel);
}