share memory使用时是不是默认都是初始化为0的?还是没有定义,要手动来初始化一下。
亲,无法保证是0的。请自己初始化一下。
还有一个问题
不同warp操作share memory时没有bank conflict。是不是可以理解为不同warp操作share memory是串行的?
是的。同一个SM里的多个warp对shared memory的访问是串行的。
您可以这样理解,一个SM里面的shared memory 一次只能为至多一个warp服务。
祝您好运!
从global memory中读取
struct A {int aa[4]; int bb[4];}
到share memory中
能不能这样写
if (tid < 8)
{
share_A.aa[tid] = global_A.aa[tid];
}
LZ您好,大致试验了一下,如此赋值是没有问题的。
需要说明:
1:我用的是结构体指针作为kernel的参数测试的,如果写成您这样,可能需要__device__声明的结构体并使用cudaMemcpyTosymbol来从主机端赋值。请您搞定相关的工作。
2:建议您将相关的代码放到.cu文件中,以免造成结构体申请空间时,对齐方式不同,造成的问题。
3:以及您的代码中tid<8应该改为tid<4,否则会写越界,这个估计是笔误。
4:以及在shared memory赋值以后,请注意加上__syncthreads();,以确保shared memory确实完成写入之后,再被其他线程使用。
大致如上,未尽事宜,请其他斑竹/网友补充。
祝您编码愉快~
关于第三点,我补充一下,楼主这样是故意的:
他有int aa[4]和int bb[4], 然后他试图访问aa[0-7]以便将aa和bb都读取。
这样是无问题的。因为int对齐到4B的没事。但对于:
struct A { int aa[5]; double bb; };然后试图访问aa[0-6]这7个4B是不行的。因为中间“可能”存在4B的padding区域。
一个可能的较好的做法是(对于这个例子)
(假设shared_A和global_A是分别指向shared memory中和global memory中struct A的指针。)
if (tid < sizeof(struct A))
{
(char *)shared_A[tid] = (char *)global_A[tid];
}
(以及,因为这里肯定是4B的整数倍大小。您还可以这样:)
if (tid < sizeof(struct A) / 4)
{
(int *)shared_A[tid] = (int *)global_A[tid];
}
都行。
请注意这个是对7# ICE大的回复的补充。
代为修正一处笔误
将文中的部分数据类型给予阐明
版主方法甚好。一会去测试下效果。
考虑到我程序本身和share memory的对齐问题。
还是用4byte的填充比较好。