28 lines
1.3 KiB
Common Lisp
28 lines
1.3 KiB
Common Lisp
//+-----------------------------------------------------------+
|
|
//| OpenCL kernel |
|
|
//| The bitonic sort kernel does an ascending sort. |
|
|
//+-----------------------------------------------------------+
|
|
//| R. Banger,K. Bhattacharyya, OpenCL Programming by Example:|
|
|
//| A comprehensive guide on OpenCL programming with examples |
|
|
//| PACKT Publishing, 2013. |
|
|
//+-----------------------------------------------------------+
|
|
__kernel void BitonicSort_GPU(__global float *data,const uint stage,const uint pass)
|
|
{
|
|
uint id=get_global_id(0);
|
|
uint distance = 1<<(stage-pass);
|
|
uint left_id =(id &(distance-1));
|
|
left_id+=(id>>(stage-pass))*(distance<<1);
|
|
uint right_id=left_id+distance;
|
|
float left_value=data[left_id];
|
|
float right_value=data[right_id];
|
|
uint same_direction=(id>>stage)&0x1;
|
|
uint temp = same_direction?right_id:temp;
|
|
right_id = same_direction?left_id:right_id;
|
|
left_id = same_direction?temp:left_id;
|
|
int compare_res=(left_value<right_value);
|
|
float greater = compare_res?right_value:left_value;
|
|
float lesser = compare_res?left_value:right_value;
|
|
data[left_id] = lesser;
|
|
data[right_id]= greater;
|
|
};
|
|
//+------------------------------------------------------------------+
|