C++ CUDA开发入门

CUDA开发笔记

文章目录

CUDA开发笔记
@[toc]
1 概述
2 环境
3 命令行编译
4 CMAKE引入CUDA
5 vscode开发CUDA
6 Qt中使用CUDA-CMake
7 QMake配置CUDA
8 核函数
9 核函数调用
9.1 核函数调用语法
9.2 执行配置参数详解
9.3 关键调用步骤
9.4 重要注意事项
9.5 调用示例分析
9.6 最佳实践建议

10 线程模型
11 示例程序
12 相关链接

更多精彩内容
👉内容导航 👈
👉Qt开发经验 👈
👉C++ 👈
👉开发工具 👈

1 概述

C++与CUDA的结合开发通常用于高性能计算，特别是在需要利用GPU进行并行计算的场景中，例如图像/视频处理、深度学习AI等。

高性能并行计算能力
GPU 拥有数千个计算核心，适用于大规模并行计算任务（如矩阵运算、物理模拟），性能远超 CPU。
适用于需要 FLOPs（浮点运算）密集的场景（如深度学习训练、科学计算）。

灵活的底层控制
允许直接管理 GPU 内存（显存）、线程调度、核函数优化，适合对性能有极致要求的场景。
支持细粒度并行（如线程块、网格的灵活划分）。

与 C++ 生态兼容
基于 C/C++ 语法扩展（通过 __global__、__device__ 等关键字），便于复用现有 C++ 代码。
提供 cuBLAS（线性代数）、cuFFT（快速傅里叶变换）等高性能库，加速开发。

跨平台支持
支持 Windows/Linux 系统，兼容 NVIDIA 全系列 GPU（如 Tesla、GeForce、Quadro）。

开发CUDA推荐使用CMake，配置简单；

不推荐使用QMake配置非常复杂。

2 环境

名称	说明
系统	windows11
编译器	msvc2017、msvc2022
CUDA版本	12.4
Cmake版本	≥ 3.8 (推荐 ≥ 3.18)

适用于 Microsoft Windows 的 CUDA 安装指南
适用于 Linux 的 CUDA 安装指南
CUDA 安装包下载地址

3 命令行编译

创建一个main.cu文件，编写CUDA代码；
使用nvcc main.cu命令编译；

4 CMAKE引入CUDA

传统方式 (已弃用)

cmakefind_package(CUDA REQUIRED) # 旧方法
cuda_add_executable(my_target ...) # 专用命令

需要显式调用 find_package(CUDA) 查找 CUDA 工具包
必须使用特殊命令（如 cuda_add_executable）处理 CUDA 文件

现代方式 (推荐)

▌ 方法一：在 project() 中声明语言

cmakeproject(my_project LANGUAGES CXX CUDA) # 声明项目支持 CUDA
add_executable(my_target main.cu) # 直接添加 .cu 文件

▌ 方法二：单独启用 CUDA

cmakeproject(my_project LANGUAGES CXX)
enable_language(CUDA) # 后续启用 CUDA 支持
add_library(my_lib SHARED kernel.cu)

5 vscode开发CUDA

test.h
```
void fun();
```

test.cu

#include <iostream>
#include "test.h"
using namespace std;__global__ void hello_gpu() 
{// 在GPU上打印Hello Worldprintf("Hello World from GPU!\n");
}void fun()
{hello_gpu<<<4, 4>>>(); // 启动内核，4个块，每个块有4个线程，执行16次hello_gpu()函数调用。cudaDeviceSynchronize(); // 等待GPU完成所有工作printf("Hello World from CPU!\n");
}

main.cpp

/********************************************************************************
* 文件名：   main.cpp
* 创建时间： 2025-03-03 14:38:56
* 开发者：   MHF
* 邮箱：     1603291350@qq.com
* 功能：     
*********************************************************************************/
#include<iostream>
#include "test.h"
using namespace std;int main()
{cout << "进入" << endl;fun();cout << "离开" << endl;return 0;
}

CMakeLists.txt

cmake_minimum_required(VERSION 3.30)
project(test1 LANGUAGES CXX CUDA)  # 新增CUDA语言支持# 设置C++标准
set(CMAKE_CXX_STANDARD 14)# 设置MSVC编译器使用UTF-8编码
if(MSVC)set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /utf-8")
endif()# 输出路径
set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/../bin)# 添加主程序
add_executable(test1 main.cu)

6 Qt中使用CUDA-CMake

这里构建工具使用CMAKE，配置非常简单；

创建一个Qt工程，构建工具使用CMake，编译器选择MSVC；
如图所示，将CMake版本修改为大于3.18，在project中添加CUDA支持；
创建一个test.cu文件，一个test.h文件，CUDA和C++代码需要分开；
在CMakeLists.txt中添加文件CUDA代码文件；

test.cu文件中代码如下所示：

#include "test.h"
#include "stdio.h"
__global__ void hello_gpu()
{// 在GPU上打印Hello Worldprintf("Hello World from GPU!\n");
}void fun()
{hello_gpu<<<4, 4>>>();   // 启动内核，4个块，每个块有4个线程，执行16次hello_gpu()函数调用。cudaDeviceSynchronize();   // 等待GPU完成所有工作printf("Hello World from CPU!\n");
}

test.h文件中声明函数；

#ifndef TEST_H
#define TEST_Hvoid fun();#endif   // TEST_H

在widget.cpp文件中调用fun()函数；
如下所示，调用fun()函数后使用GPU执行。

7 QMake配置CUDA

QT       += core guigreaterThan(QT_MAJOR_VERSION, 4): QT += widgets
CONFIG += c++17SOURCES += \main.cpp \widget.cppHEADERS += \widget.h\test.hFORMS += \widget.ui#----------------QMake CUDA配置--------------------------# CUDA 配置
CUDA_DIR = "D:/CUDA12.4.1"
CUDA_SOURCES += test.cu   # cuda源码
# CUDA_ARCH = sm_61  # 修改为你的 GPU 架构# 头文件和库路径
INCLUDEPATH += $$CUDA_DIR/include
QMAKE_LIBDIR += $$CUDA_DIR/lib/x64
LIBS += -lcudart -lcuda# NVCC 编译规则
CONFIG(debug, debug|release) {NVCC_OPTIONS += --use_fast_math -Xcompiler "/MDd"  # debug模式编译参数
} else {NVCC_OPTIONS += --use_fast_math -Xcompiler "/MD"  # release模式编译参数
}
cuda.commands = $$CUDA_DIR/bin/nvcc -c $$NVCC_OPTIONS -o ${QMAKE_FILE_OUT} ${QMAKE_FILE_NAME}
cuda.dependency_type = TYPE_C
cuda.input = CUDA_SOURCES
cuda.output = ${QMAKE_FILE_BASE}_cuda.obj
QMAKE_EXTRA_COMPILERS += cuda    # 指定附加编译器Nvcc

8 核函数

CUDA核函数（Kernel Function）是GPU并行计算的核心单元，是运行在NVIDIA GPU上的并行函数。以下是详细说明：

核函数定义

__global__ void kernelName(parameters) {// 并行代码块
}

使用 __global__ 限定符声明
返回值必须是 void
主机端调用，设备端执行

执行配置（Execution Configuration）

kernelName<<<gridDim, blockDim, sharedMemSize, stream>>>(arguments);

gridDim：网格维度（dim3类型，表示block排列）
blockDim：块维度（dim3类型，表示每个block中的thread排列）
sharedMemSize：动态共享内存大小（字节）
stream：执行流（默认为0）

线程层次结构

CUDA的线程层次结构包括线程（Thread）、块（Block）和网格（Grid）。每个块包含多个线程，而网格则包含多个块。这种层次结构使得CUDA能够高效地管理并行计算资源。

线程：是CUDA中最基本的执行单元，每个线程执行核函数的一次实例。
块：由多个线程组成，块内的线程可以共享数据，并通过共享内存进行通信。
网格：由多个块组成，网格内的块是独立的，不能直接通信，但可以通过全局内存来交换数据。

内置坐标变量：

threadIdx.x/y/z   // block内的线程坐标
blockIdx.x/y/z    // grid内的block坐标
blockDim.x/y/z    // block维度
gridDim.x/y/z     // grid维度

索引计算示例（矩阵相加）

int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
if (row < height && col < width) {C[row][col] = A[row][col] + B[row][col];
}

内存模型

CUDA的内存模型包括多种类型的内存空间，如全局内存、共享内存、常量内存和纹理内存等。这些内存空间具有不同的特性和用途，开发者需要根据实际需求选择合适的内存类型来优化性能。

全局内存：所有线程都可以访问，但访问速度相对较慢。
共享内存：每个块内的线程共享，访问速度非常快，但容量有限。
常量内存和纹理内存：用于特定类型的访问模式，可以提供优化的读取性能。

同步与原子操作

在CUDA中，线程间的同步是非常重要的。CUDA提供了一些同步原语，如__syncthreads()，用于确保块内线程在某一执行点达到同步。此外，CUDA还支持原子操作，用于在多个线程间安全地更新共享数据。

核函数特点

并行执行：数千个线程同时执行相同代码

注意事项

不能有返回值
不能是类的成员函数
不能递归调用
不能使用静态变量
不能调用主机端函数（除非__device__函数）
核函数只能访问GPU内存；
核函数不能使用变长参数；
不能使用函数指针；
核函数具有异步性。
核函数的执行时间通常较长，因此应尽量避免在核函数中进行大量的串行计算。
开发者需要仔细管理GPU上的内存资源，以避免内存泄漏或性能瓶颈。
核函数不支持C++的iostream。

动态并行（CUDA 5.0+）核函数可以启动其他核函数，但需要设备计算能力3.5+

最佳实践

每个线程处理独立计算任务
减少全局内存访问次数
使用共享内存优化数据重用
避免线程分支发散（warp divergence）

编译要求

使用NVCC编译器
需包含CUDA运行时头文件：

#include <cuda_runtime.h>

示例：向量加法核函数

__global__ void vectorAdd(float* A, float* B, float* C, int n) {int i = blockIdx.x * blockDim.x + threadIdx.x;if (i < n) {C[i] = A[i] + B[i];}
}
// 调用方式：vectorAdd<<<(n+255)/256, 256>>>(d_A, d_B, d_C, n);

9 核函数调用

以下是CUDA核函数调用的详细说明：

9.1 核函数调用语法

kernelName<<<执行配置>>>(参数列表);

9.2 执行配置参数详解

<<<gridDim, blockDim, sharedMemSize, stream>>>

参数	类型	说明	默认值
`gridDim`	`dim3`	Grid维度（Block排列方式）	`dim3(1,1,1)`
`blockDim`	`dim3`	Block维度（Thread排列方式）	必填参数
`sharedMemSize`	`size_t`	动态共享内存大小（字节）	`0`
`stream`	`cudaStream_t`	执行流（异步控制）	`0`（默认流）

9.3 关键调用步骤

配置线程层次

// 1D示例：处理N个元素，每个block 256线程
int blockSize = 256;
int gridSize = (N + blockSize - 1) / blockSize; // 向上取整
kernel<<<gridSize, blockSize>>>(...);// 2D示例：处理MxN矩阵
dim3 block(16, 16);        // 256 threads per block
dim3 grid((N+15)/16, (M+15)/16);
kernel2D<<<grid, block>>>(...);

内存准备

// 设备内存分配
float *d_data;
cudaMalloc(&d_data, size);// 主机到设备数据传输
cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice);

调用核函数

// 同步调用（默认流）
vectorAdd<<<grid, block>>>(d_A, d_B, d_C, N);// 异步调用（指定流）
cudaStream_t stream;
cudaStreamCreate(&stream);
kernel<<<grid, block, 0, stream>>>(...);

同步等待

cudaDeviceSynchronize(); // 等待所有流完成
cudaStreamSynchronize(stream); // 等待指定流

9.4 重要注意事项

线程数量限制
- 每个Block最多1024个线程（最新架构支持更多）
- Grid维度限制（x/y/z ≤ 65535）
内存访问
- 核函数参数必须指向设备内存（cudaMalloc分配）
- 不能直接访问主机内存

错误处理

cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {printf("Kernel launch error: %s\n", cudaGetErrorString(err));
}

动态并行（需计算能力≥3.5）

__global__ void parentKernel() {childKernel<<<1, 32>>>(); // 设备端启动核函数
}

9.5 调用示例分析

案例1：向量加法

// 核函数定义
__global__ void vecAdd(float *a, float *b, float *c, int n) {int i = blockIdx.x * blockDim.x + threadIdx.x;if (i < n) c[i] = a[i] + b[i];
}// 调用方式
int blockSize = 256;
int gridSize = (n + blockSize - 1) / blockSize;
vecAdd<<<gridSize, blockSize>>>(d_a, d_b, d_c, n);

案例2：矩阵乘法

// 核函数调用配置
dim3 threadsPerBlock(16, 16);
dim3 numBlocks((N + 15)/16, (M + 15)/16);
matrixMul<<<numBlocks, threadsPerBlock>>>(d_A, d_B, d_C, N, M);

9.6 最佳实践建议

线程数计算

// 推荐计算方式
blockSize = 128; // 或256（根据算法特点选择）
gridSize = (totalElements + blockSize - 1) / blockSize;

性能优化
- 使用__shared__声明共享内存
- 合并全局内存访问
- 避免线程束分化（Warp Divergence）

调试技巧

// 打印设备端信息（需重定向）
printf("Thread %d: value=%f\n", threadIdx.x, data);

10 线程模型

线程模型概念

grid：网格；
block：线程块；

线程分块是逻辑上的划分，物理上线程不分块；

配置线程：<<<grid_size，block_size>>>

最大允许线程块大小：1024；

最大允许网格大小：UINT_MAX - 1。

在这里插入图片描述

11 示例程序

main.cu

#include "test.h"
#include <cuda_runtime.h>
#include <stdio.h>
using namespace std;// 不能操作主机内存，只能使用GPU内存
__global__ void vecAdd( float* a,  float* b, float* c) 
{int i = threadIdx.x;c[i] =  a[i] + b[i];   printf("c[%d] = %f\n", i, c[i]);}void fun()
{// 在主机端分配设备内存float *d_a, *d_b, *d_c;cudaMalloc(&d_a, 10*sizeof(float));cudaMalloc(&d_b, 10*sizeof(float));cudaMalloc(&d_c, 10*sizeof(float));float a[10], b[10], c[10];for (int i = 0; i < 10; ++i) {a[i] = i;b[i] = i * 2;}// 拷贝数据到设备cudaMemcpy(d_a, a, 10*sizeof(float), cudaMemcpyHostToDevice);cudaMemcpy(d_b, b, 10*sizeof(float), cudaMemcpyHostToDevice);vecAdd<<<1, 10>>>(d_a, d_b, d_c);cudaDeviceSynchronize(); // 等待GPU完成所有工作
}