cuda-basics

CUDA编程基础

概述

CUDA (Compute Unified Device Architecture) 是NVIDIA推出的并行计算平台和编程模型。

线程层次结构

三层结构

Grid (网格)
  └─ Block (块) 
      └─ Thread (线程)

维度表示

// 1D
dim3 block(256);
dim3 grid((N + 255) / 256);

// 2D
dim3 block(16, 16);
dim3 grid((M + 15) / 16, (N + 15) / 16);

// 3D
dim3 block(8, 8, 8);
dim3 grid((M + 7) / 8, (N + 7) / 8, (K + 7) / 8);

线程索引计算

// 1D索引
int idx = blockIdx.x * blockDim.x + threadIdx.x;

// 2D索引
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;

// 全局1D索引（从2D）
int idx = row * width + col;

内存层次

内存类型对比

内存类型	位置	访问速度	大小	作用域	生命周期
Register	片上	最快	~64KB/SM	Thread	Thread
Shared Memory	片上	快	48KB-164KB	Block	Block
L1 Cache	片上	快	128KB	-	-
L2 Cache	片上	中	MB级	-	-
Global Memory	DRAM	慢	GB级	Grid	Application
Constant Memory	DRAM	中(有cache)	64KB	Grid	Application
Texture Memory	DRAM	中(有cache)	-	Grid	Application

声明方式

// Register (自动)
int local_var;

// Shared Memory
__shared__ float shared_data[256];

// Global Memory
__global__ void kernel(float* global_data) { }

// Constant Memory
__constant__ float const_data[1024];

内存访问优化

1. 合并内存访问 (Coalesced Access)

// ✅ 好：连续访问
__global__ void coalesced_read(float* data) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    float val = data[idx];  // 线程连续访问
}

// ❌ 差：跨步访问
__global__ void strided_read(float* data, int stride) {
    int idx = (blockIdx.x * blockDim.x + threadIdx.x) * stride;
    float val = data[idx];  // 线程跨步访问
}

2. 使用Shared Memory

__global__ void use_shared_memory(float* input, float* output) {
    __shared__ float tile[TILE_SIZE];
    
    int tid = threadIdx.x;
    int gid = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 从全局内存加载到共享内存
    tile[tid] = input[gid];
    __syncthreads();  // 同步
    
    // 从共享内存读取（快速）
    float val = tile[tid];
    
    // 处理...
    output[gid] = val;
}