1、假设数组u[1024][1024].
int iw = blockIdx.xblockDim.x + threadIdx.x;
int ix = blockIdx.yblockDim.y + threadIdx.y;
对u[iw][ix]的操作,是否满足合并访问呢?
2、对于二维线程,在warp内是怎么组织的呢?也就是在sm上,block内的二维线程是怎么被调度执行的。一直觉得自己不太清楚二维线程的组织,所以一般不轻易用,除非很简单的应用。
3、对于纹理操作,在deviceQuery中,有Max Texture Dimension Size(x,y,z) 1D=(65525),2D=(65535,65535),3D=(2048,2048,2048)。对于一维纹理,绑定的最大长度为65525?
先谢过版主热情的回复。
楼主您好,
(1)是否能合并访问要看warp整体,简单的说,要看warp中临近的2个线程的地址关系。而这个要具体看您的启动形状配置,而您没有给出您的线程形状,所以暂时无法判断。
但是特别的说,如果您的形状是(1,N,1) x (1,256,1)这种的,显然您的操作是可以合并的。虽然您用的都是3个内置变量的.y分量。
(2) block内的线程是按照先x再y再z的顺序组合成1维的线程,然后连续32个组合成warp然后执行的。
例如一个形状为(5,6,7)的block, 里面编号为(2,3,4)的线程实际上是1维化后的线程137(4 * 30 + 3 * 5 + 2)。
(3)对于一维纹理,
65535 * sizeof(纹元大小)才是当你使用cudaArray的时候,所能分配和绑定的最大
大小。(也就是64K个元素,而不是64KB字节)
如果你不使用cuda arrray, 而是使用普通的显存(例如有的时候你需要对一段缓冲区简单的利用上texture cache/read only cache), 那么此时你的1维纹理,将能使用相当大的显存作为后备:
此时将变成最多能分配128M个元素。
请注意前者最多可以分配/绑定64K个元素,最大一般可能是16B(int4/float4之类) * 64K = 1MB字节。
如果不够用,可以用后者,128M * 16B = 2GB,足够了。(前提是贵卡上能成功分配出2GB)
对于1D的纹理,用cuda array和线性内存做后备存储,性能上应该无区别的。建议不够就用后者。
感谢来访,周末愉快!