| //+-----------------------------------------------------------+
|
| //| 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;
|
| };
|
| //+------------------------------------------------------------------+
|
