关于constant数组赋值

加速计算 CUDA
1

斑竹好，请问对于constant数组赋值，下面这样是不是不允许？
constant double sn[64];

double h_sn[64];
for (int i = 0;i<64;i++)
{
h_sn[i] = (Xmin-Lsar/2)/V + ((Xmax+Lsar/2)/V - (Xmin-Lsar/2)/V) * i / (64- 1);
}

HANDLE_ERROR(cudaMemcpyToSymbol(sn,h_sn,sizeof(h_sn)));

对于constant数组赋值，在主机上的数组h_sn必须在定义是初始化，不可以如上这样初始化

2

LZ您好：

可以使用cudaMemcpyToSymbol()进行初始化的，不知您为何觉得不可以呢？以下是实测结果，是可以的。
（虽然和您的写法略有不同，但是您的写法也是正确的。）

附一张测试截图：

[attach]3253[/attach]

祝您编码顺利~

3

谢谢，斑竹解答。
另外就是，刚发现一个问题，同样的程序，在同样的电脑上运行，几次运行的结果不同，kernel中没有使用到shared memory,不知如何解决

4

代码如下
global void product(doubleDslow,doubleR,double* tau,double *phase,gpu_Complex *Srnm,double L,double lambda)
{
double Kr = Br/Tr;
int m = threadIdx.x + threadIdx.y * blockDim.x;
int n = blockIdx.y ;

int offset = m + n * blockDim.x*blockDim.y;

for (int k = 0;k<3;k++)
{
	Dslow[n] = sn[n]*V - Ptarget[k*3+0];
	tau[n]= 2*sn[n]/C;
	phase[offset] = PI*Kr*(tm[m]-tau[n])*(tm[m]-tau[n]);
	if ( (abs(tm[m]-tau[n])<Tr/2 )&&(abs(Dslow[n])<L/2)) 
	{
	 	Srnm[offset].r +=  cos(phase[offset]);
	 	Srnm[offset].i +=  sin(phase[offset]);

	 }
	 else
	 {
		Srnm[offset].r +=  0.0;
		Srnm[offset].i +=  0.0;
	 }
}

}

其中，sn,tm Ptarget都是constant数组

5

LZ您好：

您的问题较大可能是您的代码有问题；较小可能是您使用了某种随机数生成函数影响了结果但是您却不知道。

以及，使用shared memory是产生bug的一个途径，但是不是唯一的途径。

建议您另开一贴，详细叙述您的问题，这样才能大家一起帮您寻找问题。而像您当前这样询问，没有任何细节，是无法回答的，请见谅。

祝您好运~

6

LZ您好：

刚才在回复的时候，被您插楼了，我将人肉查看一下您的4#的代码，看看能否看出问题。

请稍后。

7

不好意思，麻烦您了

8

LZ您好：

经过观察，发现这一句：“Dslow[n] = sn[n]V - Ptarget[k3+0];”可能有问题。

因为您至少整个block（如果gridDim在x方向为1的话，那么是一个block，否则是同一行的多个block，因为这里用blockIdx.y区分）的线程都是往Dslow[n]这同一个位置写入的，一个block里面的各个线程的执行进度是不一样的，有些线程可能执行到了k=2，有些k=1，有些k=0。此时他们都按自己的进度给同一个全局的Dslow[n]赋值，但是在后面判断中使用Dslow[n]的时候却无法保证是自己写入的，因此会出现每次运行结果都不一致的情况。

请您修改您的算法实现，以解决这一问题。
以及，因为是目测的，亦无法保证无其他问题。

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

9

嗯嗯。ICE说的是。同样的，你的tau也存在此问题（如果C不是常量，感谢ICE指出这点）。

如果你需要多个线程拥有不同的副本，您可以将Dslow和tau改成局部变量。
或，如果能至少保证一个block中的值都是一样的，可以考虑在对DSlow、tau的每次写入和读取后，加入同步。

感谢来访。

10

非常感谢楼主，楼主说的应该没错。
我的算法是在gridDim在x方向一个block，y方向N个block，每个block有M个thread。其中每个block里的thread都共用一个sn,Dslow,tau.所以我只为sn,Dslow,tau分配Ndouble的大小的空间。
如果要解决这个问题，是不是为每个thread分配一个sn,Dslow,tau，即大小为NM*double的空间，还是在加一个__syncthreads();确保每一个block中的thread都同步。

11

LZ您好：

9#已经指出了2种修改方法，我再稍微解释一下：
仅就逻辑正确性而言：

1：按照您现在的写法，申请的__constant__空间和global memory中的空间不变，每个线程读入对应的值以后，保留为本地的局部变量副本，然后对这个副本进行操作即可。
比如double Dslow_n=Dslow[n],然后后面都使用自身的这个Dslow_n即可。
因为直接定义在kernel中的变量是每个线程私有的，每个线程拥有一份副本，互不影响。

2：或者在保留您原有写法不变的基础上，在“Dslow[n] = sn[n]V - Ptarget[k3+0];tau[n]= 2*sn[n]/C;”这两行的前面和后面都加上__syncthreads()以同步。

以及还有其他实现方法，且上述方法不涉及优化，仅保证逻辑正确性。

大致如此，请您参照本楼的解释理解9#的建议，并修改您的代码。

祝您编码顺利~

12

灰常感谢两位版主的帮助，最近在论坛上解决了不少问题和困惑，感谢提供这样的学习和交流的环境！:smiley:

13

您客气了，感谢您的来访。

我代表ICE表示服务您是我的荣幸。