Vector Addition
1109 字
6 分钟
Vector Addition
Vector Addition
基本常识
CPU = 一个聪明的酋长。脑子快，一次只做一件事。
GPU = 一千个笨原始人。每个都笨，但一千个笨的打赢一个聪明的。
主机 Host = CPU 世界，C++，酋长坐的位置。
设备 Device = GPU 世界，原始人的住处，不同内存，不同空间。
::: tip CPU 不能碰 GPU 内存。GPU 不能碰 CPU 内存，必须过桥传送数据，也就是 PCIe。 :::
三个魔法词
__global__ void kernel() { }// CPU 调用, GPU 执行。"原始人，冲！"__device__ void helper() { }// GPU 调, GPU 跑。原始人悄悄话。__host__ void normal() { }// CPU 调, CPU 跑。普通C++。(默认，可省略)线程层级：三层分离
Grid 大营 ├── Block 0 小部落 │ ├── Thread 0 原始人 │ ├── Thread 1 │ └── Thread 2 ... ├── Block 1 │ └── Thread 0, 1, 2 ... └── Block 2 ...::: tip 每个 block 里，32 个 thread 组成一个 warp。 :::
内置变量：不同原始人有不同的身份划分
threadIdx.x // 我是部落里第几个原始人？ (0 到 blockDim.x-1)blockIdx.x // 我部落在大营里排第几？ (0 到 gridDim.x-1)blockDim.x // 我部落有多少原始人？ (比如 256)gridDim.x // 大营有多少部落？ (比如 4)执行 <<< >>>：派遣原始人
kernel<<<部落数, 每部落原始人数>>>(参数);// 有1000个数。每个原始人处理1个。// 需要1000个原始人。每部落最多256个。int N = 1000;int threads = 256;int blocks = (N + threads - 1) / threads; // 4 个部落
add_kernel<<<blocks, threads>>>(a, b, c, N);// 4部落 × 256原始人 = 1024 个原始人。// 最后一个部落：256 - 24 = 232个干活，24个摸鱼。内存：四种石头
| 内存 | 速度 | 谁能看见 | 生命周期 | 大小 |
|---|---|---|---|---|
| 全局 | 慢（像走到山上） | 所有原始人，所有部落 | 整个程序 | 大（GB） |
| 共享 | 快（像递给旁边兄弟） | 同部落的原始人 | 一次kernel运行 | 小（48KB） |
| 局部 | 快（口袋里的石头） | 一个原始人 | 一次kernel运行 | 小 |
| 寄存器 | 瞬间（手里攥着） | 一个原始人 | 一次kernel运行 | 最小（每原始人255个） |
__global__ void example(float* global_arr) { // global_arr — 所有人都能看到。远处大石头上写着。
__shared__ float shared_arr[256]; // 只有我部落能看到。篝火旁边。
float local_val = 5.0f; // 我口袋里的。寄存器或局部内存。}搬运内存：过桥费
// 1. 分配 GPU 内存float* d_data; // "d_" 前缀 = device。原始人约定。cudaMalloc(&d_data, N * sizeof(float));
// 2. 复制 CPU → GPU（给原始人送活儿）cudaMemcpy(d_data, h_data, N * sizeof(float), cudaMemcpyHostToDevice);
// 3. 启动 kernel（原始人干活）kernel<<<blocks, threads>>>(d_data, N);
// 4. 复制 GPU → CPU（原始人带结果回来）cudaMemcpy(h_result, d_result, N * sizeof(float), cudaMemcpyDeviceToHost);
// 5. 释放 GPU 内存（打扫山洞）cudaFree(d_data);::: details 五步法总结
- cudaMalloc
- cudaMemcpy
- kernel <<<>>>
- cudaMemcpy
- cudaFree :::
__syncthreads() 同步：部落开会
__global__ void example() { // 第一阶段：每个原始人摘果子 shared_basket[threadIdx.x] = gather_berry();
__syncthreads(); // 部落里所有原始人在此停下。等每个人。 // 没摘完果子之前，谁也不许偷看篮子。
// 第二阶段：现在安全了，可以数果子 if (threadIdx.x == 0) { total = sum(shared_basket); // 小队长统计。 }}Warp：背后的真相
::: warning GPU 欺骗了你，它看似只有 256 个原始人一起干活，现实是一个 Warp = 32 个原始人，GPU 一次只执行 32 个。这 32 个必须做同一条指令。（SIMD - 单指令多数据） :::
::: tip Warp 分支发散 = 坏文明 :::
if (threadIdx.x < 16) { do_A(); // 16 个原始人走 A 路。另外 16 个站着发呆。} else { do_B(); // 16 个原始人走 B 路。另外 16 个站着发呆。}// 总时间 = A时间 + B时间。一半原始人永远在摸鱼。浪费！::: tip 好代码避免 warp 里写 if-else。32 个原始人一起齐步走。 :::
第一个程序：Vector Addition
题目
代码
#include <cuda_runtime.h>
__global__ void vector_add(const float* A, const float* B, float* C, int N) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < N) { C[i] = A[i] + B[i]; }}
// A, B, C are device pointers (i.e. pointers to memory on the GPU)extern "C" void solve(const float* A, const float* B, float* C, int N) { int threadsPerBlock = 256; int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
vector_add<<<blocksPerGrid, threadsPerBlock>>>(A, B, C, N); cudaDeviceSynchronize();}代码速查表：军师级别的原始人专用
| 干啥 | 咋写 |
|---|---|
| GPU 函数 | __global__ void name(...) |
| 启动 | name<<<部落数, 原始人数>>>(参数) |
| 我的编号 | blockIdx.x * blockDim.x + threadIdx.x |
| 部落墙 | __shared__ float arr[大小] |
| 部落等齐 | __syncthreads() |
| GPU 分配 | cudaMalloc(&ptr, 字节数) |
| 拷内存 | cudaMemcpy(目标, 源, 字节, 方向) |
| 等 GPU 干完 | cudaDeviceSynchronize() |
| GPU 释放 | cudaFree(ptr) |
Vector Addition
https://dongyanzhang.com/posts/leetgpu/vector-addition/