#define p 1
#define M 128
#define L 8
device void kalman_call(floatF,floatP0,float Q,floatR,floatS_,floaty_id,floatP_,float*Sp)
{
S_[0]= F[0]*Sp[0]; //S_
P_[0]= F[0]*P0[0]*F[0]+Q; //P_
Sp[0]=P_[0]/(P_[0]+R[0]); //K
P0[0]=P_[0]-Sp[0]P_[0]; // P=(I-KH)*P_;
Sp[0]=S_[0]+Sp[0]*(y_id[0]-S_[0]); //SOut
}
global void dct_kalman(floatR,floaty_id, float *x_id,float *x_out,unsigned int frameNum)
{
float Q=0;
unsigned int tx=threadIdx.x; //索引，线程块大小=dim3 Block(M,p,1); 网格大小=dim3 Grid(1,1);
shared float Sp[M];
shared float A[(p+1)*M];
shared float S_[M];
shared float P0[M];
shared float P_[M];
shared float F[M];
shared float autm[(L+1)*M];
Sp[tx] = y_id[tx]; //y_id[512][128]，现在只需要把第0行的数据读进Sp。
__syncthreads(); //同步，确实读入数据完成
P0[tx]=R[tx+M*L]; //R[512][128]，现在只需要把第L(L=8)行的数据读进P0。
__syncthreads(); //同步，确实读入数据完成
for(int t=L+1,n=p;t<frameNum;t++,n++) //frameNum=512
{
Q=lpc(p,A+tx,x_id+tx,autm+tx,L+1); // lpc是device函数，每个线程调用一次，单独验证过此函数，是OK的。
//__syncthreads(); //请问这里需要加同步吗？需要保证上面的device完成了吗？还是加__threadfence_block();
Q=max(Q,0.0016y_id[tx+tM]y_id[tx+tM]); //#define max(a,b) ((a > b) ? a : b)比较Q与y_id里t行每个数的大小
//syncthreads(); //请问这里需要加同步吗？需要保证上面的device完成了吗？还是加__threadfence_block();
F[tx]=A[tx+M](-1); //把上面device函数lpc算出的A数组中的第1行乘以负1读进F数组。
//__syncthreads(); //请问这里需要加同步吗？需要保证上面的device完成了吗？还是加__threadfence_block();
kalman_call(F+tx,P0+tx,Q,R+tx+tM,S+tx,y_id+tx+t*M,P+tx,Sp+tx);kalman_call是device函数，单独验证过了。
//__syncthreads(); //请问这里需要加同步吗？需要保证上面的device完成了吗？还是加__threadfence_block();
x_out[(n-p+1)*M+tx]=Sp[tx]; //把上面kalman_call算出的Sp存储回显存x_out，第一次存储在第1行，第2次存第2行
__syncthreads(); //请问这里需要加同步吗？需要保证上面的device完成了吗？还是加__threadfence_block();
}
}
这个是我第一次写的CUDA函数，是卡尔曼滤波，一共有M=128个通道要做kalman_call（每个通道之间是不相关的），每个通道之内有一个迭代关系，我就用 for(int t=L+1,n=p;t<frameNum;t++,n++)来解决，因此在for循环内让每一个通道都调用一次kalman_call，但我发现算出来的结果是不对的。不知道写法上有没有不合理之处，请版主给一些意见。