1 /*M///////////////////////////////////////////////////////////////////////////////////////
3 // IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
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.
11 // For Open Source Computer Vision Library
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.
18 // Peng Xiao, pengxiao@outlook.com
20 // Redistribution and use in source and binary forms, with or without modification,
21 // are permitted provided that the following conditions are met:
23 // * Redistribution's of source code must retain the above copyright notice,
24 // this list of conditions and the following disclaimer.
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.
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.
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.
59 #define my_comp(x,y) ((x) > (y))
61 #define my_comp(x,y) ((x) < (y))
64 /////////////////////// Bitonic sort ////////////////////////////
66 // https://github.com/HSA-Libraries/Bolt/blob/master/include/bolt/cl/sort_by_key_kernels.cl
77 const int threadId = get_global_id(0);
78 if(threadId >= count / 2)
82 const int pairDistance = 1 << (stage - passOfStage);
83 const int blockWidth = 2 * pairDistance;
85 int leftId = min( (threadId % pairDistance)
86 + (threadId / pairDistance) * blockWidth, count );
88 int rightId = min( leftId + pairDistance, count );
92 const V_T lval = vals[leftId];
93 const V_T rval = vals[rightId];
95 const K_T lkey = keys[leftId];
96 const K_T rkey = keys[rightId];
98 int sameDirectionBlockWidth = 1 << stage;
100 if((threadId/sameDirectionBlockWidth) % 2 == 1)
107 const bool compareResult = my_comp(lkey, rkey);
111 keys[rightId] = rkey;
113 vals[rightId] = rval;
118 keys[rightId] = lkey;
120 vals[rightId] = lval;
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
129 void selectionSortLocal
134 __local K_T * scratch
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
143 int offset = groupID * wg;
148 n = (groupID == (numOfGroups-1))? (count - wg*(numOfGroups-1)) : wg;
150 int clamped_i= min(i, n - 1);
152 K_T key1 = keys[clamped_i], key2;
153 V_T val1 = vals[clamped_i];
155 barrier(CLK_LOCAL_MEM_FENCE);
163 for (int j=0;j<n;++j)
166 if(my_comp(key2, key1))
167 pos++;//calculate the rank of this element in this work group
170 if(my_comp(key1, key2))
174 // key1 and key2 are same
179 for (int j=0; j< same; j++)
181 vals[pos + j] = val1;
182 keys[pos + j] = key1;
186 void selectionSortFinal
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);
201 if((offset + i ) >= count)
203 V_T val1 = vals[offset + i];
205 K_T key1 = keys[offset + i];
208 for(int j=0; j<numOfGroups-1; j++ )
210 for(int k=0; k<wg; k++)
212 key2 = keys[j*wg + k];
213 if(my_comp(key1, key2))
217 //Increment only if the value is not the same.
218 if(my_comp(key2, key1))
226 for(int k=0; k<remainder; k++)
228 key2 = keys[(numOfGroups-1)*wg + k];
229 if(my_comp(key1, key2))
233 //Don't increment if the value is the same.
234 if(my_comp(key2, key1))
240 for (int j=0; j< same; j++)
242 vals[pos + j] = val1;
243 keys[pos + j] = key1;