31 lines
1.5 KiB
Common Lisp
31 lines
1.5 KiB
Common Lisp
 | //--- by default some GPU doesn't support doubles
| ||
//--- cl_khr_fp64 directive is used to enable work with doubles
| |||
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
| |||
//+-----------------------------------------------------------+
| |||
//| 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 double *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;
| |||
double left_value=data[left_id];
| |||
double 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);
| |||
double greater = compare_res?right_value:left_value;
| |||
double lesser = compare_res?left_value:right_value;
| |||
data[left_id] = lesser;
| |||
data[right_id]= greater;
| |||
};
| |||
//+------------------------------------------------------------------+
|