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