1/*M///////////////////////////////////////////////////////////////////////////////////////
2//
3//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
4//
5//  By downloading, copying, installing or using the software you agree to this license.
6//  If you do not agree to this license, do not download, install,
7//  copy or use the software.
8//
9//
10//                           License Agreement
11//                For Open Source Computer Vision Library
12//
13// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved.
14// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
15// Third party copyrights are property of their respective owners.
16//
17// @Authors
18//    Peng Xiao, pengxiao@multicorewareinc.com
19//
20// Redistribution and use in source and binary forms, with or without modification,
21// are permitted provided that the following conditions are met:
22//
23//   * Redistribution's of source code must retain the above copyright notice,
24//     this list of conditions and the following disclaimer.
25//
26//   * Redistribution's in binary form must reproduce the above copyright notice,
27//     this list of conditions and the following disclaimer in the documentation
28//     and/or other materials provided with the distribution.
29//
30//   * The name of the copyright holders may not be used to endorse or promote products
31//     derived from this software without specific prior written permission.
32//
33// This software is provided by the copyright holders and contributors as is and
34// any express or implied warranties, including, but not limited to, the implied
35// warranties of merchantability and fitness for a particular purpose are disclaimed.
36// In no event shall the Intel Corporation or contributors be liable for any direct,
37// indirect, incidental, special, exemplary, or consequential damages
38// (including, but not limited to, procurement of substitute goods or services;
39// loss of use, data, or profits; or business interruption) however caused
40// and on any theory of liability, whether in contract, strict liability,
41// or tort (including negligence or otherwise) arising in any way out of
42// the use of this software, even if advised of the possibility of such damage.
43//
44//M*/
45
46__kernel void buildWarpPlaneMaps(__global uchar * xmapptr, int xmap_step, int xmap_offset,
47                                 __global uchar * ymapptr, int ymap_step, int ymap_offset, int rows, int cols,
48                                 __constant float * ck_rinv, __constant float * ct,
49                                 int tl_u, int tl_v, float scale, int rowsPerWI)
50{
51    int du = get_global_id(0);
52    int dv0 = get_global_id(1) * rowsPerWI;
53
54    if (du < cols)
55    {
56        int xmap_index = mad24(dv0, xmap_step, mad24(du, (int)sizeof(float), xmap_offset));
57        int ymap_index = mad24(dv0, ymap_step, mad24(du, (int)sizeof(float), ymap_offset));
58
59        float u = tl_u + du;
60        float x_ = fma(u, scale, -ct[0]);
61        float ct1 = 1 - ct[2];
62
63        for (int dv = dv0, dv1 = min(rows, dv0 + rowsPerWI); dv < dv1; ++dv, xmap_index += xmap_step,
64            ymap_index += ymap_step)
65        {
66            __global float * xmap = (__global float *)(xmapptr + xmap_index);
67            __global float * ymap = (__global float *)(ymapptr + ymap_index);
68
69            float v = tl_v + dv;
70            float y_ = fma(v, scale, -ct[1]);
71
72            float x = fma(ck_rinv[0], x_, fma(ck_rinv[1], y_, ck_rinv[2] * ct1));
73            float y = fma(ck_rinv[3], x_, fma(ck_rinv[4], y_, ck_rinv[5] * ct1));
74            float z = fma(ck_rinv[6], x_, fma(ck_rinv[7], y_, ck_rinv[8] * ct1));
75
76            if (z != 0)
77                x /= z, y /= z;
78            else
79                x = y = -1;
80
81            xmap[0] = x;
82            ymap[0] = y;
83        }
84    }
85}
86
87__kernel void buildWarpCylindricalMaps(__global uchar * xmapptr, int xmap_step, int xmap_offset,
88                                       __global uchar * ymapptr, int ymap_step, int ymap_offset, int rows, int cols,
89                                       __constant float * ck_rinv, int tl_u, int tl_v, float scale, int rowsPerWI)
90{
91    int du = get_global_id(0);
92    int dv0 = get_global_id(1) * rowsPerWI;
93
94    if (du < cols)
95    {
96        int xmap_index = mad24(dv0, xmap_step, mad24(du, (int)sizeof(float), xmap_offset));
97        int ymap_index = mad24(dv0, ymap_step, mad24(du, (int)sizeof(float), ymap_offset));
98
99        float u = (tl_u + du) * scale;
100        float x_, z_;
101        x_ = sincos(u, &z_);
102
103        for (int dv = dv0, dv1 = min(rows, dv0 + rowsPerWI); dv < dv1; ++dv, xmap_index += xmap_step,
104            ymap_index += ymap_step)
105        {
106            __global float * xmap = (__global float *)(xmapptr + xmap_index);
107            __global float * ymap = (__global float *)(ymapptr + ymap_index);
108
109            float y_ = (tl_v + dv) * scale;
110
111            float x, y, z;
112            x = fma(ck_rinv[0], x_, fma(ck_rinv[1], y_, ck_rinv[2] * z_));
113            y = fma(ck_rinv[3], x_, fma(ck_rinv[4], y_, ck_rinv[5] * z_));
114            z = fma(ck_rinv[6], x_, fma(ck_rinv[7], y_, ck_rinv[8] * z_));
115
116            if (z > 0)
117                x /= z, y /= z;
118            else
119                x = y = -1;
120
121            xmap[0] = x;
122            ymap[0] = y;
123        }
124    }
125}
126
127__kernel void buildWarpSphericalMaps(__global uchar * xmapptr, int xmap_step, int xmap_offset,
128                                     __global uchar * ymapptr, int ymap_step, int ymap_offset, int rows, int cols,
129                                     __constant float * ck_rinv, int tl_u, int tl_v, float scale, int rowsPerWI)
130{
131    int du = get_global_id(0);
132    int dv0 = get_global_id(1) * rowsPerWI;
133
134    if (du < cols)
135    {
136        int xmap_index = mad24(dv0, xmap_step, mad24(du, (int)sizeof(float), xmap_offset));
137        int ymap_index = mad24(dv0, ymap_step, mad24(du, (int)sizeof(float), ymap_offset));
138
139        float u = (tl_u + du) * scale;
140        float cosu, sinu = sincos(u, &cosu);
141
142        for (int dv = dv0, dv1 = min(rows, dv0 + rowsPerWI); dv < dv1; ++dv, xmap_index += xmap_step,
143            ymap_index += ymap_step)
144        {
145            __global float * xmap = (__global float *)(xmapptr + xmap_index);
146            __global float * ymap = (__global float *)(ymapptr + ymap_index);
147
148            float v = (tl_v + dv) * scale;
149
150            float cosv, sinv = sincos(v, &cosv);
151            float x_ = sinv * sinu;
152            float y_ = -cosv;
153            float z_ = sinv * cosu;
154
155            float x, y, z;
156            x = fma(ck_rinv[0], x_, fma(ck_rinv[1], y_, ck_rinv[2] * z_));
157            y = fma(ck_rinv[3], x_, fma(ck_rinv[4], y_, ck_rinv[5] * z_));
158            z = fma(ck_rinv[6], x_, fma(ck_rinv[7], y_, ck_rinv[8] * z_));
159
160            if (z > 0)
161                x /= z, y /= z;
162            else
163                x = y = -1;
164
165            xmap[0] = x;
166            ymap[0] = y;
167        }
168    }
169}
170