关于CUDA的二维数组

直接上代码,最奇怪的几点就是貌似

1.对设备内存的操作-指定二维指针指向的一维指针,需要在设备函数中做,在主机中做会爆炸--

2.貌似设备的形参就是传递了引用。,而不需要加&了,非常奇怪,加了&反而会错误。但是在主机代码中,是需要加& 的。我只能瞎猜,CUDA的函数里面不允许引用了。。补充一个PPT

#include <cuda_runtime.h>
#include <stdio.h>

void read_error(cudaError_t ret,int count){
    if (ret != cudaSuccess)
	{
        printf("------------%d--------------\n",count);
    printf("%s\n", cudaGetErrorString(ret));
	}
}

__global__ void matrix_add(int ** A,int ** B,int ** C,int n,int m){
    int x = blockIdx.x*blockDim.x+threadIdx.x;
    int y = blockIdx.y*blockDim.y+threadIdx.y;

    if(x<n && y<m){
        C[y][x] = A[y][x] + B[y][x];
    }
}

int divup(int a,int b){
    if(a%b==0){
        return a/b;
    }
    else{
        return a/b+1;
    }
}

void init_data(int * data,int elem){
    for(int i=0;i<elem;i++){
        data[i] = rand()%65535;
    }
}

__global__ void init_matrix(int **A, int *A_data, int n, int m){
    int tid = blockDim.x * blockIdx.x + threadIdx.x;
    if (tid<n){
        A[tid] = &A_data[tid*m];
    }
}

void Memory_host(int ** &A,int * &A_data,int n,int m){
    A = (int **)malloc(sizeof(int *)*n);
    A_data = (int *)malloc(sizeof(int)*n*m);

    for(int i=0;i<n;i++){
        A[i] = &A_data[i*n];
    }
}

void Memory_device(int ** &A,int * &A_data,int n,int m){
    cudaError_t ret = cudaMalloc((void **)&A_data,sizeof(int)*n*m);
    read_error(ret,999);
    cudaMalloc((void **)&A,sizeof(int*)*n);

    init_matrix<<<1,n>>>(A,A_data,n,m);
}


int main(){
    int m = 6,n=6,w=2,h=2;
    dim3 block(w,h,1);
    dim3 grid(divup(m,w),divup(n,h),1);
    cudaError_t ret;

    int *F;
    ret = cudaMalloc((void **)&F,sizeof(int)*n*m);
    // read_error(ret);

    int **A,*A_data;
    Memory_host(A,A_data,n,m);
    init_data(A_data,n*m);

    int **d_A,*d_A_data;
    Memory_device(d_A,d_A_data,n,m);

    ret = cudaMemcpy(d_A_data,A_data,sizeof(int)*n*m,cudaMemcpyHostToDevice);

    read_error(ret,2);

        int **B,*B_data;
        Memory_host(B,B_data,n,m);
        init_data(B_data,n*m);

        int **d_B,*d_B_data;
        Memory_device(d_B,d_B_data,n,m);
        cudaMemcpy(d_B_data,B_data,sizeof(int)*n*m,cudaMemcpyHostToDevice);

        int **C,*C_data;
        Memory_host(C,C_data,n,m);

        int **d_C,*d_C_data;
        Memory_device(d_C,d_C_data,n,m);

        matrix_add<<<grid,block>>>(d_A,d_B,d_C,n,m);

        cudaMemcpy(C_data,d_C_data,sizeof(int)*n*m,cudaMemcpyDeviceToHost);

        // printf("%d %d\n",n,m);
        for(int i=0;i<n;i++){
            for(int j=0;j<m;j++){
                printf("%d ",A[i][j]);
            }
            printf("\n");
        }
        printf("\n\n");
        for(int i=0;i<n;i++){
            for(int j=0;j<m;j++){
                printf("%d ",B[i][j]);
            }
            printf("\n");
        }
        printf("\n\n");

        for(int i=0;i<n;i++){
            for(int j=0;j<m;j++){
                printf("%d ",C[i][j]);
            }
            printf("\n");
        }
        
        free(A_data);
        cudaFree(d_A_data);
    cudaDeviceReset();
}

 

### 如何使用 `cudaMemcpy` 处理二维数组CUDA 编程中,处理二维数组通常涉及内存对齐以及数据传输的方式。为了高效地管理 GPU 上的二维数组存储结构,可以利用 `cudaMallocPitch` 和 `cudaMemcpy2D` 函数来完成这些操作。 #### 使用 `cudaMallocPitch` 分配设备端内存 当分配二维数组时,推荐使用 `cudaMallocPitch` 来获取经过优化的行间距(pitch)。这种方式能够确保每一行的数据按照硬件友好的方式进行对齐,从而提升访存性能[^2]。 以下是通过 `cudaMallocPitch` 分配二维数组的一个示例: ```cpp size_t pitch; float* d_A; // 假设宽度为 width,高度为 height int width = 1024; // 列数 int height = 768; // 行数 // 调用 cudaMallocPitch 进行内存分配 cudaMallocPitch((void**)&d_A, &pitch, width * sizeof(float), height); ``` 在这里,变量 `pitch` 是实际分配给每行字节数,可能大于理论上的 `width * sizeof(float)`,这是由于硬件对齐需求所致。 #### 数据复制到设备上 一旦完成了设备端内存分配,就可以借助 `cudaMemcpy2D` 将主机中的二维数组拷贝至设备端。假设主机上有如下定义的二维数组: ```cpp float h_A[height][width]; for(int i=0;i<height;i++) { for(int j=0;j<width;j++) { h_A[i][j] = static_cast<float>(i+j); // 初始化一些值 } } ``` 接着调用 `cudaMemcpy2D` 完成从主机向设备的数据传递过程: ```cpp cudaMemcpy2D(d_A, pitch, h_A, width * sizeof(float), width * sizeof(float), height, cudaMemcpyHostToDevice); ``` 上述代码片段表明了如何将主机上的二维数组逐行传送到具有特定 pitch 的设备端缓冲区中去。 #### 设备端访问调整后的索引方法 需要注意的是,在设备端编程时,考虑到可能存在填充的情况,因此对于原始逻辑地址 `(rowIndex,colIndex)` 需要转换成为物理偏移量形式才能正确读写元素。具体做法如下所示[^3]: ```cpp __global__ void kernelExample(float* A, size_t pitch){ int colIdx = blockIdx.x * blockDim.x + threadIdx.x; int rowIdx = blockIdx.y * blockDim.y + threadIdx.y; if(colIdx < WIDTH && rowIdx < HEIGHT){ float element = *( (float*)((char*)A + rowIdx*pitch) + colIdx ); // 对element执行某些计算... } } ``` 这里的关键在于理解 `pitch` 实际代表的是以字节计的一整行长度,而不仅仅是以浮点型数量度量的结果。所以在定位某个具体的行列组合对应的数值前,先将其转化为字符指针再加回原基础类型指针即可实现精确寻址。 #### 总结说明 综上所述,采用 `cudaMallocPitch` 及其配套工具链能有效简化复杂形状张量的操作流程,并且兼顾底层架构特性带来的潜在优势。这不仅限于简单的二维情况扩展到更高维度同样适用[^1]。 ---
评论 2
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值