1
CPU 处理 1000 个数据：
2
  1 个核心 × 1 ns/操作 × 1000 = 1000 ns
3
  8 个核心 × 1 ns/操作 × 125  = 125 ns（并行）
4

5
GPU 处理 1000 个数据：
6
  1000 个核心 × 10 ns/操作 × 1  = 10 ns（全并行）
7

8
GPU 的优势：当任务可以分解为大量独立操作时
9
GPU 的劣势：当任务需要复杂分支或串行逻辑时

1
// CUDA kernel：指定 Grid 和 Block 大小
2
__global__ void vectorAdd(float* a, float* b, float* c, int n) {
3
    int i = blockIdx.x * blockDim.x + threadIdx.x;
4
    if (i < n) {
5
        c[i] = a[i] + b[i];
6
    }
7
}
8

9
// 启动 kernel：256 个线程/Block，n/256 个 Block
10
int blockSize = 256;
11
int numBlocks = (n + blockSize - 1) / blockSize;
12
vectorAdd<<<numBlocks, blockSize>>>(a, b, c, n);

1
// 非合并访问：线程访问分散的地址
2
__global__ void badAccess(float* data) {
3
    int i = threadIdx.x * 256;  // 步长 256
4
    float val = data[i];        // 32 个线程访问 32 个分散地址
5
    // 每个线程单独访问全局内存 → 32 次内存事务
6
}
7

8
// 合并访问：线程访问连续地址
9
__global__ void goodAccess(float* data) {
10
    int i = threadIdx.x;  // 步长 1
11
    float val = data[i];  // 32 个线程访问连续 128 字节
12
    // 一次 128 字节内存事务 → 1 次访问
13
}

1
一个 SM（Streaming Multiprocessor）包含：
2
├── 64-128 个 CUDA 核心（FP32）
3
├── 32-64 个 FP64 核心
4
├── 4-8 个 Tensor 核心（矩阵加速）
5
├── 1-2 个 Warp Scheduler
6
├── 2-4 个 Dispatch Unit
7
├── 64-128 KB 寄存器文件
8
├── 48-96 KB 共享内存
9
├── 128 KB L1 缓存
10
└── 加载/存储单元
11
└── 特殊功能单元（SFU）
12

13
NVIDIA H100 (Hopper)：
14
  132 个 SM × 128 核心/SM = 16896 个 FP32 核心
15
  132 个 SM × 64 核心/SM = 8448 个 FP64 核心
16
  132 个 SM × 4 Tensor 核心/SM = 528 个 Tensor 核心

1
// 使用 CUDA Occupancy API 查询最优配置
2
int blockSize = 256;
3
int numBlocks;
4
cudaOccupancyMaxActiveBlocksPerMultiprocessor(&numBlocks,
5
    myKernel, blockSize, 0);
6
float occupancy = (numBlocks * blockSize) /
7
    (float)maxThreadsPerSM;
8
printf("Occupancy: %.1f%%\n", occupancy * 100);

1
// 使用 WMMA (Warp Matrix Multiply-Accumulate) API
2
#include <mma.h>
3
using namespace nvcuda::wmma;
4

5
__global__ void tensorCoreMatmul(half* A, half* B, float* C, int N) {
6
    fragment<matrix_a, 16, 16, 16, half, row_major> a_frag;
7
    fragment<matrix_b, 16, 16, 16, half, col_major> b_frag;
8
    fragment<accumulator, 16, 16, 16, float> c_frag;
9

10
    fill_fragment(c_frag, 0.0f);
11

12
    for (int k = 0; k < N; k += 16) {
13
        load_matrix_sync(a_frag, A + ..., N);
14
        load_matrix_sync(b_frag, B + ..., N);
15
        mma_sync(c_frag, a_frag, b_frag, c_frag);
16
    }
17
    store_matrix_sync(C + ..., c_frag, N, mem_row_major);
18
}

1
// 多 Stream 实现 CPU-GPU 传输与计算的 overlap
2
cudaStream_t stream0, stream1;
3
cudaStreamCreate(&stream0);
4
cudaStreamCreate(&stream1);
5

6
int chunkSize = N / 2;
7
// Stream 0：处理前半部分数据
8
cudaMemcpyAsync(d_A0, h_A0, chunkSize * sizeof(float),
9
                cudaMemcpyHostToDevice, stream0);
10
kernel<<<grid, block, 0, stream0>>>(d_A0, d_B0, d_C0, chunkSize);
11
cudaMemcpyAsync(d_C0, h_C0, chunkSize * sizeof(float),
12
                cudaMemcpyDeviceToHost, stream0);
13

14
// Stream 1：处理后半部分数据（与 Stream 0 并行）
15
cudaMemcpyAsync(d_A1, h_A1, chunkSize * sizeof(float),
16
                cudaMemcpyHostToDevice, stream1);
17
kernel<<<grid, block, 0, stream1>>>(d_A1, d_B1, d_C1, chunkSize);
18
cudaMemcpyAsync(d_C1, h_C1, chunkSize * sizeof(float),
19
                cudaMemcpyDeviceToHost, stream1);
20

21
cudaDeviceSynchronize();  // 等待所有 Stream 完成

1
// Unified Memory：无需手动 cudaMemcpy
2
float* data;
3
cudaMallocManaged(&data, N * sizeof(float));
4

5
// CPU 初始化
6
for (int i = 0; i < N; i++) data[i] = i * 0.5f;
7

8
// GPU 计算（驱动自动将数据迁移到 GPU）
9
kernel<<<grid, block>>>(data, N);
10
cudaDeviceSynchronize();
11

12
// CPU 读取结果（驱动自动将数据迁移回 CPU）
13
printf("result[0] = %f\n", data[0]);
14

15
cudaFree(data);

1
// GPU 上的可见性问题
2
__device__ int data = 0;
3
__device__ int flag = 0;
4

5
// Thread 0（生产者）
6
data = 42;           // 写入数据
7
__threadfence();     // 确保所有先前写操作对全局可见
8
flag = 1;            // 设置标志
9

10
// Thread 1（消费者）
11
while (flag != 1);   // 等待标志
12
// 没有 __threadfence()，可能看到 flag=1 但 data 仍是 0！
13
int val = data;      // 需要 fence 保证读到 42

1
// 基础矩阵乘法：C = A × B
2
__global__ void matmul_naive(float* A, float* B, float* C, int N) {
3
    int row = blockIdx.y * blockDim.y + threadIdx.y;
4
    int col = blockIdx.x * blockDim.x + threadIdx.x;
5

6
    if (row < N && col < N) {
7
        float sum = 0.0f;
8
        for (int k = 0; k < N; k++) {
9
            sum += A[row * N + k] * B[k * N + col];
10
        }
11
        C[row * N + col] = sum;
12
    }
13
}
14
// 性能：~1 GFLOPS（N=1024）
15
// 问题：全局内存访问 N 次，每次 ~400 ns

1
// 使用共享内存的 tiled 矩阵乘法
2
#define TILE_SIZE 16
3

4
__global__ void matmul_tiled(float* A, float* B, float* C, int N) {
5
    __shared__ float As[TILE_SIZE][TILE_SIZE];
6
    __shared__ float Bs[TILE_SIZE][TILE_SIZE];
7

8
    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
9
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
10

11
    float sum = 0.0f;
12

13
    for (int t = 0; t < N / TILE_SIZE; t++) {
14
        // 加载 tile 到共享内存
15
        As[threadIdx.y][threadIdx.x] = A[row * N + t * TILE_SIZE + threadIdx.x];
16
        Bs[threadIdx.y][threadIdx.x] = B[(t * TILE_SIZE + threadIdx.y) * N + col];
17

18
        __syncthreads();  // Block 内同步
19

20
        // 在共享内存中计算
21
        for (int k = 0; k < TILE_SIZE; k++) {
22
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
23
        }
24

25
        __syncthreads();
26
    }
27

28
    C[row * N + col] = sum;
29
}
30
// 性能：~50 GFLOPS（N=1024）
31
// 优化：共享内存减少全局内存访问次数

1
# Nsight Systems：宏观时间线分析
2
nsys profile --stats=true ./my_gpu_app
3
nsys report report1.qdrep          # 生成时间线报告
4

5
# Nsight Compute：微观 Kernel 剖析
6
ncu --set full --target-processes all ./my_gpu_app
7
ncu --launch-skip 10 --launch-count 1 ./my_gpu_app  # 跳过 warmup，只剖析第 11 次

1
-- GPU 加速的数据库查询
2
-- BlazingSQL / OmniDB 等项目
3

4
-- 传统：CPU 执行聚合
5
SELECT user_id, SUM(amount)
6
FROM orders
7
WHERE created_at >= '2026-01-01'
8
GROUP BY user_id;
9
-- 执行时间：~5 秒（1 亿行）
10

11
-- GPU 加速：将数据传输到 GPU 执行
12
-- 1. 数据从磁盘加载到 CPU 内存
13
-- 2. CPU 内存 → GPU 内存（PCIe，~12 GB/s）
14
-- 3. GPU 执行聚合（数千核心并行）
15
-- 4. 结果 GPU 内存 → CPU 内存
16
-- 执行时间：~0.5 秒（10x 加速）

1
GPU-Direct RDMA：
2
  传统：网卡 → CPU 内存 → GPU 内存（2 次拷贝）
3
  GPU-Direct：网卡 → GPU 内存（0 次额外拷贝）
4

5
  应用场景：
6
  - 高频交易：市场数据直接到 GPU 进行分析
7
  - 深度学习推理：网络接收数据直接到 GPU 处理
8
  - 视频处理：视频流直接到 GPU 编解码

1
// vecadd.cu — 完整的 CUDA 程序
2
#include <stdio.h>
3
#include <stdlib.h>
4

5
__global__ void vecAdd(const float* A, const float* B, float* C, int N) {
6
    int i = blockIdx.x * blockDim.x + threadIdx.x;
7
    if (i < N) C[i] = A[i] + B[i];
8
}
9

10
int main() {
11
    int N = 1 << 20;  // 1M 元素
12
    size_t bytes = N * sizeof(float);
13

14
    float *h_A = (float*)malloc(bytes);
15
    float *h_B = (float*)malloc(bytes);
16
    float *h_C = (float*)malloc(bytes);
17
    for (int i = 0; i < N; i++) { h_A[i] = 1.0f; h_B[i] = 2.0f; }
18

19
    float *d_A, *d_B, *d_C;
20
    cudaMalloc(&d_A, bytes);
21
    cudaMalloc(&d_B, bytes);
22
    cudaMalloc(&d_C, bytes);
23

24
    cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice);
25
    cudaMemcpy(d_B, h_B, bytes, cudaMemcpyHostToDevice);
26

27
    int blockSize = 256;
28
    int numBlocks = (N + blockSize - 1) / blockSize;
29
    vecAdd<<<numBlocks, blockSize>>>(d_A, d_B, d_C, N);
30

31
    cudaMemcpy(h_C, d_C, bytes, cudaMemcpyDeviceToHost);
32
    printf("h_C[0] = %f (expected 3.0)\n", h_C[0]);
33

34
    cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
35
    free(h_A); free(h_B); free(h_C);
36
    return 0;
37
}

1
# 编译 CUDA 程序
2
nvcc -O3 -o vecadd vecadd.cu
3

4
# 运行
5
./vecadd
6
# 输出：h_C[0] = 3.000000 (expected 3.0)
7

8
# 使用 Nsight Systems 分析时间线
9
nsys profile --stats=true ./vecadd
10
nsys report report1.qdrep
11

12
# 使用 Nsight Compute 剖析 Kernel
13
ncu --set basic ./vecadd
14
# 关注指标：Memory Throughput、Compute Throughput、Occupancy

1
# 使用 nvprof 快速查看 Kernel 执行时间（旧版工具）
2
nvprof ./vecadd
3
# 输出示例：
4
#   GPU activities:  99.3%  vecAdd  12.3us
5
#   API calls:      85.1%  cudaMalloc  320ms
6

7
# 对比不同 Block 大小的性能
8
for bs in 128 256 512 1024; do
9
    echo "Block size = $bs"
10
    ncu --metrics gpu__time_duration.sum ./vecadd
11
done