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完成所有工作
}