第一次写CUDA程序

#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+t
M,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,但我发现算出来的结果是不对的。不知道写法上有没有不合理之处,请版主给一些意见。

LZ您好:

1:我不懂卡尔曼滤波,因此无法告知您算法上哪里不合适,以下分析都基于您已经给出的信息。

2:如果您是一个线程完成一整套滤波,线程之间并无关系,也无依赖,也无交互,那么您完全可以不使用shared memory,每个线程将数据读入到本地的变量就行,在变量不是很多的情况下可以直接使用更快的寄存器存储。

3:您可以在kernel中使用threadIdx.x等内建索引变量,或者将算好的索引传进来,您这样__device__函数中写一堆S_[0]这样下标是0的元素,通过调用的时候算指针的偏移量,十分不直观。

4:无法猜出您其他没有提供实现的函数是否有问题。您提供kalman_call函数目测计算规则上并无问题,无法目测您的算法是否正确。

5:如果您需要整个block的线程都等在同一个地方,等待某个结果完成,再向下执行(比如等待shared memory写入完成),那么需要__syncthreads(),如无此需要,线程间是独立的,则无需同步。

大致如此,供您参考。

祝您编码顺利~

版主你好:
每个线程做的计算都是不相关的,独立的。按照你的意思,我只要在for循环里最后一行加一个同步就行吧。或者不加都行吧,反正每个线程只依赖本线程上一行产生的数据,不需要消费其他线程产生的数据。
for(int t=L+1,n=p;t<frameNum;t++,n++)
{
Q=lpc(p,A+tx,x_id+tx,autm+tx,L+1); //lpc
Q=max(Q,0.0016y_id[tx+tM]y_id[tx+tM]); //Q
F[tx]=A[tx+M]*(-1); //F
S_[tx]= F[tx]*Sp[tx]; //S_
P_[tx]= F[tx]*P0[tx]*F[tx]+Q; //P_
Sp[tx]=P_[tx]/(P_[tx]+*R); //K
P0[tx]=P_[tx]-Sp[tx]P_[tx]; // P=(I-KH)P_;
Sp[tx]=S_[tx]+Sp[tx]
(*y_id-S_[tx]); //SOut
x_out[(n-p+1)*M+tx]=Sp[tx]; //speech array
__syncthreads();
}

版主你好:
你上面建议的第2点,不使用shared,而直接使用寄存器,是这样吗?
global void dct_kalman(floatR,floaty_id,
float *x_id,float x_out,unsigned int frameNum)
{
float Q=0;
unsigned int tx=blockIdx.x
blockDim.x+threadIdx.x;

float Sp=0;
shared float A[(p+1)*M];
float S_=0;
float P0=0;
float P_=0;
float F=0;
shared float autm[(L+1)*M];

float Sp = y_id[tx]; //initialize the every column
syncthreads();
float P0=R[tx+ML];
__syncthreads();
for(int t=L+1,n=p+1;t<frameNum;t++,n++)
{
Q=lpc(p,A+tx,x_id+tx,autm+tx,L+1); //lpc
Q=max(Q,0.0016
y_id[tx+tM]y_id[tx+tM]); //Q
F=A[tx+M]
(-1); //F
S
= F*Sp; //S

P_= FP0F+Q; //P_
Sp=P_/(P_+R); //K
P0=P_-Sp
P_; // P=(I-K*H)P_;
Sp=S_+Sp
(*y_id-S_); //SOut
x_out[(n-p+1)*M+tx]=Sp; //speech array
_syncthreads();
}
}
寄存器存的是一个数字,当这些 S
[0]= F[0]*Sp[0];是小数组时,还是要读进共享内存吧。

版主,追问一下,如果把共享内存那里换成寄存器,虽然我开启的是一个线程块里128个线程,像上面那样写,会不会超出数量?

LZ您好:

如果每个线程的数据都仅与自己写入的shared memory的数据有关的话,可以不加同步,一个线程自己的访问总是前后顺序得以保证的。

以及,此时可以不使用shared memory,不再赘述。

祝您好运~

LZ您好:

代码的细节请您自己检查一下。

您之前是在shared memory中每个线程只使用自己对应的一个变量,现在则是直接将该变量定义在kernel中,一般会使用该线程自己的寄存器。

如果是小的数组,并且满足每个线程都独立的情况下,如果在编译时能确定下标的访问规律,那么编译器会选择使用多个寄存器来实现,不过编译时不能确定,则会放入local memory中。

以及,如果您寄存器使用量过多,编译器会自动评估并将一部分变量交换到local memory中。

大致如此,供您参考。

LZ您好:

具体一个线程使用多少寄存器和您的具体实现以及编译器的优化有关,128个线程的block,如果用满各个架构下的单thread的寄存器最大可用量(注意,这个值是随架构不同而不同的),在早期sm1.0和1.1的卡上会超过单个SM的寄存器资源,其他架构上均总是可以执行的。

您不妨先写出来再profiling看看。

祝您好运~

谢谢版主了,这里明白了,可以不加同步了,因为这里每个线程的数据都仅与自己写入的shared memory的数据有关。

版主,如果在shared memory中每个线程只使用自己对应的一个变量,这样改为使用寄存器,我会改
例如, S_[tx]= F[tx]Sp[tx]; 就改为 float S_=0; float F=0; float Sp=0; S_= FSp; 那如果现在不是一个线程对应一个变量,而是变成了一个小矩阵,S_[2][1]=F[2][2]Sp[2][1] (这是里矩阵相乘),用shared我才能写出来,存为这样__shared__ float S_[2M];shared float F[2*(2M)];shared float Sp[2M]; 再做矩阵相乘。用线程就不会了,请版主指教一下用线程怎么写?

LZ您好:

您在kernel里面直接定义小的数组就可以了,和常规C语言中是一样一样的。

如果数组元素很少,您直接写成多个变量也是可以的。

如前所述,如果您使用小的数组,并且在编译的时候,下标访问规律是确定的,那么编译器将安排数组使用寄存器,否则将数组整个放入local memory。

大致如此,祝您编码顺利~

恩,大概明白了,我试试,谢谢版主。

不客气的,欢迎您常来~