CUDA編程筆記（1）--最簡單的核函數

本文用來記錄cuda編程的一些筆記以及知識

本筆記運行在windows系統，vs編譯器中，cuda版本是12.6

先看一下最基本的代碼例子：

#include<iostream>
#include<cstdio>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"//它聲明了你在 CUDA kernel 函數中會使用的一些 內置變量
__global__ void kernel() {printf("hello world");
}
int main() {kernel <<<1, 1 >>> ();cudaError_t err = cudaDeviceSynchronize();if (err != cudaSuccess) {std::cerr << "CUDA Error: " << cudaGetErrorString(err) << std::endl;}return 0;
}

下面我會介紹所有出現的東西，以讓大家對cuda編程有一個最基本的理解

#include "cuda_runtime.h"

? 提供 CUDA 運行時 API（Runtime API）的聲明

該頭文件包含了調用 CUDA 核函數所需的所有基本功能聲明，例如：

🚀 設備管理類函數

cudaGetDevice()

cudaSetDevice()

cudaGetDeviceCount()

🧠 內存管理函數

cudaMalloc()：在 GPU 上分配內存

cudaFree()：釋放 GPU 上的內存

cudaMemcpy()：主機和設備之間的數據拷貝

🧮 執行控制函數

cudaDeviceSynchronize()：等待所有 GPU 上的任務完成

cudaDeviceReset()：復位設備（常用于程序退出前）

cudaGetLastError()：獲取上一次 CUDA 操作的錯誤

? 支持 kernel 函數的 <<<>>> 調用語法

沒有這個頭文件，編譯器就不會正確識別：

kernel<<<numBlocks, threadsPerBlock>>>(...);

這種語法就是 CUDA 的核函數啟動方式，必須在包含 cuda_runtime.h 后才有效。

? 加載所有 device API 所需的類型和結構體

比如：

cudaError_t：CUDA 錯誤碼的類型

cudaMemcpyKind：用于標識拷貝方向（如 cudaMemcpyHostToDevice）

?#include "device_launch_parameters.h"

? 提供 CUDA 核函數啟動時需要的一些內置變量聲明

這個頭文件聲明了你在 CUDA 核函數（__global__ 函數）中常用的以下 線程索引相關變量：

變量名說明
threadIdx 當前線程在 block 中的索引（例如 threadIdx.x）
blockIdx 當前 block 在 grid 中的索引
blockDim 當前 block 中線程的維度
gridDim 當前 grid 的 block 數量

這些變量雖然在 CUDA 編譯器中是“內置”的，但如果沒有包含適當的頭文件，有些編譯器（或靜態分析工具）會報聲明缺失的警告。這個頭文件就是為了避免此類問題。

? 為什么要顯式 include 它？

雖然在大多數情況下 #include <cuda_runtime.h> 已經間接包含了它，但顯式寫出來有幾個好處：

可讀性更好：清晰知道 kernel 中用到哪些內置變量；

避免 IDE 報錯：有些開發工具（如 Visual Studio）如果沒有包含這個頭文件，在編輯器中會提示這些變量未聲明；

提高代碼可移植性：一些平臺或舊版本 CUDA SDK 不自動包含它。

變量名	說明
`threadIdx`	當前線程在 block 中的索引（例如 `threadIdx.x`）
`blockIdx`	當前 block 在 grid 中的索引
`blockDim`	當前 block 中線程的維度
`gridDim`	當前 grid 的 block 數量

核函數

? 一、核函數的基本語法

🔹 定義核函數
__global__ void kernel_name(parameter_list) { // GPU 上執行的代碼 }
__global__：CUDA 的函數修飾符，表示該函數：

在 GPU 上執行

由 CPU 主機代碼調用

kernel_name：核函數的名稱。

parameter_list：參數列表，可以傳遞普通數據指針（如設備內存地址）或標量值。

🔹 調用核函數（使用 <<< >>> 執行配置）
kernel_name<<<numBlocks, threadsPerBlock>>>(parameter_values);
numBlocks：grid 中的 block 數量（可以是一維、二維或三維）。

threadsPerBlock：每個 block 中的線程數量（同樣支持一維、二維或三維）。

parameter_values：傳遞給核函數的參數。

? 示例（最常見的一維）：
__global__ void addKernel(int* data) { 
int idx = blockIdx.x * blockDim.x + threadIdx.x; data[idx] += 1; 
} 
int main() {addKernel<<<10, 256>>>(device_ptr); // 共啟動 2560 個線程 
}
? 二、核函數中常用的內置變量

變量名類型說明
threadIdx dim3 當前線程在其所屬 block 中的索引
blockIdx dim3 當前 block 在 grid 中的索引
blockDim dim3 當前 block 中線程的數量
gridDim dim3 當前 grid 中 block 的數量

注意：這些變量都是結構體 dim3 類型，支持 .x, .y, .z 維度訪問。

? 示例：全局索引計算
int globalIndex = blockIdx.x * blockDim.x + threadIdx.x;
🚧 三、核函數的注意事項

1. 核函數只能由 CPU 調用，不能被設備端其他函數調用

核函數必須使用 __global__ 修飾；

不能在 GPU 設備端用另一個核函數調用它；

如果你想在設備端調用，可以使用 __device__ 或 __host__ __device__ 修飾的函數。

2. 核函數返回類型必須是 void
__global__ void foo() {}// ? 合法 
__global__ int foo() { return 1; } ///? 不合法
3. 不能使用 std::cout，只能使用 printf()

核函數中不支持 C++ 的標準流輸出。

使用 #include <cstdio> 和 printf() 進行調試。

4. 核函數中的數組必須放在共享內存或全局內存中
__global__ void kernel() { 
int arr[10]; // ? 合法，小數組在寄存器中 // 大數組建議用 `__shared__` 或 `cudaMalloc`}
5. 核函數必須同步（如 cudaDeviceSynchronize()）

如果主機代碼緊接著要訪問 GPU 數據，必須使用：
cudaDeviceSynchronize(); 
6. 線程索引越界要自己判斷！

GPU 不會自動防止訪問越界；

通常你要在核函數開頭寫判斷邏輯：

7. 設備函數（__device__）可以在核函數中調用
__device__ int square(int x) { return x * x; }__global__ void kernel(int* out) {int idx = threadIdx.x;out[idx] = square(idx);
}
? 總結：核函數語法關鍵點

元素用法示例
定義方式 __global__ void kernel(...)
啟動語法 kernel<<<blocks, threads>>>(...)
內置索引變量 blockIdx, threadIdx, blockDim, gridDim
索引計算 int idx = blockIdx.x * blockDim.x + threadIdx.x
輸出方式 printf()
同步函數 cudaDeviceSynchronize()

變量名	類型	說明
`threadIdx`	dim3	當前線程在其所屬 block 中的索引
`blockIdx`	dim3	當前 block 在 grid 中的索引
`blockDim`	dim3	當前 block 中線程的數量
`gridDim`	dim3	當前 grid 中 block 的數量

元素	用法示例
定義方式	`__global__ void kernel(...)`
啟動語法	`kernel<<<blocks, threads>>>(...)`
內置索引變量	`blockIdx`, `threadIdx`, `blockDim`, `gridDim`
索引計算	`int idx = blockIdx.x * blockDim.x + threadIdx.x`
輸出方式	`printf()`
同步函數	`cudaDeviceSynchronize()`

grid，block，thread

? 一、三者之間的關系

Grid（網格） └── 多個 Block（線程塊） ? ? ?└── 多個 Thread（線程）

Grid：一次 kernel 啟動時生成的所有線程塊的集合。

Block：由多個線程組成的執行單元，GPU 以 block 為基本調度單位。

Thread：最基本的并行執行單元。

? 二、它們的作用與結構

1. Thread（線程）

執行最小單位，每個線程有自己的局部變量。

每個線程可以根據自己的 ID 來處理不同的數據元素。

CUDA 中線程 ID：

threadIdx.x, threadIdx.y, threadIdx.z

2. Block（線程塊）

一個線程塊內的線程可以：

使用共享內存（__shared__）進行數據共享

通過 __syncthreads() 進行同步

每個線程塊最多可有 1024 個線程（受 GPU 限制）

CUDA 中 block ID：

blockIdx.x, blockIdx.y, blockIdx.z

CUDA 中 block 內線程維度：

blockDim.x, blockDim.y, blockDim.z

3. Grid（網格）

Grid 是所有 Block 的集合。

所有的 Block 并行執行。

Grid 維度：

gridDim.x, gridDim.y, gridDim.z

? 三、線程索引計算：獲取線程全局 ID

一般我們會用下面公式獲取線程在整個 Grid 中的全局 ID（以一維為例）：

int global_id = blockIdx.x * blockDim.x + threadIdx.x;

二維情形：

int row = blockIdx.y * blockDim.y + threadIdx.y;

int col = blockIdx.x * blockDim.x + threadIdx.x;

? 四、維度示意圖（以 2D 舉例）
GridDim (2x2)
┌────────────┬────────────┐
│ Block(0,0) │ Block(1,0) │
│ Threads ? ?│ Threads ? ?│
└────────────┴────────────┘
┌────────────┬────────────┐
│ Block(0,1) │ Block(1,1) │
│ Threads ? ?│ Threads ? ?│
└────────────┴────────────┘
每個 Block 里還有：
Threads (e.g., 16x16)
┌──────┐
│ t0 ? │
│ t1 ? │
│ ... ?│
└──────┘
? 五、核函數調用中的 <<<>>> 配置

kernel<<<gridDim, blockDim>>>(...);

gridDim：Grid 的尺寸（可以是 int 或 dim3）

blockDim：每個 Block 中線程的數量

例如：

dim3 grid(4, 4); // 4x4 個 Block

dim3 block(8, 8); // 每個 Block 8x8 個線程（共 64 個線程）

kernel<<<grid, block>>>(...);

總線程數：4×4×8×8 = 1024

? 六、一些常見注意事項

注意點說明
每個 Block 中線程總數限制通常為 1024（可查詢設備屬性）
線程越界處理你必須手動判斷是否越界，否則可能非法訪問內存
使用共享內存只能在線程塊內部共享，線程間需同步
不同 Block 不共享變量 Block 之間不能通信（除非使用 global memory）
線程的層級結構是固定的 Grid → Block → Thread 是一層層嵌套的

? 七、小結表格

層級名稱定義方式標識符（內置變量）
1 Grid <<<gridDim, ...>>> gridDim, blockIdx
2 Block <<<..., blockDim>>> blockDim, threadIdx
3 Thread 在核函數中自動生成 threadIdx

注意點	說明
每個 Block 中線程總數限制	通常為 1024（可查詢設備屬性）
線程越界處理	你必須手動判斷是否越界，否則可能非法訪問內存
使用共享內存	只能在線程塊內部共享，線程間需同步
不同 Block 不共享變量	Block 之間不能通信（除非使用 global memory）
線程的層級結構是固定的	Grid → Block → Thread 是一層層嵌套的

層級	名稱	定義方式	標識符（內置變量）
1	Grid	`<<<gridDim, ...>>>`	`gridDim`, `blockIdx`
2	Block	`<<<..., blockDim>>>`	`blockDim`, `threadIdx`
3	Thread	在核函數中自動生成	`threadIdx`

?在示例中，<<<1,1>>>表示我們只啟動了一個線程，因此hello world只會輸出一次；我們將代碼改為如下：

#include<iostream>
#include<cstdio>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"//它聲明了你在 CUDA kernel 函數中會使用的一些 內置變量
__global__ void kernel() {/*int current_thread_id = blockIdx.x * blockDim.x + threadIdx.x;printf("current thread id is :%d\n", current_thread_id);*/printf("hello world\n");
}
int main() {kernel <<<4, 8 >>> ();cudaError_t err = cudaDeviceSynchronize();if (err != cudaSuccess) {std::cerr << "CUDA Error: " << cudaGetErrorString(err) << std::endl;}return 0;
}

grid中有4個block，每個block擁有8個thread，所以總共有32個線程，也就會輸出32個hello world。

計算線程的唯一id

線程在線性空間中的全局編號（global thread id）

int global_id =threadIdx.x+ threadIdx.y * blockDim.x+ threadIdx.z * blockDim.x * blockDim.y+ blockIdx.x * blockDim.x * blockDim.y * blockDim.z+ blockIdx.y * gridDim.x * blockDim.x * blockDim.y * blockDim.z+ blockIdx.z * gridDim.x * gridDim.y * blockDim.x * blockDim.y * blockDim.z;

分解解釋：

線程在線程塊內的編號（local thread offset）：

local_id = threadIdx.x+ threadIdx.y * blockDim.x+ threadIdx.z * blockDim.x * blockDim.y;

線程塊的編號（block offset）：?

block_id = blockIdx.x+ blockIdx.y * gridDim.x+ blockIdx.z * gridDim.x * gridDim.y;

每個 Block 中的線程數量：

threads_per_block = blockDim.x * blockDim.y * blockDim.z;

?所以也可以寫成：

global_id = block_id * threads_per_block + local_id;

如果你處理的是三維數組索引

除了線性 global_id，你還可以直接求出每個線程在全局的 (x, y, z) 坐標：

int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int z = blockIdx.z * blockDim.z + threadIdx.z;

類比：想象一組三維數組任務

比如你現在有個數據張量：float data[64][64][64]，你希望：

每個線程負責處理這個數組的一個元素，比如 data[z][y][x]。
CUDA 會把這些任務分給多個線程，你就用上面那行代碼求出：每個線程該處理哪個 x,y,z 的位置。

具體例子

假設你啟動 CUDA kernel去處理一個[8][8][8]的數組（當然也可以是一維的）

kernel<<<dim3(2,2,2), dim3(4,4,4)>>>();

這表示：

Grid 有 2x2x2=8 個 Block
每個 Block 有 4x4x4=64 個線程
所以總線程數 = 8 × 64 = 512

那么：

第一個 Block（blockIdx=(0,0,0)）內的線程 threadIdx=(0,0,0) 處理的是 (x=0, y=0, z=0)
第二個 Block（blockIdx=(1,0,0)）內的 threadIdx=(0,0,0) 處理的是 (x=4, y=0, z=0)（x 方向偏移了一個 Block）
第三個 Block（blockIdx=(0,1,0)）內的 threadIdx=(0,0,0) 處理的是 (x=0, y=4, z=0)（y 方向偏移了一個 Block）
第四個 Block（blockIdx=(1,1,0)）內的 threadIdx=(1,2,3) 處理的是 (x=5, y=6, z=3)

cudaDeviceSynchronize的作用

cudaDeviceSynchronize() 是 CUDA 編程中的一個重要函數，它的主要作用是：

阻塞當前 CPU 線程，直到 GPU 上之前所有的任務全部完成。

詳細來說：

CUDA 的操作（如內核函數啟動、內存拷貝等）是異步執行的，調用它們時 CPU 不會等待 GPU 完成任務就直接繼續執行后面的代碼。
cudaDeviceSynchronize() 會讓 CPU 阻塞，直到 GPU 上所有之前發起的任務（比如內核執行、數據傳輸等）都執行完畢，確保 GPU 任務完成后再繼續執行 CPU 代碼。
它常用于調試、性能測量或者需要確保 GPU 計算結果已經完成才能進行下一步處理的場景。