1/**
2* \brief Image warping kernel function.
3* \param[in] input Input image object.
4* \param[out] output scaled output image object.
5* \param[in] warp_config: image warping parameters
6*/
7
8#ifndef WARP_Y
9#define WARP_Y 1
10#endif
11
12// 8 bytes for each Y pixel
13#define PIXEL_X_STEP   8
14
15typedef struct {
16    int frame_id;
17    int width;
18    int height;
19    float trim_ratio;
20    float proj_mat[9];
21} CLWarpConfig;
22
23__kernel void
24kernel_image_warp_8_pixel (
25    __read_only image2d_t input,
26    __write_only image2d_t output,
27    CLWarpConfig warp_config)
28{
29    // dest coordinate
30    int d_x = get_global_id(0);
31    int d_y = get_global_id(1);
32
33    int out_width = get_image_width (output);
34    int out_height = get_image_height (output);
35
36    const sampler_t sampler = CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR;
37
38    // source coordinate
39    float s_x = 0.0f;
40    float s_y = 0.0f;
41    float warp_x = 0.0f;
42    float warp_y = 0.0f;
43    float w = 0.0f;
44
45    float t_x = 0.0f;
46    float t_y = 0.0f;
47
48    float16 pixel = 0.0f;
49    float* output_pixel = (float*)(&pixel);
50    int i = 0;
51
52    t_y = d_y;
53#pragma unroll
54    for (i = 0; i < PIXEL_X_STEP; i++) {
55        t_x = (float)(PIXEL_X_STEP * d_x + i);
56
57        s_x = warp_config.proj_mat[0] * t_x +
58              warp_config.proj_mat[1] * t_y +
59              warp_config.proj_mat[2];
60        s_y = warp_config.proj_mat[3] * t_x +
61              warp_config.proj_mat[4] * t_y +
62              warp_config.proj_mat[5];
63        w = warp_config.proj_mat[6] * t_x +
64            warp_config.proj_mat[7] * t_y +
65            warp_config.proj_mat[8];
66        w = w != 0.0f ? 1.0f / w : 0.0f;
67
68        warp_x = (s_x * w) / (float)(PIXEL_X_STEP * out_width);
69        warp_y = (s_y * w) / (float)out_height;
70
71#if WARP_Y
72        output_pixel[i] = read_imagef(input, sampler, (float2)(warp_x, warp_y)).x;
73#else
74        float2 temp = read_imagef(input, sampler, (float2)(warp_x, warp_y)).xy;
75        output_pixel[2 * i] = temp.x;
76        output_pixel[2 * i + 1] = temp.y;
77#endif
78    }
79
80#if WARP_Y
81    write_imageui(output, (int2)(d_x, d_y), convert_uint4(as_ushort4(convert_uchar8(pixel.lo * 255.0f))));
82#else
83    write_imageui(output, (int2)(d_x, d_y), as_uint4(convert_uchar16(pixel * 255.0f)));
84#endif
85
86}
87
88__kernel void
89kernel_image_warp_1_pixel (
90    __read_only image2d_t input,
91    __write_only image2d_t output,
92    CLWarpConfig warp_config)
93{
94    // dest coordinate
95    int d_x = get_global_id(0);
96    int d_y = get_global_id(1);
97
98    int out_width = get_image_width (output);
99    int out_height = get_image_height (output);
100
101    const sampler_t sampler = CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR;
102
103    // source coordinate
104    float s_x = warp_config.proj_mat[0] * d_x +
105                warp_config.proj_mat[1] * d_y +
106                warp_config.proj_mat[2];
107    float s_y = warp_config.proj_mat[3] * d_x +
108                warp_config.proj_mat[4] * d_y +
109                warp_config.proj_mat[5];
110    float w = warp_config.proj_mat[6] * d_x +
111              warp_config.proj_mat[7] * d_y +
112              warp_config.proj_mat[8];
113    w = w != 0.0f ? 1.0f / w : 0.0f;
114
115    float warp_x = (s_x * w) / (float)out_width;
116    float warp_y = (s_y * w) / (float)out_height;
117
118    float4 pixel = read_imagef(input, sampler, (float2)(warp_x, warp_y));
119
120    write_imagef(output, (int2)(d_x, d_y), pixel);
121}
122
123