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 #if kercn != 3 47 #define storepix(val, addr) *(__global T *)(addr) = val 48 #define TSIZE (int)sizeof(T) 49 #define scalar scalar_ 50 #else 51 #define storepix(val, addr) vstore3(val, 0, (__global T1 *)(addr)) 52 #define TSIZE ((int)sizeof(T1)*3) 53 #define scalar (T)(scalar_.x, scalar_.y, scalar_.z) 54 #endif 55 56 __kernel void setIdentity(__global uchar * srcptr, int src_step, int src_offset, int rows, int cols, 57 ST scalar_) 58 { 59 int x = get_global_id(0); 60 int y0 = get_global_id(1) * rowsPerWI; 61 62 if (x < cols) 63 { 64 int src_index = mad24(y0, src_step, mad24(x, TSIZE, src_offset)); 65 66 #if kercn == cn 67 #pragma unroll 68 for (int y = y0, i = 0, y1 = min(rows, y0 + rowsPerWI); i < rowsPerWI; ++y, ++i, src_index += src_step) 69 if (y < y1) 70 storepix(x == y ? scalar : (T)(0), srcptr + src_index); 71 #elif kercn == 4 && cn == 1 72 if (y0 < rows) 73 { 74 storepix(x == y0 >> 2 ? (T)(scalar, 0, 0, 0) : (T)(0), srcptr + src_index); 75 if (++y0 < rows) 76 { 77 src_index += src_step; 78 storepix(x == y0 >> 2 ? (T)(0, scalar, 0, 0) : (T)(0), srcptr + src_index); 79 80 if (++y0 < rows) 81 { 82 src_index += src_step; 83 storepix(x == y0 >> 2 ? (T)(0, 0, scalar, 0) : (T)(0), srcptr + src_index); 84 85 if (++y0 < rows) 86 { 87 src_index += src_step; 88 storepix(x == y0 >> 2 ? (T)(0, 0, 0, scalar) : (T)(0), srcptr + src_index); 89 } 90 } 91 } 92 } 93 #else 94 #error "Incorrect combination of cn && kercn" 95 #endif 96 } 97 } 98