modules/ocl/src/opencl/arithm_bitwise_not.cl

   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 //    Jiang Liyuan, jlyuan001.good@163.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 defined (DOUBLE_SUPPORT)
  47 #ifdef cl_khr_fp64
  48 #pragma OPENCL EXTENSION cl_khr_fp64:enable
  49 #elif defined (cl_amd_fp64)
  50 #pragma OPENCL EXTENSION cl_amd_fp64:enable
  51 #endif
  52 #endif
  53
  54 ///////////////////////////////////////////////////////////////////////////////////////////////////////
  55 ////////////////////////////////////////////BITWISE_NOT////////////////////////////////////////////////
  56 ///////////////////////////////////////////////////////////////////////////////////////////////////////
  57
  58 __kernel void arithm_bitwise_not_D0 (__global uchar *src1, int src1_step, int src1_offset,
  59                                      __global uchar *dst,  int dst_step,  int dst_offset,
  60                                      int rows, int cols, int dst_step1)
  61 {
  62     int x = get_global_id(0);
  63     int y = get_global_id(1);
  64
  65     if (x < cols && y < rows)
  66     {
  67         x = x << 2;
  68         int src1_index = mad24(y, src1_step, x + src1_offset);
  69
  70         int dst_end    = mad24(y, dst_step, dst_offset + dst_step1);
  71         int dst_index  = mad24(y, dst_step, dst_offset + x);
  72
  73         uchar4 src1_data = vload4(0, src1 + src1_index);
  74         uchar4 dst_data = vload4(0, dst + dst_index);
  75         uchar4 tmp_data = ~src1_data;
  76
  77         dst_data.x = dst_index + 0 < dst_end ? tmp_data.x : dst_data.x;
  78         dst_data.y = dst_index + 1 < dst_end ? tmp_data.y : dst_data.y;
  79         dst_data.z = dst_index + 2 < dst_end ? tmp_data.z : dst_data.z;
  80         dst_data.w = dst_index + 3 < dst_end ? tmp_data.w : dst_data.w;
  81
  82         vstore4(dst_data, 0, dst + dst_index);
  83     }
  84 }
  85
  86
  87 __kernel void arithm_bitwise_not_D1 (__global char *src1, int src1_step, int src1_offset,
  88                                      __global char *dst,  int dst_step,  int dst_offset,
  89                                      int rows, int cols, int dst_step1)
  90 {
  91     int x = get_global_id(0);
  92     int y = get_global_id(1);
  93
  94     if (x < cols && y < rows)
  95     {
  96         x = x << 2;
  97         int src1_index = mad24(y, src1_step, x + src1_offset);
  98
  99         int dst_end    = mad24(y, dst_step, dst_offset + dst_step1);
 100         int dst_index  = mad24(y, dst_step, dst_offset + x);
 101
 102         char4 src1_data = vload4(0, src1 + src1_index);
 103         char4 dst_data = vload4(0, dst + dst_index);
 104         char4 tmp_data = ~src1_data;
 105
 106         dst_data.x = dst_index + 0 < dst_end ? tmp_data.x : dst_data.x;
 107         dst_data.y = dst_index + 1 < dst_end ? tmp_data.y : dst_data.y;
 108         dst_data.z = dst_index + 2 < dst_end ? tmp_data.z : dst_data.z;
 109         dst_data.w = dst_index + 3 < dst_end ? tmp_data.w : dst_data.w;
 110
 111         vstore4(dst_data, 0, dst + dst_index);
 112     }
 113 }
 114
 115
 116 __kernel void arithm_bitwise_not_D2 (__global ushort *src1, int src1_step, int src1_offset,
 117                                      __global ushort *dst,  int dst_step,  int dst_offset,
 118                                      int rows, int cols, int dst_step1)
 119
 120 {
 121     int x = get_global_id(0);
 122     int y = get_global_id(1);
 123
 124     if (x < cols && y < rows)
 125     {
 126         x = x << 2;
 127
 128 #ifdef dst_align
 129 #undef dst_align
 130 #endif
 131 #define dst_align ((dst_offset >> 1) & 3)
 132         int src1_index = mad24(y, src1_step, (x << 1) + src1_offset - (dst_align << 1));
 133
 134         int dst_start  = mad24(y, dst_step, dst_offset);
 135         int dst_end    = mad24(y, dst_step, dst_offset + dst_step1);
 136         int dst_index  = mad24(y, dst_step, dst_offset + (x << 1) & (int)0xfffffff8);
 137
 138         ushort4 src1_data = vload4(0, (__global ushort *)((__global char *)src1 + src1_index));
 139
 140         ushort4 dst_data = *((__global ushort4 *)((__global char *)dst + dst_index));
 141         ushort4 tmp_data = ~ src1_data;
 142
 143         dst_data.x = ((dst_index + 0 >= dst_start) && (dst_index + 0 < dst_end)) ? tmp_data.x : dst_data.x;
 144         dst_data.y = ((dst_index + 2 >= dst_start) && (dst_index + 2 < dst_end)) ? tmp_data.y : dst_data.y;
 145         dst_data.z = ((dst_index + 4 >= dst_start) && (dst_index + 4 < dst_end)) ? tmp_data.z : dst_data.z;
 146         dst_data.w = ((dst_index + 6 >= dst_start) && (dst_index + 6 < dst_end)) ? tmp_data.w : dst_data.w;
 147
 148         *((__global ushort4 *)((__global char *)dst + dst_index)) = dst_data;
 149     }
 150 }
 151
 152
 153
 154 __kernel void arithm_bitwise_not_D3 (__global short *src1, int src1_step, int src1_offset,
 155                                      __global short *dst,  int dst_step,  int dst_offset,
 156                                      int rows, int cols, int dst_step1)
 157
 158 {
 159     int x = get_global_id(0);
 160     int y = get_global_id(1);
 161
 162     if (x < cols && y < rows)
 163     {
 164         x = x << 2;
 165
 166 #ifdef dst_align
 167 #undef dst_align
 168 #endif
 169 #define dst_align ((dst_offset >> 1) & 3)
 170         int src1_index = mad24(y, src1_step, (x << 1) + src1_offset - (dst_align << 1));
 171
 172         int dst_start  = mad24(y, dst_step, dst_offset);
 173         int dst_end    = mad24(y, dst_step, dst_offset + dst_step1);
 174         int dst_index  = mad24(y, dst_step, dst_offset + (x << 1) & (int)0xfffffff8);
 175
 176         short4 src1_data = vload4(0, (__global short *)((__global char *)src1 + src1_index));
 177
 178         short4 dst_data = *((__global short4 *)((__global char *)dst + dst_index));
 179         short4 tmp_data = ~ src1_data;
 180
 181         dst_data.x = ((dst_index + 0 >= dst_start) && (dst_index + 0 < dst_end)) ? tmp_data.x : dst_data.x;
 182         dst_data.y = ((dst_index + 2 >= dst_start) && (dst_index + 2 < dst_end)) ? tmp_data.y : dst_data.y;
 183         dst_data.z = ((dst_index + 4 >= dst_start) && (dst_index + 4 < dst_end)) ? tmp_data.z : dst_data.z;
 184         dst_data.w = ((dst_index + 6 >= dst_start) && (dst_index + 6 < dst_end)) ? tmp_data.w : dst_data.w;
 185
 186         *((__global short4 *)((__global char *)dst + dst_index)) = dst_data;
 187     }
 188 }
 189
 190
 191
 192 __kernel void arithm_bitwise_not_D4 (__global int *src1, int src1_step, int src1_offset,
 193                                      __global int *dst,  int dst_step,  int dst_offset,
 194                                      int rows, int cols, int dst_step1)
 195 {
 196     int x = get_global_id(0);
 197     int y = get_global_id(1);
 198
 199     if (x < cols && y < rows)
 200     {
 201         int src1_index = mad24(y, src1_step, (x << 2) + src1_offset);
 202         int dst_index  = mad24(y, dst_step,  (x << 2) + dst_offset);
 203
 204         int data1 = *((__global int *)((__global char *)src1 + src1_index));
 205         int tmp  = ~ data1;
 206
 207         *((__global int *)((__global char *)dst + dst_index)) = tmp;
 208     }
 209 }
 210
 211 __kernel void arithm_bitwise_not_D5 (__global char *src, int src_step, int src_offset,
 212                                      __global char *dst, int dst_step, int dst_offset,
 213                                      int rows, int cols, int dst_step1)
 214 {
 215     int x = get_global_id(0);
 216     int y = get_global_id(1);
 217
 218     if (x < cols && y < rows)
 219     {
 220         int src_index = mad24(y, src_step, (x << 2) + src_offset);
 221         int dst_index = mad24(y, dst_step, (x << 2) + dst_offset);
 222
 223         char4 data;
 224
 225         data = *((__global char4 *)((__global char *)src + src_index));
 226         data = ~ data;
 227
 228         *((__global char4 *)((__global char *)dst + dst_index)) = data;
 229     }
 230 }
 231
 232 #if defined (DOUBLE_SUPPORT)
 233 __kernel void arithm_bitwise_not_D6 (__global char *src, int src_step, int src_offset,
 234                                      __global char *dst, int dst_step, int dst_offset,
 235                                      int rows, int cols, int dst_step1)
 236 {
 237     int x = get_global_id(0);
 238     int y = get_global_id(1);
 239
 240     if (x < cols && y < rows)
 241     {
 242         int src_index = mad24(y, src_step, (x << 3) + src_offset);
 243         int dst_index = mad24(y, dst_step,  (x << 3) + dst_offset);
 244
 245         char8 data;
 246
 247         data = *((__global char8 *)((__global char *)src + src_index));
 248         data = ~ data;
 249
 250         *((__global char8 *)((__global char *)dst + dst_index)) = data;
 251     }
 252 }
 253 #endif