动态分配的二维数组

版主，您好：我有一个程序实现过程大概这样：
首先，在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呢？

LZ您好：

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

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

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

http://cudazone.nvidia.cn/forum/ ... &extra=page%3D1

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

ice 发表于 2013-11-5 22:12
LZ您好：

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

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

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

#define M 2000
#define N 3000

__global__ void matrixAdd(float *A, float *B, float *C)
{
        int i=threadIdx.x+blockIdx.x*32;
        int j=threadIdx.y+blockIdx.y*32;
        C[j*N+i]=A[j*N+i]+B[j*N+i];

}
int main()
{

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

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

        cudaMemcpy(A, a, sizeof(float)*M*N, cudaMemcpyHostToDevice);
        cudaMemcpy(B, b, sizeof(float)*M*N, 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)*M*N, cudaMemcpyDeviceToHost);
        printf("done!\n");
        for(int r=0;r<M;r++){
                for(int co=0;co<N;co++){
                        if (c[r*N+co]!=a[r*N+co]+b[r*N+co])
                        {        printf("wrong!\n");
                        }

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

        return 0;
}

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

xlwhlb 发表于 2013-11-5 22:46
哦 我发现是在kernel函数里面少一个线程的判断，导致的错误，是这样吧

LZ您好：

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

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

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

大致如此，祝您好运~

ice 发表于 2013-11-5 23:01
LZ您好：

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

感谢版主回复

1. cudaError_t test_kernel;
   
2. MatAdd<<<>>>();
   
3. test_kernel=cudaDeviceSynchronize();
   
4. 
   

复制代码

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

xlwhlb 发表于 2013-11-6 10:34
感谢版主回复是这样加吧应该，但是是显示(null)，不是success;
还有前面看到您回复过帖子，说cuda5.0，计 ...

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计算能力的编译参数，请您去掉或者修改。

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

ice 发表于 2013-11-6 11:01
LZ您好:

1：您可以如下检查kernel的执行情况：

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

xlwhlb 发表于 2013-11-6 11:21
您说的第一个问题解决了
但是怎么修改编译参数？我用的Vs2010，都是默认安装的
多谢版主 ...

LZ您好：

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

祝您好运~

ice 发表于 2013-11-6 11:31
LZ您好：

请右键单击项目名称选择属性——CDUA C/C++——寻找GPU Architecture选项并修改，修改时，请不 ...

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