CUDA如何使用二维数组的疑问

加速计算 CUDA
1

这段代码是网上给出的CUDA使用二维数组的例子，我有几个问题

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <stdlib.h>

#define ROWS 32
#define COLS 16

#define CHECK(res) if(res!=cudaSuccess){exit(-1);}

__global__ void Kerneltest(int **da, unsigned int rows, unsigned int cols)
{
unsigned int row = blockDim.y*blockIdx.y + threadIdx.y;
unsigned int col = blockDim.x*blockIdx.x + threadIdx.x;
if (row < rows && col < cols)
{
da[row][col] = row*cols + col;
}
}

int main(int argc, char **argv)
{
int **da = NULL;
int **ha = NULL;
int *dc = NULL;
int *hc = NULL;
cudaError_t res;
int r, c;
bool is_right=true;

res = cudaMalloc((void**)(&da), ROWS*sizeof(int*));CHECK(res)
res = cudaMalloc((void**)(&dc), ROWS*COLS*sizeof(int));CHECK(res)
ha = (int**)malloc(ROWS*sizeof(int*));
hc = (int*)malloc(ROWS*COLS*sizeof(int));

for (r = 0; r < ROWS; r++)
{
ha[r] = dc + r*COLS;
}
res = cudaMemcpy((void*)(da), (void*)(ha), ROWS*sizeof(int*), cudaMemcpyHostToDevice);CHECK(res)
dim3 dimBlock(16,16);
dim3 dimGrid((COLS+dimBlock.x-1)/(dimBlock.x), (ROWS+dimBlock.y-1)/(dimBlock.y));
Kerneltest<<<dimGrid, dimBlock>>>(da, ROWS, COLS);
res = cudaMemcpy((void*)(hc), (void*)(dc), ROWS*COLS*sizeof(int), cudaMemcpyDeviceToHost);CHECK(res)

for (r = 0; r < ROWS; r++)
{
for (c = 0; c < COLS; c++)
{
printf("%4d ", hc[r*COLS+c]);
if (hc[r*COLS+c] != (r*COLS+c))
{
is_right = false;
}
}
printf("\n");
}
printf("the result is %s!\n", is_right? "right":"false");

cudaFree((void*)da);
cudaFree((void*)dc);
free(ha);
free(hc);
getchar();
return 0;
}

for (r = 0; r < ROWS; r++)
{
ha[r] = dc + rCOLS;
}
这段代码不理解，ha[r]不是指向主机端内存的行的首地址吗？dc不是指向设备端显存的地址吗？为什么可以ha[r]=dc+rCOLS;?

另外，主机端二维数组可以这样动态分配，

int r, c;
int **arr = (int**)malloc(ROWS*sizeof(int*));
int *data = (int*)malloc(COLS*ROWS*sizeof(int));
for (r = 0; r < ROWS; r++)
{
arr[r] = data + r*COLS;
}

free(arr);
free(data);

是不是可以这样理解：data 是整个二维数组的首地址，为data分配了COLSROWSsizeof(int)的空间；arr[r]是ROWS个行指针，它指向的首地址是data + r*COLS，分配的空间大小为COLS。我对这段代码的理解也不太确定，我printf数组data[i*COLS+j]和arr[i][j]的值是一样的。
希望大家帮忙详细解释一下~

2

您好，让您久等了。
大致查看了您贴的代码，就两个问题答复如下：

1：“ha[r]不是指向主机端内存的行的首地址吗？dc不是指向设备端显存的地址吗？为什么可以ha[r]=dc+r*COLS;?”根据您的代码，da,dc,ha,hc都是存放在主机端的指针变量，因此主机端代码访问这个变量是没有问题的。其中，da和dc使用cuda 的API申请了位于显存上的global memory的存储空间，并保存了该空间的地址；ha和hc则是申请了host端内存空间，并保存了地址。

换句话说，dc自身是在host端保存着的，dc的内容是device端的某个地址，这个地址在host端是不能拿来寻址使用的，但是dc作为host端的一个变量，我们可以随意操作。我们将dc所保存的device端空间的首地址拿出来，加上固定的偏移量COLS的倍数，得到逻辑上每一行的首地址，并将这个首地址的数值保存在ha里面，此时，ha依然是保存在host端的，但是其内容则是device端二维数组每一行的首地址，ha是一个保存了device端地址的指针数组。这里，ha和dc是类似的，他们都保存在host端，内容都是device端的地址数据。

之后，在40行的地方，使用cudaMemcpy将ha的内容复制给da。这样da这个变量本身在host端，da通过cudaMalloc申请到了device上的空间，这是一个指针数组，da变量保存了数组的首地址。cudaMemcpy之后，device上的da指向的这个指针数组得到了ha数组的内容，也就是得到了dc数组各个行的首地址。

同时还应该注意到，kernel拿到da之后，da指向的是device端global memory里面的地址，device端访问是没问题的；da是一个指针数组，拿到da的元素以后，是一个指针，这个指针依然指向device端的地址（实际上是dc和dc添加偏移量之后的各个行地址），所以还可以继续访问。这就是原代码中17行的使用方法。

总之，您需要搞清楚指针变量自身是存放于host端的，其指向的地址是device端的地址。然后将一个二级指针作为参数传给kernel函数，通过该二级指针/指针数组自己手工维护了对二维数组的寻址。

2：您的理解基本是正确的，data是实际分配保存数据的二维数组的首地址，申请的空间大小是COLSROWSsizeof(int)；arr是一个二级指针/指针数组，它的每个元素保存了data每行的地址，arr[0]作为arr的第一个元素，指向data的首地址，对应于r==0的情况。arr申请的总的空间大小为ROWSsizeof(int)。

但是，需要注意的是，arr自身作为指针的话，是指向保存自身这个指针数组的存储空间的首地址的；而arr作为指针数组，其元素依然是指针，分别指向data的各个行的首地址。

同时，还需要说明的一点在于，这里指针访问其实并不涉及空间分配，data的总的分配空间是COLSROWSsizeof(int)，这个是已经分配好的，按照我们的操作，每行逻辑上享有的空间是COLS*sizeof(int)字节。（写超了的话，会写到逻辑上的下一行，或者写越界到数组以外的地方，引发未知问题）
同时说一下，这个做法就是通过一个额外的指针数组手工维护二维数组寻址的做法。

两个问题回答如上，供您参考。

欢迎莅临cudazone，祝您编码愉快~

3

补充：

前文说了您给出的代码实际上是手工维护二维数组寻址的做法。那么这个方法有神马好处和缺点呢？

一般认为好处是使用形式直观好用，并且指针数组里面各个指针实际上可以指向各个不同的不连续的存储空间。比如我们malloc了data_a和data_b，这两个数组存储是不连续的，可以让arr[0]~arr[n-1]指向data_a各行，arr[n]~arr[m-1]指向data_b各行，然后arr访问，看上去就像一个连续的空间一样。

而此方法的缺点在于，访问得到一次数值，实际上需要访问两次，先访问指针数组，得到行地址，然后再访问一次，得到最终的数值。这样在访存上有一定代价。

其实也可以简单地使用一维的数组，对于所有需要访问arr[b][a]的地方，都直接改写为data[b*COLS+a]即可，或者您也可以写一个__device__函数完成此类转换，或者也可以用宏。

大致如上，祝您编码愉快~

4

非常感谢！更系统的理解了指针。是这样的，我想在要解决这么一个问题C100 2的排列组合问题，那么就有100*（100-1）/2=4950种情况，我要把这么多情况放到一个数组里面，4950是行的数量，每一行假设有200个元素，也就是200列，那么我就写了如下的循环

	for(int i=0;i<XCOL-1;i++)
	{
		for(int j=i+1;j<XCOL;j++)
		{
			for(int k=0;k<XROW;k++)
			{
				EpiXMatrix[((2*XCOL*i-i*i-i)/2)+(j-i-1)][k]=TranXMatrix[i][k]*TranXMatrix[j][k];
			}
		}
	}

因为这4950个行的顺序是这样的，先是1和2-100,然后是2和3-100,…,最后是99和100，是不是这种情况用二维数组好表示一些？不知道我的描述你理解什么意思了没有。

5

十分抱歉，未能理解您的叙述和代码实现意图，无法建议。

请看明白的网友补充。

6

十分抱歉，未能理解您的叙述和代码实现意图，无法建议。

请看明白的网友补充。

7

十分抱歉，未能理解您的叙述和代码实现意图，无法建议。

请看明白的网友补充。

8

手机发帖连击了，见谅。

9

没有关系，我描述的有些问题~

10

今天电脑上网，我来大致猜测一下：
1：既然用C100 2表示，那么实际上是一个组合问题而不是排列问题，排列问题一般用P表示。

2：C100 2=4950没错，问题是您如何将这个问题放到数组里面，以及您的问题是否使用数组最合适。由于您没有详细说明，这些问题，不可而知。比如您在文中提到，使用一个4950行的数组来实现，那么应该是每种组合使用一行，那么这一行都放些神马东西呢？您没说，只是说假设有200元素，请问这些元素是神马？和前面的组合有神马关系？

3：您给出了一个三重循环的代码，没有解释实现意图。如果您的k循环完整执行一次是一行的话，那么j循环完整执行一次就是一个二维的目标了，以及i再循环，就成三维的了。而您在文中一直说的是二维情况，那您是使用重复赋值么？您在k循环体中EpiXMatrix的第一个元素的位置那样一大串运算是神马意图呢？

4：后面您又开始叙述您的4950行了，按照您的叙述，您的每行的元素数量似乎是不相等的。以及您一直没有解释您的二维数组里面“200元素”是神马东西。

以上是脑力所及的所有猜测，但是基本上无法搞清楚您的具体问题，所以目前无法给您任何建议。

祝您2013年元旦快乐！

11

好久没有上论坛了，明天要做一个很重要的演讲，这些天一直忙着这个演讲，明天之后继续好好学习CUDA~

12

欢迎随时归来~

13

还是有一点不解，res = cudaMemcpy((void*)(da), (void*)(ha), ROWSsizeof(int), cudaMemcpyHostToDevice);为什么不直接写cudaMemcpy(da, ha, ROWSsizeof(int), cudaMemcpyHostToDevice);，而要在前面加上（void*）呢？另外，如果我定义一个float **da的话，那又该怎么进行复制呢？

14

cudaMemcpy()函数的前两个参数的类型根据定义是void类型的，所以可以对其他类型的指针强制转换(void)(da)。同时，如果不强制转换，会自动转换的，所以第二种写法也是可以的。

如果您定义了一个float **da，这是一个二级指针，如果它指向的那个指针已经分配好了空间什么的，可以用它指向的那个指针作为cudaMalloc()的指针参数。

---------------------------------修改错误的分割线----------------------------------

13#询问的是“如果我定义一个float **da的话，那又该怎么进行复制呢？”
上面在回答的时候，一时笔误，将cudaMemcpy()误写为cudaMalloc()（已经用红色标出），深表歉意。

cudaMemcpy()的话，需要用da指向的指针作为参数，也就是*da。

如果是cudaMalloc()，那么假定有float *dda=NULL；float **da=&dda；此时已经用二级指针da保存了dda的地址，那么可以直接写为 cudaMalloc（da，…）；

15

。。。上述程序中，da也是一个二级指针啊？
我用以下方式写，是否正确呢？
float **dp, **du, **p, **u;
float *dc, *dd, *c, *d;

p = (float **)malloc(N*sizeof(float*));  
u = (float **)malloc(N*sizeof(float*));
c = (float *)malloc(N*N*sizeof(float)); 
d = (float *)malloc(N*N*sizeof(float)); 

cudaMalloc((void**)(&dp), N*sizeof(float*));
cudaMalloc((void**)(&du), N*sizeof(float*));
cudaMalloc((void**)(&dc), N*sizeof(float));
cudaMalloc((void**)(&dd), N*sizeof(float));

cudaMemcpy((void*)(dp), (void*)(p), Nsizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy((void*)(du), (void*)(u), Nsizeof(float), cudaMemcpyHostToDevice);

kernel<<<>>>();

cudaMemcpy((void*)(c), (void*)(dc), NNsizeof(float),cudaMemcpyDeviceToHost);
cudaMemcpy((void*)(d), (void*)(dd), NNsizeof(float),cudaMemcpyDeviceToHost);

cudaFree((void*)dp);
cudaFree((void*)du);
cudaFree((void*)dc);
cudaFree((void*)dd);
free(p);
free(u);
free(c);
free(d);

16

这样写了之后总是有以下错误：

17

虽然此贴是去年12月份的，不该由我回复，但是作为ICE的朋友，我需要提示楼主一句：

你的图和你的代码无法对应，缺乏很多信息。建议你补齐带行号的代码，然后重新询问ICE. 否则这是对ICE的不尊重。

18

你需要知道一些基本的，例如int **pp的pp是一个指向int *p的指针，而p则是指向int的指针；
或者例如int (**qq)[20];的qq是一个指向int (*q)[20];的指针，而q则是指向一个int [20];的数组的首地址的指针之类的常识。
在你有了这样的常识的情况下，你需要直接提问：你试图“如何用”，但是“结果不对”，提出“为何这么用不对”。

而不是，你随意写出一个float **p;乱用。 然后各大版主轮番上阵告诉你，你可能哪里乱用了，以及为何不能这样乱用的地方。

这地儿是版主回答问题的。但版主无义务去猜测你的问题是什么，然后根据猜测回答。这颠覆你一个提问者的本分。

请支持版主的工作，下次直接问你想要的东西，哪怕再初级，版主们也会全力以赴的。谢谢。

19

实在不好意思，以后知道该怎么提问了！:slight_smile:

20

您客气了。啄木鸟很好看！谢谢!