线程分配问题

请教:对于1024*1024二维数组,应该如何指定行和列的索引,使其能够依次访问数组的每一个元素?

LZ能详细叙述下您的问题所在么?单纯访问这样一个二维数组是很容易的事。

好的。问题是这样的:
对于1024*1024的一个二维数组,我要从row=2,col=2开始,依次对每一个数组元素进行相应的计算,想请教一下如何进行线程的分配,得到的行和列的索引能满足要求?

LZ您好,您的叙述中有“依次”一个词,如果您需要访问这些数据并进行计算的时候,维持某种顺序性,那么请考虑使用CPU计算,GPU上多个线程往往是不保证先后顺序的。

如果您只是需要对这些元素每个都计算一遍,元素间也没有什么顺序逻辑依赖,那么您可以安排多个线程来完成您的计算。
此时

1:您只需要考虑使用多少线程,完成这些数据的处理,即可,并无特定的需要遵循的限制。您只需要安排好线程编号和寻址之间的关系即可。

2:如果您需要一个示例性的安排,那么,简单地取每个block 1024个线程,总共有1024个blocks。假定您的数组是行优先存储的,那么安排每个block计算一行的数据,blockIdx.x*1024+threadIdx.x即是每个线程需要计算的对应数据地址。并在kernel最前面加上判断,让row=0,1,col=0,1(按照c语言下标习惯,如果您文中本意是从1开始的,请做适当调整)的对应线程直接返回不干活即可。

祝您编码顺利~

多谢版主,这个问题现在搞明白了!

[

不客气的,积跬步以致千里,祝您取得更大的进步。

还是想请教一个问题:对于1536*1536的一个二维网格,我进行如下分配:
dim3 dimGrid(1536/16,1536/16);
dim3 dimBlock(16,16,1);
行和列的索引分别为:
int row = blockIdx.x * blockDim.x + threadIdx.x;
int col = blockIdx.y * blockDim.y + threadIdx.y;
这样在计算的时候能访问到每一个网格吗?(不考虑访问顺序,元素间也没有什么顺序逻辑依赖)

[

您好,目测是可以的,您不妨一试。

祝您节日愉快!

苦逼啊,劳动节放假还得学习……我试了一下,总感觉当 blockIdx.x=0, blockIdx.y=0时,没有行和列索引号分别为threadIdx.x和threadIdx.y(除了二者都等于0时)的线程在参与计算。
额。。这个有点绕,不知在我上述这种情况下,一般怎么进行grid和block的划分会合适呢?

看不懂楼主的问题。楼主真心明白自己在说什么?

任何一个block, 只要它存在,总是会包含线程的。
只要它包含>1个线程,那么总有一个线程索引分量不是0的。

所以楼主说,特定的block, 只有一个线程(threadIdx.x/y都是0的那个)参与了计算,在楼主上文给出(16,16)的block shape的情况下,我认为这不可能。

除非您的硬件突然出现了问题,导致您的block(包含16*16=256个线程)中的任何线程,都得到自己的threadIdx.x = 0和自己的threadIdx.y = 0。

但此概率小于您命中明晚的双色球的概率。

请三思以确定,您的发言没有背叛您的思想。谢谢合作。

额。看来是我的表述很有问题。我再试着阐述一下:
对于一个1536*1536的网格,应如何设定 执行配置 ,使得row和col得到的索引能够对每一个网格都进行一次计算?

您已经在7#进行了正确的尝试:

dim3 grid_shape(96,96);
dim3 block_shape(16,16);

然后用
int col = blockIdx.x * blockDim.x + threadIdx.x;
int row = blockIdx.y * blockDim.y + threadIdx.y;
这是可以对应的。但是:
(1)您在上文给予了否认。说只有神马0号线程有用,能参与计算。
(2)然后我指出block里的256的线程都好好的参与了计算。
(3)然后您重新表示不知道如何设置形状了。
(4)然后本文我表示上文的形状就可以。
(5)然后根据您的上下文,您应该继续表示:(1)“只有0号线程才能计算”了。
(6)然后我应该继续表达(2)
(7)然后您应该继续表达(3)
(8)然后我应该继续表达(4)

死循环中。

除非一个人作出改变。

多谢版主耐心的讲解,由于没有翻页导致没有看到版主11#的回复,然后就有了12#的又一次提问,非常抱歉,这个问题现在清楚了!还有一点,看编程手册中Block和Thread是按列存储的,在cuda中二维数组是按行还是按列存储的?这个还真有点绕……

没关系,没看到现在看到就行。

block安排和线程安排不是什么“按列”安排的,而是这样的:
一个block里,按照threadIdx.x/y/z的顺序,组成warp(以及里面的threads)。
一个grid里,按照blockIdx.x/y/z的顺序,组成block(以及里面的warps和warps里的threads)。

需要注意的是,warp不能跨block边界组合,32个block, 分别有31个线程,那么将每个block组成1个warp, 一共32个warps; 而不能连续的2个blocks合并组合成31个warps。

关于您的第二个问题,在“在cuda中二维数组是按行还是按列存储的?”,其实是这样的,
在CUDA C/C++中,有type array;是按低维度优先存储的。

而不能说“列”或者“行”,因为这取决与你的应用。
当你用array[y]的形式访问的时候,低维度优先存储的该数组则在你的理解中是按“行优先”存储的。
而当你用array[y]的形式访问的时候,低维度优先存储的该数组则在你的理解中是按“列优先”存储的。

所以第二个问题的答案是,语言/编译器只负责保证低维度优先。而在应用中是判断“行存储”还是“列存储”则要看该应用是如何用的,也就是取决于人,而和语言、编译器无关了。

再次感谢版主们多次详细的解答,的确也从中学习了不少!