GPU 内存架构
┌─────────────────────────────────────────┐
│ GPU 显存 │
├─────────────────────────────────────────┤
│ ┌─────────────┐ ┌─────────────┐ │
│ │ Global │ │ Shared │ │
│ │ Memory │ │ Memory │ │
│ │ (VRAM) │ │ (SMEM) │ │
│ └─────────────┘ └─────────────┘ │
│ │
│ ┌─────────────┐ ┌─────────────┐ │
│ │ Constant │ │ Local │ │
│ │ Memory │ │ Memory │ │
│ └─────────────┘ └─────────────┘ │
└─────────────────────────────────────────┘
↓ ↑
CPU (通过 PCIe) GPU 计算单元
内存类型对比
| 内存类型 |
位置 |
延迟 |
大小 |
用途 |
| Global |
VRAM |
高 |
大 |
输入/输出数据 |
| Shared |
SMEM |
低 |
小 |
线程块内通信 |
| Constant |
缓存 |
中 |
中 |
只读数据 |
| Local |
寄存器/VRAM |
高 |
小 |
线程私有 |
| Register |
SM |
最低 |
极小 |
线程私有 |
CUDA 内存管理 (rust-cuda)
use cuda_sys::ffi::*;
let mut d_ptr: *mut f32 = std::ptr::null_mut();
unsafe {
cudaMalloc(&mut d_ptr as *mut *mut f32, size * std::mem::size_of::<f32>())
};
unsafe {
cudaMemcpy(
d_ptr as *mut c_void,
h_ptr as *const c_void,
size * std::mem::size_of::<f32>(),
cudaMemcpyHostToDevice
);
};
let mut h_result: Vec<f32> = vec![0.0; size];
unsafe {
cudaMemcpy(
h_result.as_mut_ptr() as *mut c_void,
d_ptr as *const c_void,
size * std::mem::size_of::<f32>(),
cudaMemcpyDeviceToHost
);
};
unsafe {
cudaFree(d_ptr as *mut c_void);
};
零拷贝内存
let mut h_ptr: *mut f32 = std::ptr::null_mut();
unsafe {
cudaMallocHost(&mut h_ptr as *mut *mut f32, size * std::mem::size_of::<f32>())
};
let stream: cudaStream_t = std::ptr::null_mut();
unsafe {
cudaMemcpyAsync(
d_ptr as *mut c_void,
h_ptr as *const c_void,
size * std::mem::size_of::<f32>(),
cudaMemcpyHostToDevice,
stream
);
};
unsafe {
cudaStreamSynchronize(stream);
};
统一内存 (Unified Memory)
let mut unified_ptr: *mut f32 = std::ptr::null_mut();
unsafe {
cudaMallocManaged(&mut unified_ptr as *mut *mut f32, size * std::mem::size_of::<f32>());
};
unsafe {
for i in 0..size {
*unified_ptr.add(i) = i as f32;
}
};
launch_kernel(unified_ptr, size);
unsafe {
println!("Result: {}", unified_ptr.add(0).read());
};
unsafe {
cudaFree(unified_ptr as *mut c_void);
};
内存合并访问
__global__ void bad_access(float* data) {
int idx = threadIdx.x + blockIdx.x * 32;
float value = data[idx * 32];
}
__global__ void coalesced_access(float* data) {
int idx = threadIdx.x + blockIdx.x * blockDim.x;
float value = data[idx];
}
共享内存使用
__global__ void shared_memory_reduce(float* input, float* output) {
__shared__ float sdata[256];
int tid = threadIdx.x;
int idx = blockIdx.x * blockDim.x + threadIdx.x;
sdata[tid] = input[idx];
__syncthreads();
for (int s = blockDim.x / 2; s > 0; s >>= 1) {
if (tid < s) {
sdata[tid] += sdata[tid + s];
}
__syncthreads();
}
if (tid == 0) {
output[blockIdx.x] = sdata[0];
}
}
内存对齐
const size_t ALIGNMENT = 256;
struct alignas(256) AlignedData {
float4 position;
float4 normal;
};
assert(((uintptr_t)ptr % ALIGNMENT) == 0);
性能优化检查表
| 优化项 |
检查点 |
| 内存合并 |
线程访问连续内存 |
| 共享内存 |
减少全局内存访问 |
| 内存对齐 |
256 字节对齐 |
| 异步操作 |
计算与传输重叠 |
| 固定内存 |
使用页锁定内存 |
| 批处理 |
减少内核启动开销 |
与其他技能关联
rust-gpu
│
├─► rust-performance → 性能优化
├─► rust-unsafe → 底层内存操作
└─► rust-embedded → no_std 设备