【详细】CUDA开发学习教程清单
以下是一份系统化的CUDA开发学习教程清单,涵盖基础到进阶内容,并附带关键代码示例:
📚 一、基础环境搭建(Ubuntu示例)
1. 安装CUDA Toolkit
# 下载.run文件后执行安装
sudo sh cuda_12.3.2_524.91.01_linux.run
# 配置环境变量
echo 'export PATH=/usr/local/cuda/bin:$PATH' >> ~/.bashrc
echo 'export LD_LIBRARY_PATH=/usr/local/cuda/lib64:$LD_LIBRARY_PATH' >> ~/.bashrc
source ~/.bashrc
# 验证安装
nvcc --version # 输出版本信息即成功
2. 验证GPU支持
#include <cuda_runtime.h>
#include <stdio.h>
int main() {int deviceCount;cudaGetDeviceCount(&deviceCount);printf("可用GPU数量: %d\n", deviceCount);return 0;
}
// 编译: nvcc device_query.cu -o device_query
⚙️ 二、核心编程模型
1. Hello GPU(首个核函数)
#include <stdio.h>
__global__ void helloGPU() {printf("Hello from GPU thread %d!\n", threadIdx.x);
}
int main() {helloGPU<<<1, 5>>>(); // 1个块,5个线程cudaDeviceSynchronize(); // 等待GPU执行return 0;
}
// 输出5行问候语
2. 向量加法(线程组织)
__global__ void vectorAdd(int *a, int *b, int *c, int n) {int idx = blockIdx.x * blockDim.x + threadIdx.x; // 全局索引if (idx < n) c[idx] = a[idx] + b[idx];
}
int main() {int n = 1000;int *d_a, *d_b, *d_c;cudaMalloc(&d_a, n*sizeof(int));cudaMalloc(&d_b, n*sizeof(int));cudaMalloc(&d_c, n*sizeof(int));// 数据初始化及传输省略...dim3 blockSize(256); // 每块256线程dim3 gridSize((n + 255)/256); // 计算所需块数vectorAdd<<<gridSize, blockSize>>>(d_a, d_b, d_c, n);cudaMemcpy(h_c, d_c, n*sizeof(int), cudaMemcpyDeviceToHost);// 释放资源省略...
}
🧠 三、内存管理进阶
1. 共享内存加速矩阵乘法
__global__ void matMul(float *A, float *B, float *C, int N) {__shared__ float sA[16][16]; // 块内共享内存__shared__ float sB[16][16];int tx = threadIdx.x, ty = threadIdx.y;int row = blockIdx.y * blockDim.y + ty;int col = blockIdx.x * blockDim.x + tx;float sum = 0;for (int i = 0; i < N; i += blockDim.x) {sA[ty][tx] = A[row*N + i + tx]; // 协作加载数据块sB[ty][tx] = B[(i+ty)*N + col];__syncthreads(); // 同步线程确保数据就绪for (int k = 0; k < blockDim.x; k++) sum += sA[ty][k] * sB[k][tx];__syncthreads();}C[row*N+col] = sum;
}
// 调用示例: dim3 block(16,16); matMul<<<grid, block>>>(A, B, C, N);
🚀 四、实战应用场景
1. OpenCV图像处理(CUDA加速)
#include <opencv2/cudaimgproc.hpp>
int main() {cv::Mat src_host = cv::imread("image.jpg");cv::cuda::GpuMat src, gray;src.upload(src_host); // 上传到GPUcv::cuda::cvtColor(src, gray, cv::COLOR_BGR2GRAY); // GPU灰度转换cv::Mat gray_host;gray.download(gray_host); // 下载结果
}
2. 并行归约求最大值
__global__ void reduceMax(float *d_in, float *d_out) {extern __shared__ float sdata[];int tid = threadIdx.x;int i = blockIdx.x * blockDim.x + tid;sdata[tid] = d_in[i];__syncthreads();for (int s = blockDim.x/2; s>0; s>>=1) {if (tid < s) sdata[tid] = fmaxf(sdata[tid], sdata[tid+s]);__syncthreads();}if (tid == 0) d_out[blockIdx.x] = sdata[0];
}
// 调用: reduceMax<<<grid, block, smem_size>>>(d_input, d_output);
📖 五、推荐学习资源
-
官方文档
- CUDA C++ Programming Guide - 架构详解
- CUDA Best Practices Guide - 性能优化手册
-
中文教程
- 谭升博客《CUDA入门教程》:逐行注释的GitHub代码
- 知乎专栏《CUDA编程入门》:从Hello World到内存优化
- 《》中文版:系统化概念+工业级案例
-
开源项目实践
- LightSeq : Transformer推理加速库
- CUDA Samples : NVIDIA官方示例集
💡 六、调试与优化工具
工具名称 | 用途 | 命令示例 |
---|---|---|
cuda-gdb | GPU线程级调试 | cuda-gdb ./my_program |
nsys | 性能分析(时间线) | nsys profile ./my_program |
ncu | 内核指令级分析 | ncu --set full ./my_program |
学习路径建议:从向量加法 → 矩阵运算 → 共享内存优化 → 结合深度学习框架。所有代码需用
nvcc
编译,注意GPU架构匹配(如RTX 4090需-arch=sm_89
)。遇到内存错误时,使用cuda-memcheck ./program
检测越界访问。