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 // Jin Ma jin (at) 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 #ifndef cn 47 #define cn 1 48 #endif 49 50 #define sz (int)sizeof(float) 51 #define src_elem_at(_src, y, step, x) *(__global const float *)(_src + mad24(y, step, (x) * sz)) 52 #define dst_elem_at(_dst, y, step, x) *(__global float *)(_dst + mad24(y, step, (x) * sz)) 53 54 __kernel void buildMotionMaps(__global const uchar * forwardMotionPtr, int forwardMotion_step, int forwardMotion_offset, 55 __global const uchar * backwardMotionPtr, int backwardMotion_step, int backwardMotion_offset, 56 __global const uchar * forwardMapPtr, int forwardMap_step, int forwardMap_offset, 57 __global const uchar * backwardMapPtr, int backwardMap_step, int backwardMap_offset, 58 int rows, int cols) 59 { 60 int x = get_global_id(0); 61 int y = get_global_id(1); 62 63 if (x < cols && y < rows) 64 { 65 int forwardMotion_index = mad24(forwardMotion_step, y, (int)sizeof(float2) * x + forwardMotion_offset); 66 int backwardMotion_index = mad24(backwardMotion_step, y, (int)sizeof(float2) * x + backwardMotion_offset); 67 int forwardMap_index = mad24(forwardMap_step, y, (int)sizeof(float2) * x + forwardMap_offset); 68 int backwardMap_index = mad24(backwardMap_step, y, (int)sizeof(float2) * x + backwardMap_offset); 69 70 float2 forwardMotion = *(__global const float2 *)(forwardMotionPtr + forwardMotion_index); 71 float2 backwardMotion = *(__global const float2 *)(backwardMotionPtr + backwardMotion_index); 72 __global float2 * forwardMap = (__global float2 *)(forwardMapPtr + forwardMap_index); 73 __global float2 * backwardMap = (__global float2 *)(backwardMapPtr + backwardMap_index); 74 75 float2 basePoint = (float2)(x, y); 76 77 forwardMap[0] = basePoint + backwardMotion; 78 backwardMap[0] = basePoint + forwardMotion; 79 } 80 } 81 82 __kernel void upscale(__global const uchar * srcptr, int src_step, int src_offset, int src_rows, int src_cols, 83 __global uchar * dstptr, int dst_step, int dst_offset, int scale) 84 { 85 int x = get_global_id(0); 86 int y = get_global_id(1); 87 88 if (x < src_cols && y < src_rows) 89 { 90 int src_index = mad24(y, src_step, sz * x * cn + src_offset); 91 int dst_index = mad24(y * scale, dst_step, sz * x * scale * cn + dst_offset); 92 93 __global const float * src = (__global const float *)(srcptr + src_index); 94 __global float * dst = (__global float *)(dstptr + dst_index); 95 96 #pragma unroll 97 for (int c = 0; c < cn; ++c) 98 dst[c] = src[c]; 99 } 100 } 101 102 103 inline float diffSign1(float a, float b) 104 { 105 return a > b ? 1.0f : a < b ? -1.0f : 0.0f; 106 } 107 108 inline float3 diffSign3(float3 a, float3 b) 109 { 110 float3 pos; 111 pos.x = a.x > b.x ? 1.0f : a.x < b.x ? -1.0f : 0.0f; 112 pos.y = a.y > b.y ? 1.0f : a.y < b.y ? -1.0f : 0.0f; 113 pos.z = a.z > b.z ? 1.0f : a.z < b.z ? -1.0f : 0.0f; 114 return pos; 115 } 116 117 __kernel void diffSign(__global const uchar * src1, int src1_step, int src1_offset, 118 __global const uchar * src2, int src2_step, int src2_offset, 119 __global uchar * dst, int dst_step, int dst_offset, int rows, int cols) 120 { 121 int x = get_global_id(0); 122 int y = get_global_id(1); 123 124 if (x < cols && y < rows) 125 *(__global float *)(dst + mad24(y, dst_step, sz * x + dst_offset)) = 126 diffSign1(*(__global const float *)(src1 + mad24(y, src1_step, sz * x + src1_offset)), 127 *(__global const float *)(src2 + mad24(y, src2_step, sz * x + src2_offset))); 128 } 129 130 __kernel void calcBtvRegularization(__global const uchar * src, int src_step, int src_offset, 131 __global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols, 132 int ksize, __constant float * c_btvRegWeights) 133 { 134 int x = get_global_id(0) + ksize; 135 int y = get_global_id(1) + ksize; 136 137 if (y < dst_rows - ksize && x < dst_cols - ksize) 138 { 139 src += src_offset; 140 141 #if cn == 1 142 const float srcVal = src_elem_at(src, y, src_step, x); 143 float dstVal = 0.0f; 144 145 for (int m = 0, count = 0; m <= ksize; ++m) 146 for (int l = ksize; l + m >= 0; --l, ++count) 147 { 148 dstVal += c_btvRegWeights[count] * (diffSign1(srcVal, src_elem_at(src, y + m, src_step, x + l)) 149 - diffSign1(src_elem_at(src, y - m, src_step, x - l), srcVal)); 150 } 151 152 dst_elem_at(dst, y, dst_step, x) = dstVal; 153 #elif cn == 3 154 __global const float * src0ptr = (__global const float *)(src + mad24(y, src_step, 3 * sz * x + src_offset)); 155 float3 srcVal = (float3)(src0ptr[0], src0ptr[1], src0ptr[2]), dstVal = 0.f; 156 157 for (int m = 0, count = 0; m <= ksize; ++m) 158 { 159 for (int l = ksize; l + m >= 0; --l, ++count) 160 { 161 __global const float * src1ptr = (__global const float *)(src + mad24(y + m, src_step, 3 * sz * (x + l) + src_offset)); 162 __global const float * src2ptr = (__global const float *)(src + mad24(y - m, src_step, 3 * sz * (x - l) + src_offset)); 163 164 float3 src1 = (float3)(src1ptr[0], src1ptr[1], src1ptr[2]); 165 float3 src2 = (float3)(src2ptr[0], src2ptr[1], src2ptr[2]); 166 167 dstVal += c_btvRegWeights[count] * (diffSign3(srcVal, src1) - diffSign3(src2, srcVal)); 168 } 169 } 170 171 __global float * dstptr = (__global float *)(dst + mad24(y, dst_step, 3 * sz * x + dst_offset + 0)); 172 dstptr[0] = dstVal.x; 173 dstptr[1] = dstVal.y; 174 dstptr[2] = dstVal.z; 175 #else 176 #error "Number of channels should be either 1 of 3" 177 #endif 178 } 179 } 180
