modules/ocl/src/opencl/kernel_sort_by_key.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, 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@outlook.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 oclMaterials 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 K_T
  47 #define K_T float
  48 #endif
  49
  50 #ifndef V_T
  51 #define V_T float
  52 #endif
  53
  54 #ifndef IS_GT
  55 #define IS_GT false
  56 #endif
  57
  58 #if IS_GT
  59 #define my_comp(x,y) ((x) > (y))
  60 #else
  61 #define my_comp(x,y) ((x) < (y))
  62 #endif
  63
  64 /////////////////////// Bitonic sort ////////////////////////////
  65 // ported from
  66 // https://github.com/HSA-Libraries/Bolt/blob/master/include/bolt/cl/sort_by_key_kernels.cl
  67 __kernel
  68     void bitonicSort
  69     (
  70         __global K_T * keys,
  71         __global V_T * vals,
  72         int count,
  73         int stage,
  74         int passOfStage
  75     )
  76 {
  77     const int threadId = get_global_id(0);
  78     if(threadId >= count / 2)
  79     {
  80         return;
  81     }
  82     const int pairDistance = 1 << (stage - passOfStage);
  83     const int blockWidth   = 2 * pairDistance;
  84
  85     int leftId = min( (threadId % pairDistance)
  86                    + (threadId / pairDistance) * blockWidth, count );
  87
  88     int rightId = min( leftId + pairDistance, count );
  89
  90     int temp;
  91
  92     const V_T lval = vals[leftId];
  93     const V_T rval = vals[rightId];
  94
  95     const K_T lkey = keys[leftId];
  96     const K_T rkey = keys[rightId];
  97
  98     int sameDirectionBlockWidth = 1 << stage;
  99
 100     if((threadId/sameDirectionBlockWidth) % 2 == 1)
 101     {
 102         temp = rightId;
 103         rightId = leftId;
 104         leftId = temp;
 105     }
 106
 107     const bool compareResult = my_comp(lkey, rkey);
 108
 109     if(compareResult)
 110     {
 111         keys[rightId] = rkey;
 112         keys[leftId]  = lkey;
 113         vals[rightId] = rval;
 114         vals[leftId]  = lval;
 115     }
 116     else
 117     {
 118         keys[rightId] = lkey;
 119         keys[leftId]  = rkey;
 120         vals[rightId] = lval;
 121         vals[leftId]  = rval;
 122     }
 123 }
 124
 125 /////////////////////// Selection sort ////////////////////////////
 126 //kernel is ported from Bolt library:
 127 //https://github.com/HSA-Libraries/Bolt/blob/master/include/bolt/cl/sort_kernels.cl
 128 __kernel
 129     void selectionSortLocal
 130     (
 131         __global K_T * keys,
 132         __global V_T * vals,
 133         const int count,
 134         __local  K_T * scratch
 135     )
 136 {
 137     int          i  = get_local_id(0); // index in workgroup
 138     int numOfGroups = get_num_groups(0); // index in workgroup
 139     int groupID     = get_group_id(0);
 140     int         wg  = get_local_size(0); // workgroup size = block size
 141     int n; // number of elements to be processed for this work group
 142
 143     int offset   = groupID * wg;
 144     int same     = 0;
 145
 146     vals      += offset;
 147     keys      += offset;
 148     n = (groupID == (numOfGroups-1))? (count - wg*(numOfGroups-1)) : wg;
 149
 150     int clamped_i= min(i, n - 1);
 151
 152     K_T key1 = keys[clamped_i], key2;
 153     V_T val1 = vals[clamped_i];
 154     scratch[i] = key1;
 155     barrier(CLK_LOCAL_MEM_FENCE);
 156
 157     if(i >= n)
 158     {
 159         return;
 160     }
 161
 162     int pos = 0;
 163     for (int j=0;j<n;++j)
 164     {
 165         key2  = scratch[j];
 166         if(my_comp(key2, key1))
 167             pos++;//calculate the rank of this element in this work group
 168         else
 169         {
 170             if(my_comp(key1, key2))
 171                 continue;
 172             else
 173             {
 174                 // key1 and key2 are same
 175                 same++;
 176             }
 177         }
 178     }
 179     for (int j=0; j< same; j++)
 180     {
 181         vals[pos + j] = val1;
 182         keys[pos + j] = key1;
 183     }
 184 }
 185 __kernel
 186     void selectionSortFinal
 187     (
 188         __global K_T * keys,
 189         __global V_T * vals,
 190         const int count
 191     )
 192 {
 193     const int          i  = get_local_id(0); // index in workgroup
 194     const int numOfGroups = get_num_groups(0); // index in workgroup
 195     const int groupID     = get_group_id(0);
 196     const int         wg  = get_local_size(0); // workgroup size = block size
 197     int pos = 0, same = 0;
 198     const int offset = get_group_id(0) * wg;
 199     const int remainder = count - wg*(numOfGroups-1);
 200
 201     if((offset + i ) >= count)
 202         return;
 203     V_T val1 = vals[offset + i];
 204
 205     K_T key1 = keys[offset + i];
 206     K_T key2;
 207
 208     for(int j=0; j<numOfGroups-1; j++ )
 209     {
 210         for(int k=0; k<wg; k++)
 211         {
 212             key2 = keys[j*wg + k];
 213             if(my_comp(key1, key2))
 214                 break;
 215             else
 216             {
 217                 //Increment only if the value is not the same.
 218                 if(my_comp(key2, key1))
 219                     pos++;
 220                 else
 221                     same++;
 222             }
 223         }
 224     }
 225
 226     for(int k=0; k<remainder; k++)
 227     {
 228         key2 = keys[(numOfGroups-1)*wg + k];
 229         if(my_comp(key1, key2))
 230             break;
 231         else
 232         {
 233             //Don't increment if the value is the same.
 234             if(my_comp(key2, key1))
 235                 pos++;
 236             else
 237                 same++;
 238         }
 239     }
 240     for (int j=0; j< same; j++)
 241     {
 242         vals[pos + j] = val1;
 243         keys[pos + j] = key1;
 244     }
 245 }