First commit [09/03/2018]
This commit is contained in:
Pierre8rTeam
@@ -0,0 +1,31 @@
|
||||
//--- 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;
|
||||
};
|
||||
//+------------------------------------------------------------------+
|
||||
@@ -0,0 +1,156 @@
|
||||
//--- 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
|
||||
//+------------------------------------------------------------------+
|
||||
//| fft_init OpenCL kernel for Fast Fourier Transfrom |
|
||||
//+------------------------------------------------------------------+
|
||||
//| Matthew Scarpino, "OpenCL in Action: How to accelerate graphics |
|
||||
//| and computations", Manning, 2012, Chapter 14. |
|
||||
//+------------------------------------------------------------------+
|
||||
__kernel void fft_init(__global double2 *in_data,
|
||||
__global double2 *out_data,
|
||||
__local double2 *l_data,
|
||||
uint points_per_group,uint size,int dir)
|
||||
{
|
||||
uint4 br,index;
|
||||
uint points_per_item,g_addr,l_addr,i,fft_index,stage,N2;
|
||||
double2 x1,x2,x3,x4,sum12,diff12,sum34,diff34;
|
||||
|
||||
points_per_item=points_per_group/get_local_size(0);
|
||||
l_addr = get_local_id(0)*points_per_item;
|
||||
g_addr = get_group_id(0)*points_per_group + l_addr;
|
||||
//--- load data from bit-reversed addresses and perform 4-point FFTs
|
||||
for(i=0; i<points_per_item; i+=4)
|
||||
{
|
||||
index=(uint4)(g_addr,g_addr+1,g_addr+2,g_addr+3);
|
||||
fft_index=size/2;
|
||||
stage=1;
|
||||
N2 =(uint)log2((double)size)-1;
|
||||
br =(index<< N2) & fft_index;
|
||||
br|=(index>> N2) & stage;
|
||||
//--- bit-reverse addresses
|
||||
while(N2>1)
|
||||
{
|
||||
N2-=2;
|
||||
fft_index>>=1;
|
||||
stage<<=1;
|
||||
br |= (index << N2) & fft_index;
|
||||
br |= (index >> N2) & stage;
|
||||
}
|
||||
//--- load global data
|
||||
x1 = in_data[br.s0];
|
||||
x2 = in_data[br.s1];
|
||||
x3 = in_data[br.s2];
|
||||
x4 = in_data[br.s3];
|
||||
|
||||
sum12=x1+x2;
|
||||
diff12= x1-x2;
|
||||
sum34 = x3+x4;
|
||||
diff34=(double2)(x3.s1-x4.s1,x4.s0-x3.s0)*dir;
|
||||
l_data[l_addr]=sum12+sum34;
|
||||
l_data[l_addr+1] = diff12 + diff34;
|
||||
l_data[l_addr+2] = sum12 - sum34;
|
||||
l_data[l_addr+3] = diff12 - diff34;
|
||||
l_addr += 4;
|
||||
g_addr += 4;
|
||||
}
|
||||
//--- perform initial stages of the FFT - each of length N2*2
|
||||
for(N2=4; N2<points_per_item; N2<<=1)
|
||||
{
|
||||
l_addr=get_local_id(0)*points_per_item;
|
||||
for(fft_index=0; fft_index<points_per_item; fft_index+=2*N2)
|
||||
{
|
||||
x1=l_data[l_addr];
|
||||
l_data[l_addr]+=l_data[l_addr+N2];
|
||||
l_data[l_addr+N2]=x1-l_data[l_addr+N2];
|
||||
for(i=1; i<N2; i++)
|
||||
{
|
||||
x3.s0=cos(M_PI_F*i/N2);
|
||||
x3.s1=dir*sin(M_PI_F*i/N2);
|
||||
x2=(double2)(l_data[l_addr+N2+i].s0*x3.s0+l_data[l_addr+N2+i].s1*x3.s1,l_data[l_addr+N2+i].s1*x3.s0-l_data[l_addr+N2+i].s0*x3.s1);
|
||||
l_data[l_addr+N2+i]=l_data[l_addr+i]-x2;
|
||||
l_data[l_addr+i]+=x2;
|
||||
}
|
||||
l_addr+=2*N2;
|
||||
}
|
||||
}
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
//--- perform FFT with other items in group - each of length N2*2
|
||||
stage=2;
|
||||
for(N2=points_per_item; N2<points_per_group; N2<<=1)
|
||||
{
|
||||
br.s0=(get_local_id(0)+(get_local_id(0)/stage)*stage) *(points_per_item/2);
|
||||
size = br.s0 % (N2*2);
|
||||
for(i=br.s0; i<br.s0+points_per_item/2; i++)
|
||||
{
|
||||
x3.s0=cos(M_PI_F*size/N2);
|
||||
x3.s1=dir*sin(M_PI_F*size/N2);
|
||||
x2=(double2)(l_data[N2+i].s0*x3.s0+l_data[N2+i].s1*x3.s1,l_data[N2+i].s1*x3.s0-l_data[N2+i].s0*x3.s1);
|
||||
l_data[N2+i]=l_data[i]-x2;
|
||||
l_data[i]+=x2;
|
||||
size++;
|
||||
}
|
||||
stage<<=1;
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
}
|
||||
//--- store results in global memory
|
||||
l_addr = get_local_id(0)*points_per_item;
|
||||
g_addr = get_group_id(0)*points_per_group + l_addr;
|
||||
for(i=0; i<points_per_item; i+=4)
|
||||
{
|
||||
out_data[g_addr]=l_data[l_addr];
|
||||
out_data[g_addr+1] = l_data[l_addr+1];
|
||||
out_data[g_addr+2] = l_data[l_addr+2];
|
||||
out_data[g_addr+3] = l_data[l_addr+3];
|
||||
g_addr += 4;
|
||||
l_addr += 4;
|
||||
}
|
||||
}
|
||||
//+------------------------------------------------------------------+
|
||||
//| fft_stage OpenCL kernel for Fast Fourier Transfrom |
|
||||
//+------------------------------------------------------------------+
|
||||
//| Matthew Scarpino, "OpenCL in Action: How to accelerate graphics |
|
||||
//| and computations", Manning, 2012, Chapter 14. |
|
||||
//+------------------------------------------------------------------+
|
||||
__kernel void fft_stage(__global double2 *g_data,uint stage,uint points_per_group,int dir)
|
||||
{
|
||||
uint points_per_item,addr,N,ang,i;
|
||||
double c,s;
|
||||
double2 input1,input2,w;
|
||||
|
||||
points_per_item=points_per_group/get_local_size(0);
|
||||
addr=(get_group_id(0)+(get_group_id(0)/stage)*stage)*(points_per_group/2)+get_local_id(0)*(points_per_item/2);
|
||||
N=points_per_group*(stage/2);
|
||||
ang=addr%(N*2);
|
||||
|
||||
for(i=addr; i<addr+points_per_item/2; i++)
|
||||
{
|
||||
c = cos(M_PI_F*ang/N);
|
||||
s = dir*sin(M_PI_F*ang/N);
|
||||
input1 = g_data[i];
|
||||
input2 = g_data[i+N];
|
||||
w=(double2)(input2.s0*c+input2.s1*s,input2.s1*c-input2.s0*s);
|
||||
g_data[i]=input1+w;
|
||||
g_data[i+N]=input1-w;
|
||||
ang++;
|
||||
}
|
||||
}
|
||||
//+------------------------------------------------------------------+
|
||||
//| fft_scale OpenCL kernel for Fast Fourier Transfrom |
|
||||
//+------------------------------------------------------------------+
|
||||
//| Matthew Scarpino, "OpenCL in Action: How to accelerate graphics |
|
||||
//| and computations", Manning, 2012, Chapter 14. |
|
||||
//+------------------------------------------------------------------+
|
||||
__kernel void fft_scale(__global double2 *g_data,uint points_per_group,uint scale)
|
||||
{
|
||||
uint points_per_item,addr,i;
|
||||
|
||||
points_per_item=points_per_group/get_local_size(0);
|
||||
addr=get_group_id(0)*points_per_group+get_local_id(0)*points_per_item;
|
||||
|
||||
for(i=addr; i<addr+points_per_item; i++)
|
||||
{
|
||||
g_data[i]/=scale;
|
||||
}
|
||||
}
|
||||
//+------------------------------------------------------------------+
|
||||
@@ -0,0 +1,69 @@
|
||||
//--- 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 for matrix multiplication |
|
||||
//| using global work groups |
|
||||
//+-----------------------------------------------------------+
|
||||
//| http://gpgpu-computing4.blogspot.ru/2009/09/ |
|
||||
//| /matrix-multiplication-2-opencl.html |
|
||||
//+-----------------------------------------------------------+
|
||||
__kernel void MatrixMult_GPU1(__global double *matrix_a,
|
||||
__global double *matrix_b,
|
||||
__global double *matrix_c,
|
||||
int rows_a,int cols_a,int cols_b)
|
||||
{
|
||||
int i=get_global_id(0);
|
||||
int j=get_global_id(1);
|
||||
double sum=0.0;
|
||||
for(int k=0; k<cols_a; k++)
|
||||
{
|
||||
sum+=matrix_a[cols_a*i+k]*matrix_b[cols_b*k+j];
|
||||
}
|
||||
matrix_c[cols_b*i+j]=sum;
|
||||
}
|
||||
#define BLOCK_SIZE 10
|
||||
//+-----------------------------------------------------------+
|
||||
//| OpenCL kernel for matrix multiplication |
|
||||
//| using local groups with common local memory |
|
||||
//+-----------------------------------------------------------+
|
||||
//| http://gpgpu-computing4.blogspot.ru/2009/10/ |
|
||||
//| /matrix-multiplication-3-opencl.html |
|
||||
//+-----------------------------------------------------------+
|
||||
__kernel void MatrixMult_GPU2(__global double *matrix_a,
|
||||
__global double *matrix_b,
|
||||
__global double *matrix_c,
|
||||
int rows_a,int cols_a,int cols_b)
|
||||
{
|
||||
int group_i=get_group_id(0);
|
||||
int group_j=get_group_id(1);
|
||||
|
||||
int i=get_local_id(0);
|
||||
int j=get_local_id(1);
|
||||
|
||||
int offset_b=BLOCK_SIZE*group_i;
|
||||
int offset_a_start=cols_a*BLOCK_SIZE*group_j;
|
||||
double sum=(float)0.0;
|
||||
|
||||
for(int offset_a=offset_a_start;
|
||||
offset_a<offset_a_start+cols_a;
|
||||
offset_a+=BLOCK_SIZE,
|
||||
offset_b+=BLOCK_SIZE*cols_b)
|
||||
{
|
||||
__local double submatrix_a[BLOCK_SIZE][BLOCK_SIZE];
|
||||
__local double submatrix_b[BLOCK_SIZE][BLOCK_SIZE];
|
||||
|
||||
submatrix_a[i][j]=matrix_a[offset_a+cols_a*i+j];
|
||||
submatrix_b[i][j]=matrix_b[offset_b+cols_b*i+j];
|
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
for(int k=0; k<BLOCK_SIZE; k++)
|
||||
sum+=submatrix_a[i][k]*submatrix_b[k][j];
|
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
}
|
||||
int offset_c=BLOCK_SIZE*(cols_b*group_j+group_i);
|
||||
matrix_c[offset_c+cols_b*i+j]=sum;
|
||||
};
|
||||
//+------------------------------------------------------------------+
|
||||
@@ -0,0 +1,51 @@
|
||||
//--- 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
|
||||
//+------------------------------------------------------------------+
|
||||
//| Morlet wavelet function |
|
||||
//+------------------------------------------------------------------+
|
||||
double Morlet(const double t)
|
||||
{
|
||||
return exp(-t*t*0.5)*cos(M_2_PI*t);
|
||||
}
|
||||
//+------------------------------------------------------------------+
|
||||
//| OpenCL kernel function |
|
||||
//+------------------------------------------------------------------+
|
||||
__kernel void Wavelet_GPU(__global double *data,int datacount,int x_size,int y_size,__global double *result)
|
||||
{
|
||||
size_t i = get_global_id(0);
|
||||
size_t j = get_global_id(1);
|
||||
double a1=(double)10e-10;
|
||||
double a2=(double)15.0;
|
||||
double da=(a2-a1)/(double)y_size;
|
||||
double db=((double)datacount-(double)0.0)/x_size;
|
||||
double a=a1+j*da;
|
||||
double b=0+i*db;
|
||||
uint norm=1;
|
||||
double B=(double)1.0; //Morlet
|
||||
double B_inv=(double)1.0/B;
|
||||
double a_inv=(double)1.0/a;
|
||||
double dt=(double)1.0;
|
||||
double coef=(double)0.0;
|
||||
if(norm==0)
|
||||
coef=sqrt(a_inv);
|
||||
else
|
||||
{
|
||||
for(int k=0; k<datacount; k++)
|
||||
{
|
||||
double arg=(dt*k-b)*a_inv;
|
||||
arg=-B_inv*arg*arg;
|
||||
coef=coef+exp(arg);
|
||||
}
|
||||
}
|
||||
double sum=(float)0.0;
|
||||
for(int k=0; k<datacount; k++)
|
||||
{
|
||||
double arg=(dt*k-b)*a_inv;
|
||||
sum+=data[k]*Morlet(arg);
|
||||
}
|
||||
sum=sum/coef;
|
||||
uint pos=(int)(j*x_size+i);
|
||||
result[pos]=sum;
|
||||
};
|
||||
//+------------------------------------------------------------------+
|
||||
Reference in New Issue
Block a user