【CUDA代码实践03】m维网格n维线程块对二维矩阵的索引

为了方便下次找到文章，也方便联系我给大家提供帮助，欢迎大家点赞👍、收藏📂和关注🔔！一起讨论技术问题💻，一起学习成长📚！如果你有任何问题或想法，随时留言，我会尽快回复哦😊！

---

近年来，人工智能（AI）技术，尤其是大模型的快速发展，打开了全新的时代大门。对于想要在这个时代迅速成长并提升自身能力的个人而言，学会利用AI辅助学习已经成为一种趋势。不论是国内的文心一言、豆包，还是国外的ChatGPT、Claude，它们都能成为我们编程学习的有力助手。利用AI进行编程学习将大大提升自己的编程学习效率，这里给大家推荐一个我自己在用的集成ChatGPT和Claude的网站（国内可用，站点稳定）：传送门

抓住AI时代每一个机会，加速自己成长，提高自己的核心价值！

一、数据存储方式

• 数据在内存中是以线性、以行为主的方式存储
• 本篇文章中，16x8的二维数组，在内存中一段连续的128个地址存储该数组
  
  代码结构
  

先放上头文件及main文件
common.cuh

#ifndef COMMON_CUH
#define COMMON_CUH#include "cuda_runtime.h"
#include <stdio.h>// 声明外部函数，它们将在其他文件中实现。
// 这些函数定义了 CUDA 的网格和块结构，分别表示
// 2维网格和2维线程块、2维网格和1维线程块、1维网格和1维线程块。
extern void grid2D_block2D();
extern void grid2D_block1D();
extern void grid1D_block1D();// ErrorCheck 是一个内联函数，用于检查 CUDA 函数的返回错误码。
// 如果有错误发生，它将打印错误代码、错误名称、错误描述、文件名和行号。
// 此函数的目的是帮助调试 CUDA 错误。
inline cudaError_t ErrorCheck(cudaError_t error_code, const char* filename, int lineNumber) {if (error_code != cudaSuccess) {printf("CUDA error:\ncode=%d, name=%s, description=%s\nfile=%s,line=%d\n",error_code, cudaGetErrorName(error_code), cudaGetErrorString(error_code), filename, lineNumber);return error_code;  // 返回错误码以便调用方了解错误情况。}return error_code;  // 如果没有错误，返回相同的错误码。
}// setGPU 是一个内联函数，用于设置 GPU 设备。
// 它首先获取系统中可用的 CUDA 兼容 GPU 数量。
// 如果没有找到可用的 GPU，程序将退出，否则设置设备并显示相应信息。
inline void setGPU() {int iDeviceCount = 0;  // 存储系统中可用 GPU 的数量// 获取设备数量并检查返回的错误码。cudaError_t error = ErrorCheck(cudaGetDeviceCount(&iDeviceCount), __FILE__, __LINE__);// 如果没有 GPU 或发生错误，则终止程序。if (error != cudaSuccess || iDeviceCount == 0) {printf("No CUDA compatible GPU found\n");exit(-1);  // 返回非零值，表示错误。} else {printf("The count of GPUs is %d.\n", iDeviceCount);  // 显示找到的 GPU 数量。}// 设置设备 ID 为 0 的 GPUint iDevice = 0;error = ErrorCheck(cudaSetDevice(iDevice), __FILE__, __LINE__);if (error != cudaSuccess) {printf("cudaSetDevice failed!\n");exit(-1);  // 设置失败时终止程序。} else {printf("cudaSetDevice success!\n");  // 成功设置 GPU 后的确认信息。}
}#endif // COMMON_CUH

main.cu

#include "cuda_runtime.h"
#include <stdio.h>
#include "./common.cuh" // 包含自定义的通用 CUDA 工具，例如 setGPU 和 ErrorCheckint main() {// grid2D_block2D();  // 使用 2维网格和 2维线程块的函数，已注释掉// grid2D_block1D();  // 使用 2维网格和 1维线程块的函数，已注释掉grid1D_block1D();    // 使用 1维网格和 1维线程块的函数，执行矩阵加法return 0; // 返回 0 表示程序执行成功
}

二、二维网格二维线程块

二维网格和二维线程块对二维矩阵进行索引，每个线程可负责一个矩阵元素的计算任务

//
// Created by Administrator on 2024/10/25.
//
#include "common.cuh"// 定义一个 CUDA 内核函数 addMatrix，用于对两个矩阵进行元素逐一相加。
// A、B 是输入矩阵，C 是输出矩阵，nx 和 ny 分别是矩阵的列数和行数。
__global__ void addMatrix(int *A, int *B, int *C, const int nx, const int ny) {int ix = blockIdx.x * blockDim.x + threadIdx.x; // 确定线程在 x 方向上的索引int iy = blockIdx.y * blockDim.y + threadIdx.y; // 确定线程在 y 方向上的索引unsigned int idx = iy * nx + ix; // 计算该线程对应矩阵中的一维索引// 仅当索引在矩阵范围内时执行加法运算，以避免越界访问if (ix < nx && iy < ny) {C[idx] = A[idx] + B[idx];}
}// 定义一个函数 grid2D_block2D 来设置并调用 CUDA 内核
// 该函数配置并使用二维网格和二维块结构
void grid2D_block2D(void) {setGPU(); // 设置 GPU// 初始化矩阵大小和字节数int nx = 16; // 列数int ny = 8;  // 行数int nxy = nx * ny; // 矩阵元素总数size_t stBytesCount = nxy * sizeof(int); // 矩阵所需的总字节数// 在主机（CPU）上分配内存int *ipHost_A, *ipHost_B, *ipHost_C;ipHost_A = (int *)malloc(stBytesCount); // 矩阵 AipHost_B = (int *)malloc(stBytesCount); // 矩阵 BipHost_C = (int *)malloc(stBytesCount); // 矩阵 C// 初始化 A 和 B 的值，C 初始化为零if (ipHost_A != NULL && ipHost_B != NULL && ipHost_C != NULL) {for (int i = 0; i < nxy; i++) {ipHost_A[i] = i;        // 矩阵 A 的元素值设为 iipHost_B[i] = i + 1;    // 矩阵 B 的元素值设为 i+1}memset(ipHost_C, 0, stBytesCount); // 矩阵 C 的元素初始化为 0}else {printf("fail to malloc memory.\n");exit(-1); // 如果内存分配失败，退出程序}// 在设备（GPU）上分配内存int *ipDevice_A, *ipDevice_B, *ipDevice_C;ErrorCheck(cudaMalloc((int **)&ipDevice_A, stBytesCount), __FILE__, __LINE__);ErrorCheck(cudaMalloc((int **)&ipDevice_B, stBytesCount), __FILE__, __LINE__);ErrorCheck(cudaMalloc((int **)&ipDevice_C, stBytesCount), __FILE__, __LINE__);if (ipDevice_A != NULL && ipDevice_B != NULL && ipDevice_C != NULL) {// 将 A 和 B 从主机复制到设备ErrorCheck(cudaMemcpy(ipDevice_A, ipHost_A, stBytesCount, cudaMemcpyHostToDevice), __FILE__, __LINE__);ErrorCheck(cudaMemcpy(ipDevice_B, ipHost_B, stBytesCount, cudaMemcpyHostToDevice), __FILE__, __LINE__);ErrorCheck(cudaMemcpy(ipDevice_C, ipHost_C, stBytesCount, cudaMemcpyHostToDevice), __FILE__, __LINE__);}else {printf("fail to malloc memory.\n");free(ipHost_A);free(ipHost_B);free(ipHost_C);exit(-1); // 如果设备内存分配失败，退出程序}// 设置线程块和网格维度dim3 block(4, 4); // 定义每个块的尺寸（4x4 线程块）dim3 grid((nx + block.x - 1) / block.x, (ny + block.y - 1) / block.y); // 定义网格尺寸printf("Thread config : grid (%d, %d) block (%d, %d)\n", grid.x, grid.y, block.x, block.y);// 启动 CUDA 内核addMatrix<<<grid, block>>>(ipDevice_A, ipDevice_B, ipDevice_C, nx, ny);ErrorCheck(cudaDeviceSynchronize(), __FILE__, __LINE__);// 将结果从设备复制回主机ErrorCheck(cudaMemcpy(ipHost_C, ipDevice_C, stBytesCount, cudaMemcpyDeviceToHost), __FILE__, __LINE__);// 输出前 10 个元素的加法结果，验证计算是否正确for (int i = 0; i < 10; i++) {printf("idx=%2d\tmatrix_A:%d\tmatrix_B:%d\tresult=%d\n", i + 1, ipHost_A[i], ipHost_B[i], ipHost_C[i]);}// 释放主机和设备上的内存free(ipHost_A);free(ipHost_B);free(ipHost_C);ErrorCheck(cudaFree(ipDevice_A), __FILE__, __LINE__);ErrorCheck(cudaFree(ipDevice_B), __FILE__, __LINE__);ErrorCheck(cudaFree(ipDevice_C), __FILE__, __LINE__);ErrorCheck(cudaDeviceReset(), __FILE__, __LINE__);return;
}

三、二维网格一维线程块

二维网格和一维线程块对二维矩阵进行索引
每个线程可负责一个矩阵元素的计算任务
与二维网格二维线程块的情况极为相似

//
// Created by Administrator on 2024/10/25.
//
#include "common.cuh"// 定义一个 CUDA 内核函数 addMatrix_21D，用于执行矩阵相加操作。
// 与标准的 2维线程块不同，此内核使用 2维网格和 1维线程块配置。
// A、B 是输入矩阵，C 是输出矩阵，nx 和 ny 分别是矩阵的列数和行数。
__global__ void addMatrix_21D(int *A, int *B, int *C, const int nx, const int ny)
{int ix = blockIdx.x * blockDim.x + threadIdx.x; // 计算线程在 x 方向上的索引int iy = threadIdx.y;                           // 线程在 y 方向上的索引unsigned int idx = iy * nx + ix; // 将二维索引转换为一维索引// 仅在索引位于矩阵范围内时执行加法操作，避免越界访问if (ix < nx && iy < ny){C[idx] = A[idx] + B[idx];}
}// 定义一个函数 grid2D_block1D 来配置并调用 CUDA 内核。
// 该函数使用二维网格和一维块的配置。
void grid2D_block1D(void)
{setGPU(); // 设置 GPU 设备// 设置矩阵的尺寸int nx = 16; // 矩阵列数int ny = 8;  // 矩阵行数int nxy = nx * ny; // 矩阵总元素数size_t stBytesCount = nxy * sizeof(int); // 矩阵所需的总字节数// 在主机（CPU）上分配内存int *ipHost_A, *ipHost_B, *ipHost_C;ipHost_A = (int *)malloc(stBytesCount); // 分配矩阵 A 的内存ipHost_B = (int *)malloc(stBytesCount); // 分配矩阵 B 的内存ipHost_C = (int *)malloc(stBytesCount); // 分配结果矩阵 C 的内存// 初始化矩阵 A 和 B 的数据if(ipHost_A != NULL && ipHost_B != NULL && ipHost_C != NULL){for(int i = 0; i < nxy; i++){ipHost_A[i] = i;        // A 的每个元素为 iipHost_B[i] = i + 1;    // B 的每个元素为 i+1}} else {printf("fail to malloc memory.\n");exit(-1); // 如果内存分配失败，退出程序}// 在设备（GPU）上分配内存int *ipDevice_A, *ipDevice_B, *ipDevice_C;ErrorCheck(cudaMalloc((int **)&ipDevice_A, stBytesCount), __FILE__, __LINE__);ErrorCheck(cudaMalloc((int **)&ipDevice_B, stBytesCount), __FILE__, __LINE__);ErrorCheck(cudaMalloc((int **)&ipDevice_C, stBytesCount), __FILE__, __LINE__);if(ipDevice_A != NULL && ipDevice_B != NULL && ipDevice_C != NULL){// 将主机内存复制到设备ErrorCheck(cudaMemcpy(ipDevice_A, ipHost_A, stBytesCount, cudaMemcpyHostToDevice), __FILE__, __LINE__);ErrorCheck(cudaMemcpy(ipDevice_B, ipHost_B, stBytesCount, cudaMemcpyHostToDevice), __FILE__, __LINE__);ErrorCheck(cudaMemcpy(ipDevice_C, ipHost_C, stBytesCount, cudaMemcpyHostToDevice), __FILE__, __LINE__);} else {// 如果分配失败，释放内存并退出程序free(ipHost_A);free(ipHost_B);free(ipHost_C);exit(-1);}// 定义线程块和网格的尺寸dim3 block(4); // 每个块有 4 个线程dim3 grid((nx + block.x - 1) / block.x, (ny + block.y - 1) / block.y); // 配置网格尺寸printf("Thread config: grid (%d, %d), block(%d, %d).\n", grid.x, grid.y, block.x, block.y);// 启动 CUDA 内核addMatrix_21D<<<grid, block>>>(ipDevice_A, ipDevice_B, ipDevice_C, nx , ny);ErrorCheck(cudaDeviceSynchronize(), __FILE__, __LINE__);// 将结果从设备复制回主机ErrorCheck(cudaMemcpy(ipHost_C, ipDevice_C, stBytesCount, cudaMemcpyDeviceToHost), __FILE__, __LINE__);// 输出前 10 个元素的加法结果，进行验证for(int i = 0; i < 10; i++){printf("idx=%2d\tmatrix_A:%d\tmatrix_B:%d\tresult=%d\n", i + 1, ipHost_A[i], ipHost_B[i], ipHost_C[i]);}// 释放主机和设备上的内存free(ipHost_A);free(ipHost_B);free(ipHost_C);ErrorCheck(cudaFree(ipDevice_A), __FILE__, __LINE__);ErrorCheck(cudaFree(ipDevice_B), __FILE__, __LINE__);ErrorCheck(cudaFree(ipDevice_C), __FILE__, __LINE__);ErrorCheck(cudaDeviceReset(), __FILE__, __LINE__);return;
}

四、一维网格一维线程块

一维网格和一维线程块对二维矩阵进行索引
每个线程负责矩阵一列的运算
编写核函数时，需要使用循环

//
// Created by Administrator on 2024/10/28.
//
#include "common.cuh"// 定义一个 CUDA 内核函数 addMatrix_11D，使用 1D 网格和 1D 块来执行矩阵相加。
// A、B 是输入矩阵，C 是输出矩阵，nx 和 ny 分别是矩阵的列数和行数。
__global__ void addMatrix_11D(int *A, int *B, int *C, const int nx, const int ny)
{int ix = blockIdx.x * blockDim.x + threadIdx.x; // 计算线程在 x 方向的索引// 确保索引在矩阵范围内if (ix < nx){// 在 y 方向上循环遍历for (int iy = 0; iy < ny; iy++){unsigned int idx = iy * nx + ix; // 将二维索引转换为一维索引C[idx] = A[idx] + B[idx]; // 将 A 和 B 对应位置相加并存储在 C 中}}
}// 定义一个函数 grid1D_block1D 来配置并调用 CUDA 内核，使用 1D 网格和 1D 块
void grid1D_block1D(void)
{printf("grid1D_block1D\n");setGPU(); // 设置 GPU 设备// 定义矩阵的尺寸int nx = 16; // 矩阵的列数int ny = 8;  // 矩阵的行数int nxy = nx * ny; // 矩阵的总元素数size_t stBytesCount = nxy * sizeof(int); // 矩阵所需的总字节数// 在主机（CPU）上分配内存int *ipHost_A, *ipHost_B, *ipHost_C;ipHost_A = (int *)malloc(stBytesCount); // 分配矩阵 A 的内存ipHost_B = (int *)malloc(stBytesCount); // 分配矩阵 B 的内存ipHost_C = (int *)malloc(stBytesCount); // 分配结果矩阵 C 的内存// 初始化矩阵 A 和 B 的数据if(ipHost_A != NULL && ipHost_B != NULL && ipHost_C != NULL){for(int i = 0; i < nxy; i++){ipHost_A[i] = i;       // 矩阵 A 的元素设为 iipHost_B[i] = i + 1;   // 矩阵 B 的元素设为 i+1}} else {printf("fail to malloc memory.\n");exit(-1); // 如果分配失败，退出程序}// 在设备（GPU）上分配内存int *ipDevice_A, *ipDevice_B, *ipDevice_C;ErrorCheck(cudaMalloc((int **)&ipDevice_A, stBytesCount), __FILE__, __LINE__);ErrorCheck(cudaMalloc((int **)&ipDevice_B, stBytesCount), __FILE__, __LINE__);ErrorCheck(cudaMalloc((int **)&ipDevice_C, stBytesCount), __FILE__, __LINE__);if(ipDevice_A != NULL && ipDevice_B != NULL && ipDevice_C != NULL){// 将主机内存复制到设备ErrorCheck(cudaMemcpy(ipDevice_A, ipHost_A, stBytesCount, cudaMemcpyHostToDevice), __FILE__, __LINE__);ErrorCheck(cudaMemcpy(ipDevice_B, ipHost_B, stBytesCount, cudaMemcpyHostToDevice), __FILE__, __LINE__);} else {// 如果分配失败，释放内存并退出程序free(ipHost_A);free(ipHost_B);free(ipHost_C);exit(-1);}// 设置线程块和网格的尺寸dim3 block(4); // 每个块包含 4 个线程dim3 grid((nx + block.x - 1) / block.x); // 设置 1D 网格的维度printf("Thread config: grid (%d, %d), block(%d, %d).\n", grid.x, grid.y, block.x, block.y);// 启动 CUDA 内核addMatrix_11D<<<grid, block>>>(ipDevice_A, ipDevice_B, ipDevice_C, nx , ny);ErrorCheck(cudaDeviceSynchronize(), __FILE__, __LINE__);// 将结果从设备复制回主机ErrorCheck(cudaMemcpy(ipHost_C, ipDevice_C, stBytesCount, cudaMemcpyDeviceToHost), __FILE__, __LINE__);// 输出前 10 个元素的加法结果，验证计算正确性for(int i = 0; i < 10; i++){printf("idx=%2d\tmatrix_A:%d\tmatrix_B:%d\tresult=%d\n", i + 1, ipHost_A[i], ipHost_B[i], ipHost_C[i]);}// 释放主机和设备上的内存free(ipHost_A);free(ipHost_B);free(ipHost_C);ErrorCheck(cudaFree(ipDevice_A), __FILE__, __LINE__);ErrorCheck(cudaFree(ipDevice_B), __FILE__, __LINE__);ErrorCheck(cudaFree(ipDevice_C), __FILE__, __LINE__);ErrorCheck(cudaDeviceReset(), __FILE__, __LINE__);return;
}