动态分配的二维数组

加速计算 CUDA
1

版主，您好：我有一个程序实现过程大概这样：
首先，在CPU端 定义float** host_a,然后并动态分配内存(Nx,Nz取值大概1000~2000左右，)得到二维数组host_a[Nx][Nz]，
然后，将host_a数据copy到device端，利用GPU对其进行处理运算；
然后，再传回到CPU端的host_b[Nx][Nz]

但是host端动态分配二维数组可以做到，但是device端该怎么定义以及实现拷贝数据呢？

GPU端要定义的是 float** device_a 数据类型吗？那怎么cudaMemcpy呢？

2

LZ您好：

首先您这个用法不是“二维数组”，而是一个一维的数组里面元素是指针，然后每个指针指向一个申请到的一维缓冲区。一维的缓冲区之间是不连续的。

如果您确实要如此使用，那么在device端也需要类似申请，然后每个一维的缓冲区单独复制。

一般推荐的做法是申请连续的缓冲区使用，您可以参考以前的讨论帖：

http://cudazone.nvidia.cn/forum/forum.php?mod=viewthread&tid=7056&extra=page%3D1

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

3

[quote=“system”, post: 2, topic: 3278]

LZ您好：

首先您这个用法不是“二维数组”，而是一个一维的数组里面元素是指针，然后每个指针指向一个申请 …[/quote]

多谢版主回答
按照一维缓冲区的做法是不是下面这样，定义一维的，
但是运行出来，好像结果GPU结果不对，是为什么呢

include <stdio.h>
include “cuda_runtime.h”
include “device_launch_parameters.h”
include

define M 2000
define N 3000

global void matrixAdd(float A, float B, float C)
{
int i=threadIdx.x+blockIdx.x32;
int j=threadIdx.y+blockIdx.y32;
C[jN+i]=A[jN+i]+B[jN+i];

}
int main()
{

float a,b,c;
cudaMallocHost((void)&a ,NMsizeof(float));
cudaMallocHost((void**)&b ,NMsizeof(float));
cudaMallocHost((void**)&c ,NM*sizeof(float));
float *A,*B,*C;

cudaMalloc((void **) &A, sizeof(float)MN);
cudaMalloc((void **) &B, sizeof(float)MN);
cudaMalloc((void **) &C, sizeof(float)MN);
for (int j=0;j<M;j++)
{
for (int i=0;i<N;i++)
{
a[jN+i]=i+j;
b[jN+i]=i+2*j;
}
}

cudaMemcpy(A, a, sizeof(float)MN, cudaMemcpyHostToDevice);
cudaMemcpy(B, b, sizeof(float)MN, cudaMemcpyHostToDevice);

dim3 threadsPerBlock(32, 32);
dim3 dimGrid( (N + threadsPerBlock.x - 1)/threadsPerBlock.x ,(M + threadsPerBlock.y - 1)/threadsPerBlock.y ) ;
matrixAdd<<<dimGrid, threadsPerBlock>>>(A,B,C);

cudaMemcpy(c, C, sizeof(float)MN, cudaMemcpyDeviceToHost);
printf(“done!\n”);
for(int r=0;r<M;r++){
for(int co=0;co<N;co++){
if (c[rN+co]!=a[rN+co]+b[r*N+co])
{ printf(“wrong!\n”);
}

  }

}
cudaFree(A);
cudaFree(B);
cudaFree(C);

return 0;
}

4

哦 我发现是在kernel函数里面少一个线程的判断，导致的错误，是这样吧

5

LZ您好：

1：您kernel里面使用了自己维护的一维化寻址，目测是正确的。

2：您host端Grid（实际是blocks per grid）参数使用了快速向上取整，但是您的kernel内部并无判断以让多余的线程提前返回。所以您的kernel在执行中会访存越界，从而挂掉。

您可以在您的kernel后面添加cudaDeviceSynchronize(),并查看返回值，以确定您的kernel的运行情况。
或者用nsight调试您的代码。

大致如此，祝您好运~

6

感谢版主回复

cudaError_t test_kernel;
MatAdd<<<>>>();
test_kernel=cudaDeviceSynchronize();

是这样加吧应该，但是是显示b，不是success[/b];
还有前面看到您回复过帖子，说cuda5.0，计算能力2.0以上，可以在kernel函数中加printf,我的为什么会报错

7

LZ您好:

1：您可以如下检查kernel的执行情况：
cudaError_t test_kernel;
MatAdd<<<>>>();
test_kernel=cudaDeviceSynchronize();
printf(“%s \n”,cudaGetErrorString(test_kernel));

如果打印出来“no error”表示您kernel执行成功。（注意这个字符串对应的cudaError_t的枚举值为cudasuccess。）

2：请您按照2.0以上的编译参数编译您的代码。CUDA toolkit默认项目类型的编译参数中有1.0计算能力的编译参数，请您去掉或者修改。

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

8

您说的第一个问题解决了
但是怎么修改编译参数？我用的Vs2010，都是默认安装的
多谢版主

9

LZ您好：

请右键单击项目名称选择属性——CDUA C/C++——寻找GPU Architecture选项并修改，修改时，请不要继承原来的值。

祝您好运~

10

您好，没有找到GPU Archeitecture选项呢
common,device,host,command line里面都没有
还是我没有配置好？？
怎么上传不了图片啊

11

LZ您好：

刚才实机查看了一下，在：CUDAC /C++——Device——Code Generation这里，选择编辑。

如果您的GPU是计算能力2.0的（GTX 480，470，465，580，570;TESLA 2050，2070，2075等）请选为compute_20,sm_20
2.1的（其他fermi核心的geforce卡）请选为compute_20,sm_21
3.0的（一般的kepler核心的geforce卡以及tesla K10）请选为compute_30,sm_30
3.5的（TESLA K20/K20X,geforce TITAN,780,780Ti以及GK208核心的geforce 630等）请选为compute_35,sm_35

前面提到的GPU Architecture是VS 2008中的叫法，给您带来迷惑深表歉意。

祝您好运~

12

感谢版主详实的回复，而且这句话，

给您带来迷惑深表歉意。
受不起啊，再次表示感谢:)

13

不客气的，欢迎您常来论坛~