在CUDA如何使用二维数组（**[M][N]）

来源：互联网发布：mac版腾讯视频设置在哪编辑：程序博客网时间：2024/05/20 18:00

总有些童鞋想知道怎么在CUDA中使用二维数组（[M][N]这种类型），其实这个很简单，不过你要完全弄明白，必须对指针，地址等概念非常清楚才行。写这篇博客解决下大家这个问题：

1、首先讲述一下在一般C语言中如何使用二维数组。

[cpp] view plaincopy
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);  

代码中的arr实个二维数组变量了，你可以在for循环之后arr[i][j]的方式使用它。

2、告诉你如何在CUDA中使用二维数组可以类比1中的方法，不过你要清楚几点，这几点在代码之后说明。

[cpp] view plaincopy
#include <stdio.h>  
#include <stdlib.h>  
#include <cuda_runtime.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;  
}  

在CUDA中使用二维数组的几点说明：

1）da是一个二维变量，一定更不可以在33行的时候把ha改成da！一定要记住显存和内存是相互独立的，主机端的程序不可以直接操作显存！必须通过CUDA 提供的API函数来操作！

2）注意在内存申请时强制类型转换(void**)（&），怎么把***的变量转成**了！！这主要是API借口决定的，最好自己显式转换格式，避免不必要的麻烦。

3）看见数据拷贝的函数了吗，类型、类型、还是类型。

4）别忘了释放内存和显存！看见没，还是类型。

5）很希望这篇博客能帮到大家，可是我真的不推荐大家在GPU上使用二维数组！真的！！为什么呢？终归是效率惹的祸！显存的访问总是慢的。二维访存，可是连续访问了两次啊。要是老这样做，不但执行效率低，而且写代码也慢。如果对内存的概念不熟悉，千万别趟这趟浑水。看懂这段代码，就当是学习一下或者理解下内存、显存与内存独立的概念和规则吧。

附上执行结果：

转自：http://blog.csdn.net/bendanban/article/details/7669624

0 0