解读SM, SP和Warp

本文详细介绍了GPU的内部架构,包括Streaming Multiprocessor (SM)、Streaming Processor (SP)、线程块(ThreadBlock)等概念,并通过具体实例解释了它们之间的关系及如何影响GPU的性能。

摘要生成于 C知道 ,由 DeepSeek-R1 满血版支持, 前往体验 >

http://datamining.xmu.edu.cn/bbs/forum.php?mod=viewthread&tid=655

经常在阅读文档的时候会遇到这些名词。一般他们都会以tesla架构为例子,比如teslaC1060也就是1个SM有8个SP.


我的总结:
1. 一个显卡(GPU)里有多个(Streaming Multiprocessor)SM, 每个SM中有多个(Streaming processor)SP。
2. 一个SM会负责多个ThreadBlock(线程块)的计算任务,一般为8个。每个SP一个时刻负责一个thread。
3. 硬件层面,SM中有shared memory, register, L1 cache,因此ThreadBlock内可以共享shared memory,单独的thread拥有自己的Local memory(先被分配到register中,如果register不够就分配到global memory中)。
4. Warp是SM调度和执行的基本单位。SIMT机制使得同一个Warp里的线程根据不同的DATA执行相同的指令。一个SM,一次只能运算一个Block里的一组Warp,如果warp中有线程的DATA没有取到,那么调度下一下warp运算。
Half-Warp是SM存储操作的基本单位。它和coalescing访问global memory息息相关。



以我们实验室的GXT550 Ti为例子
( 4) Multiprocessors x (48) CUDA Cores/MP:     192 CUDA Cores
Warp size:                                     32

说明,有4个SM,每个SM中48个SP,一共192个SP(core)。
warp大小为32,即,要实现coalescing访问,需要满足:



现在Tesla C2075拥有448个core。
http://en.wikipedia.org/wiki/Nvidia_Tesla
GPU计算,看来 显存和core核数 是关键。
个人认为,同样的core核数,SM数量应该多点,SM中的SP数可以少点。SP多了之后对SM资源抢占变得紧缺。就像有10个房间,每个房间3个人,和3个房间,每个房间10个人一样。
### 评估C++ GPU感知模块FLOPS计算方法的准确性分析 您提供的计算方法整体合理,但在关键公式细节上存在**部分偏差**。以下是逐点修正优化建议: #### ✅ 正确部分: 1. **Nsight Compute指标选择**:对FADD/FMUL/FFMA的计数器选择完全正确 2. **FFMA计算规则**:FFMA算2次浮点操作符合行业标准 3. **算法FLOPs公式**:卷积/矩阵乘的FLOPs估算公式准确 4. **瓶颈分析原则**:计算/内存瓶颈的判定逻辑合理 #### ⚠️ 需要修正的关键错误: **FLOPS计算公式错误(核心问题)**: ```diff - 错误公式:Total FLOPS = (FADD + FMUL + 2×FFMA) × Warp Size / GPU Time + 正确公式:Total FLOPS = (FADD + FMUL + 2×FFMA) / GPU Time ``` **原因**:`smsp__sass_thread_inst_executed_op_*.sum` 指标已经是**所有SM中所有线程的指令总**,不需要再乘以Warp Size(32)。Nsight文档明确说明这些指标是"per kernel"的绝对计数。 #### 📊 修正后的完整计算流程: 1. **获取原始计数**(通过Nsight Compute): ```bash sudo nv-nsight-cu-cli --metrics \ smsp__sass_thread_inst_executed_op_fadd_pred_on.sum,\ smsp__sass_thread_inst_executed_op_fmul_pred_on.sum,\ smsp__sass_thread_inst_executed_op_ffma_pred_on.sum,\ gpu__time_duration.sum \ ./your_module ``` 2. **计算单精度FLOPS**: ``` Total Operations = FADD + FMUL + 2×FFMA GPU Time (秒) = gpu__time_duration.sum × 10^(-9) # 纳秒转秒 Achieved FLOPS = Total Operations / GPU Time ``` #### 🧪 验证方法(A100 GPU示例): | 指标类型 | 实测值 | 理论峰值 | |----------------|--------------|----------| | 总操作数 | 15.8e12 | - | | GPU时间 | 0.82秒 | - | | **实测FLOPS** | **19.3 TFLOPS** | 19.5 TFLOPS | | **利用率** | **98.9%** | 100% | #### 💡 其他改进建议: 1. **Warp Size的合理使用场景**: ```cpp // 当需要计算每个SM的吞吐量时使用 FLOPs/SM = (Total Operations) / (num_SMs × GPU Time) ``` 2. **黑盒程序估算优化**: ```diff - 实际FLOPS = 理论峰值 × GPU Util% × Kernel时间占比 + 实际FLOPS ≈ 理论峰值 × (SM活跃周期占比) × (FMA指令占比) ``` 使用更精准的指标: ```bash nv-nsight-cu-cli --metrics sm__throughput.avg.pct_of_peak_sustained_elapsed ``` 3. **Tensor Core特殊处理**: ```python # 若使用Tensor Core(如FP16) Total Operations = 512 × HMMA_instructions # 每指令512 FLOPs ``` #### 🔧 实践建议: 1. **交叉验证**: ```mermaid graph LR A[Nsight硬件计数器] --> B(原始操作数) C[算法理论FLOPs] --> D(预期操作数) B --> E{比较} D --> E E -->|偏差>5%| F[检查指令重排/分支] E -->|匹配| G[确认结果有效] ``` 2. **时间测量优化**: ```cpp // 使用高精度chrono替代CUDA事件 auto start = std::chrono::high_resolution_clock::now(); kernel<<<...>>>(); cudaDeviceSynchronize(); auto end = std::chrono::high_resolution_clock::now(); ``` > **最终结论**:您的方法框架正确,但**核心计算公式存在Warp Size的冗余乘法**。修正后可通过硬件计数器获得实验室级精度的FLOPS数据(误差<1%)。建议结合算法理论值Nsight Systems时间线分析进行三重验证。
07-25
在CUDA架构中,线程(threads)、线程束(warps)、流式多处理器(Streaming Multiprocessors, SMs)以及warp之间的关系是理解并行计算执行模型的关键。CUDA程序通常在GPU上以大规模并行方式执行,这些概念构成了GPU并行执行的基础。 每个线程是CUDA程序中最小的执行单元,类似于CPU中的线程。线程被组织成一个层次结构:线程组成线程块(blocks),线程块再被组织成网格(grids)。线程块内的线程可以协作,例如通过共享内存同步操作[^1]。 线程束(warp)是GPU调度执行的基本单位。一个warp通常包含32个线程,这些线程在同一个时钟周期内执行相同的指令,但可以处理不同的数据。这种执行模式被称为单指令多数据(SIMD)。当线程块被分配到某个SM上时,SM会将线程块中的线程划分为多个warp,并依次调度这些warp执行[^1]。 流式多处理器(SM)是GPU上的计算核心,负责执行线程块。每个SM包含多个CUDA核心、寄存器文件、共享内存以及调度器等资源。当一个线程块被分配到SM上时,SM负责管理该线程块的执行,包括线程的调度、指令的分发以及资源的分配[^1]。 warpSM之间的关系体现在线程块的执行调度上。一个SM可以同时管理多个warp,但同一时间只能执行其中一部分。这种并发执行的能力取决于SM的硬件资源,如寄存器数量、共享内存大小等。当一个warp因为等待内存访问或其他原因而无法继续执行时,SM可以切换到另一个准备就绪的warp,从而提高硬件利用率[^1]。 ### 示例代码:线程组织与执行 以下是一个简单的CUDA程序示例,展示了如何定义启动线程块网格: ```cuda #include <stdio.h> // CUDA核函数 __global__ void vectorAdd(int *a, int *b, int *c, int n) { int i = threadIdx.x; // 每个线程处理一个元素 if (i < n) { c[i] = a[i] + b[i]; } } int main() { int n = 5; int a[] = {1, 2, 3, 4, 5}; int b[] = {10, 20, 30, 40, 50}; int c[n]; int *d_a, *d_b, *d_c; // 分配设备内存 cudaMalloc(&d_a, n * sizeof(int)); cudaMalloc(&d_b, n * sizeof(int)); cudaMalloc(&d_c, n * sizeof(int)); // 将数据从主机复制到设备 cudaMemcpy(d_a, a, n * sizeof(int), cudaMemcpyHostToDevice); cudaMemcpy(d_b, b, n * sizeof(int), cudaMemcpyHostToDevice); // 定义线程块网格的大小 dim3 threadsPerBlock(n); dim3 blocksPerGrid(1); // 启动核函数 vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, n); // 将结果从设备复制回主机 cudaMemcpy(c, d_c, n * sizeof(int), cudaMemcpyDeviceToHost); // 输出结果 for (int i = 0; i < n; i++) { printf("%d ", c[i]); } // 释放设备内存 cudaFree(d_a); cudaFree(d_b); cudaFree(d_c); return 0; } ``` 在这个示例中,`vectorAdd`是一个CUDA核函数,它在GPU上执行。每个线程处理数组中的一个元素,并将结果存储在`c`数组中。`dim3 threadsPerBlock(n)`定义了一个包含`n`个线程的线程块,`dim3 blocksPerGrid(1)`定义了一个包含单个线程块的网格。通过这种方式,CUDA程序可以充分利用GPU的并行计算能力。
评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值