各位版大，现在CUDA的kernal函数还不能用静态变量吗？还有就是不同的Block除了通过kernal函数同步，还有没有其他的方法呢？
楼主您好！
（1）还不能使用静态局部变量的。
（2）还不能在grid内执行blocks们同步的。
（3）关于您询问的实现方式，我还不知道如何能实现上述2者。抱歉。
以及，关于变通的“静态局部变量”实现。您看这样行不：
（1）写成__device__变量，作为数组，以便让每个线程都能找到自己的副本。
（2）将原本的kernel的static type variable = ****; 和后续的使用代码，拆分成2个部分。一个kernel专门进行初始化计算。一个kernel则后续使用。
（3）小心的注意自己的多个此kernel的同时执行问题（例如将此kernel的多次启动总是串行的）。
以及，如果可以，建议不要采纳此建议。而是修改算法。去除对“静态局部变量的使用”。
一般认为，使用这玩意，不是很好的编码习惯。一些语言，已经去掉了对此的支持了（以便提高可读性和减少BUG)，例如C#.
谢谢千军版主，您已经解答了我的问题了，用__device__ 声名变量，实现一个初始化kernal，这就相当于静态变量。不过这样一来全局变量就很多了，显得有点小乱，呵呵。
你错了。不仅仅是这样。还需要注意维护多个线程的多个副本（相当于TLS了），所以较为麻烦。建议改用其他方式，不用这玩意。你说呢？
的确不好啊，折腾了我一晚上，麻烦老大看看下面的小程序有什么问题：
device int *dev_a;
device int *dev_b;
device int *dev_c;
global void Add(){
int i = threadIdx.x;
dev_c[i]=dev_a[i]+dev_b[i];
}
int main(){
int a[30];
int b[30];
int *c;
c=(int *)malloc(30*30*sizeof(int ));
for (int i=0;i<30;i++){
a[i]=1;
b[i]=1;
}
int error0 = cudaMalloc((void **)&dev_a,30*sizeof(int));
int error1 = cudaMalloc((void **)&dev_b,30*sizeof(int));
int error2 = cudaMalloc((void **)&dev_c,30*sizeof(int));
int error3 = cudaMemcpy(dev_a,a,30*sizeof(int),cudaMemcpyHostToDevice);
int error4 = cudaMemcpy(dev_b,b,30*sizeof(int),cudaMemcpyHostToDevice);
Add<<<1,30>>>();//dev_a,dev_b,dev_c);
int error5 = cudaMemcpy(c,dev_c,30*sizeof(int),cudaMemcpyDeviceToHost);
return 0；
}
其中error5返回值为30，,查文档说是进程已关闭，不明白什么意思。
楼主您在贵贴里犯了小错误。我为您指出：
（1）您不能在host code里直接使用__device__变量(您在cudaMalloc中直接使用了)。
（2）您不能直接对2个指针进行加法。
（3）因为您在(1)中的错误，导致您的dev_a,dev_b, dev_c均没有被初始化，以及kernel会失败。
建议：
（1）您可以在host上定义一个指针变量，然后用cudaMalloc分配缓冲区并使该指针指向该缓冲区。最后通过cudaMemcpyToSymbol将指针的值传到__device__的对应的指针变量上。
（2）您可以使用：dev_c[ i ] = dev_a[ i ] + dev_b[ i ];
（3）修正(1)后。您可以直接在cudaMemcpy的参数里，使用在(1)中定义的，3个cudaMalloc函数返回过的指针。
以及，可能有其他问题，但请先修正这3点。
以及，新问题请新发主题贴。不要和其他帖子混在一起。