C++ CUDA开发入门

CUDA开发笔记

---

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

---

1 概述

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

1. 高性能并行计算能力
   • GPU 拥有数千个计算核心，适用于大规模并行计算任务（如矩阵运算、物理模拟），性能远超 CPU。
   • 适用于需要 FLOPs（浮点运算）密集的场景（如深度学习训练、科学计算）。
2. 灵活的底层控制
   • 允许直接管理 GPU 内存（显存）、线程调度、核函数优化，适合对性能有极致要求的场景。
   • 支持细粒度并行（如线程块、网格的灵活划分）。
3. 与 C++ 生态兼容
   • 基于 C/C++ 语法扩展（通过 __global__、__device__ 等关键字），便于复用现有 C++ 代码。
   • 提供 cuBLAS（线性代数）、cuFFT（快速傅里叶变换）等高性能库，加速开发。
4. 跨平台支持
   • 支持 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 World
      printf("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
  * 邮箱：     [e-mail]
  * 功能：     
  *********************************************************************************/
  #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，配置非常简单；

1. 创建一个Qt工程，构建工具使用CMake，编译器选择MSVC；

2. 如图所示，将CMake版本修改为大于3.18，在project中添加CUDA支持；

   

3. 创建一个test.cu文件，一个test.h文件，CUDA和C++代码需要分开；

4. 在CMakeLists.txt中添加文件CUDA代码文件；

   

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

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

6. test.h文件中声明函数；

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

7. 在widget.cpp文件中调用fun()函数；

   

8. 如下所示，调用fun()函数后使用GPU执行。

   

7 QMake配置CUDA

QT       += core gui

greaterThan(QT_MAJOR_VERSION, 4): QT += widgets
CONFIG += c++17

SOURCES += \
    main.cpp \
    widget.cpp

HEADERS += \
    widget.h\
    test.h

FORMS += \
    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 关键调用步骤

1. 配置线程层次

   // 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>>>(...);
   

2. 内存准备

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

3. 调用核函数

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

4. 同步等待

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

---

9.4 重要注意事项

1. 线程数量限制

   • 每个Block最多1024个线程（最新架构支持更多）
   • Grid维度限制（x/y/z ≤ 65535）

2. 内存访问

   • 核函数参数必须指向设备内存（cudaMalloc分配）
   • 不能直接访问主机内存

3. 错误处理

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

4. 动态并行（需计算能力≥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 最佳实践建议

1. 线程数计算

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

2. 性能优化

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

3. 调试技巧

   // 打印设备端信息（需重定向）
   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完成所有工作
  }
  
  

12 相关链接

• CUDA 工具包文档 12.4 更新 1
• CUDA 工具包 12.4 Update 1 下载 |NVIDIA 开发者
• 目录 — 快速入门指南 12.8 文档
• CUDA 运行时 API ：： CUDA 工具包文档
• CUDA Runtime API
• CUDA Toolkit Archive | NVIDIA Developer
• CUDA 工具包文档 12.8
• CUDA示例
• 适用于 Microsoft Windows 的 CUDA 安装指南
• 适用于 Linux 的 CUDA 安装指南
• CUDA 安装包下载地址