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, Institute Of Software Chinese Academy Of Science, 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// Zhang Ying, zhangying913@gmail.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#ifdef OP_MAX_EIGEN_VAL 47 48__kernel void maxEigenVal(__global const uchar * srcptr, int src_step, int src_offset, int cols, 49 int total, __global uchar * dstptr 50#ifdef HAVE_MASK 51 , __global const uchar * maskptr, int mask_step, int mask_offset 52#endif 53 ) 54{ 55 int lid = get_local_id(0); 56 int gid = get_group_id(0); 57 int id = get_global_id(0); 58 59 __local float localmem_max[WGS2_ALIGNED]; 60 float maxval = -FLT_MAX; 61 62 for (int grain = groupnum * WGS; id < total; id += grain) 63 { 64 int src_index = mad24(id / cols, src_step, mad24((id % cols), (int)sizeof(float), src_offset)); 65#ifdef HAVE_MASK 66 int mask_index = mad24(id / cols, mask_step, id % cols + mask_offset); 67 if (maskptr[mask_index]) 68#endif 69 maxval = max(maxval, *(__global const float *)(srcptr + src_index)); 70 } 71 72 if (lid < WGS2_ALIGNED) 73 localmem_max[lid] = maxval; 74 barrier(CLK_LOCAL_MEM_FENCE); 75 76 if (lid >= WGS2_ALIGNED && total >= WGS2_ALIGNED) 77 localmem_max[lid - WGS2_ALIGNED] = max(maxval, localmem_max[lid - WGS2_ALIGNED]); 78 barrier(CLK_LOCAL_MEM_FENCE); 79 80 for (int lsize = WGS2_ALIGNED >> 1; lsize > 0; lsize >>= 1) 81 { 82 if (lid < lsize) 83 { 84 int lid2 = lsize + lid; 85 localmem_max[lid] = max(localmem_max[lid], localmem_max[lid2]); 86 } 87 barrier(CLK_LOCAL_MEM_FENCE); 88 } 89 90 if (lid == 0) 91 *(__global float *)(dstptr + (int)sizeof(float) * gid) = localmem_max[0]; 92} 93 94__kernel void maxEigenValTask(__global float * dst, float qualityLevel, 95 __global int * cornersptr) 96{ 97 float maxval = -FLT_MAX; 98 99 #pragma unroll 100 for (int x = 0; x < groupnum; ++x) 101 maxval = max(maxval, dst[x]); 102 103 dst[0] = maxval * qualityLevel; 104 cornersptr[0] = 0; 105} 106 107#elif OP_FIND_CORNERS 108 109#define GET_SRC_32F(_y, _x) *(__global const float *)(eigptr + (_y) * eig_step + (_x) * (int)sizeof(float) ) 110 111__kernel void findCorners(__global const uchar * eigptr, int eig_step, int eig_offset, 112#ifdef HAVE_MASK 113 __global const uchar * mask, int mask_step, int mask_offset, 114#endif 115 __global uchar * cornersptr, int rows, int cols, 116 __constant float * threshold, int max_corners) 117{ 118 int x = get_global_id(0); 119 int y = get_global_id(1); 120 121 __global int* counter = (__global int*) cornersptr; 122 __global float2 * corners = (__global float2 *)(cornersptr + (int)sizeof(float2)); 123 124 if (y < rows && x < cols 125#ifdef HAVE_MASK 126 && mask[mad24(y, mask_step, x + mask_offset)] 127#endif 128 ) 129 { 130 ++x, ++y; 131 float val = GET_SRC_32F(y, x); 132 133 if (val > threshold[0]) 134 { 135 float maxVal = val; 136 maxVal = max(GET_SRC_32F(y - 1, x - 1), maxVal); 137 maxVal = max(GET_SRC_32F(y - 1, x ), maxVal); 138 maxVal = max(GET_SRC_32F(y - 1, x + 1), maxVal); 139 140 maxVal = max(GET_SRC_32F(y , x - 1), maxVal); 141 maxVal = max(GET_SRC_32F(y , x + 1), maxVal); 142 143 maxVal = max(GET_SRC_32F(y + 1, x - 1), maxVal); 144 maxVal = max(GET_SRC_32F(y + 1, x ), maxVal); 145 maxVal = max(GET_SRC_32F(y + 1, x + 1), maxVal); 146 147 if (val == maxVal) 148 { 149 int ind = atomic_inc(counter); 150 if (ind < max_corners) 151 { 152 // pack and store eigenvalue and its coordinates 153 corners[ind].x = val; 154 corners[ind].y = as_float(y | (x << 16)); 155 } 156 } 157 } 158 } 159} 160 161#endif 162