当前位置：首页 > news >正文

Blackwell 和 Hopper 架构的 GPGPU 新功能全面综述

news 2025/8/17 7:11:28

NVIDIA 的 Blackwell 和 Hopper 架构代表了 GPU 计算的最新进展，引入了多项创新功能来加速 AI 和高性能计算工作负载。本文将深入探讨这些架构的关键新特性，包括 Tensor Memory Access (TMA)、跨 SM 的共享内存访问、FP8 计算支持等，并提供可运行的代码示例。

1. Tensor Memory Access (TMA)

TMA 是 Hopper 架构引入的一项重要功能，它允许高效地在全局内存和共享内存之间传输张量数据块。

TMA 技术解析

TMA 的主要优势包括：

单条指令即可描述整个张量传输

支持 1D-5D 的张量传输

自动处理边界条件

与异步拷贝引擎协同工作，实现计算与数据传输重叠

TMA 代码示例

#include <cuda.h>
#include <cuda_runtime.h>
#include <stdio.h>// 简单的 TMA 示例 (需要 CUDA 12+ 和 Hopper 架构)
__global__ void tma_example(float* global_data) {// 声明共享内存__shared__ float shared_tile[32][32];// 创建 TMA 描述符 (实际代码会更复杂)uint64_t tma_desc;// 这里应该填充 TMA 描述符的各个字段// 执行 TMA 拷贝asm volatile ("cp.async.bulk.tensor.2d.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], [%2], %3;":: "r"(shared_tile), "r"(tma_desc), "r"(global_data), "r"(32*32*sizeof(float)));// 等待 TMA 完成asm volatile ("mbarrier.arrive.bulk.complete_tx::bytes %0, %1;" :: "r"(shared_tile), "r"(32*32*sizeof(float)));// 使用共享内存数据...
}int main() {float* d_data;cudaMalloc(&d_data, 1024*1024*sizeof(float));tma_example<<<1, 32>>>(d_data);cudaDeviceSynchronize();cudaFree(d_data);return 0;
}

2. 跨 SM 的共享内存访问 (Cluster Shared Memory)

Hopper 架构引入了集群概念，允许同一集群内的不同 SM 直接访问彼此的共享内存。

技术解析

集群由多个 SM 组成 (通常 2-8 个)

集群内的 SM 可以低延迟访问彼此的共享内存

需要显式启用集群模式

使用 __shared__ 修饰符和新的寻址模式

代码示例

#include <cuda.h>
#include <cuda_runtime.h>
#include <stdio.h>// 集群共享内存示例
__global__ void cluster_shared_mem_example() {// 声明集群共享内存__shared__ int cluster_shared[1024];// 只有集群中的第一个块初始化共享内存if (clusterIdx.x == 0 && threadIdx.x == 0) {for (int i = 0; i < 1024; i++) {cluster_shared[i] = i;}}// 集群同步__syncthreads_cluster();// 所有集群内的块都可以访问相同的数据if (threadIdx.x == 0) {printf("Block %d sees cluster_shared[10] = %d\n", blockIdx.x, cluster_shared[10]);}
}int main() {// 启动一个包含2个块的集群cluster_shared_mem_example<<<dim3(2,1,1), 32>>>();cudaDeviceSynchronize();return 0;
}

3. FP8 计算支持

Blackwell 和 Hopper 都增强了 FP8 支持，特别适合 AI 训练和推理。

FP8 技术解析

两种 FP8 格式: E4M3 (4位指数, 3位尾数) 和 E5M2 (5位指数, 2位尾数)

硬件加速的 FP8 矩阵乘积累加运算

与 Tensor Core 紧密集成

自动或手动缩放保持精度

FP8 代码示例

#include <cuda.h>
#include <cuda_runtime.h>
#include <cuda_fp8.h>
#include <stdio.h>// FP8 矩阵乘法示例
__global__ void fp8_matmul(const __nv_fp8_e4m3* a, const __nv_fp8_e4m3* b, float* c, int m, int n, int k) {using namespace nvcuda;// 每个线程块处理一个矩阵块int row = blockIdx.y * blockDim.y + threadIdx.y;int col = blockIdx.x * blockDim.x + threadIdx.x;if (row < m && col < n) {float sum = 0.0f;for (int i = 0; i < k; ++i) {// 将FP8转换为float进行计算float a_val = __half2float(__nv_cvt_fp8_to_halfraw(a[row * k + i], __NV_E4M3));float b_val = __half2float(__nv_cvt_fp8_to_halfraw(b[i * n + col], __NV_E4M3));sum += a_val * b_val;}c[row * n + col] = sum;}
}int main() {const int m = 32, n = 32, k = 32;// 分配和初始化FP8矩阵__nv_fp8_e4m3 *d_a, *d_b;float *d_c;cudaMalloc(&d_a, m*k*sizeof(__nv_fp8_e4m3));cudaMalloc(&d_b, k*n*sizeof(__nv_fp8_e4m3));cudaMalloc(&d_c, m*n*sizeof(float));// 初始化代码省略...dim3 blocks((n+15)/16, (m+15)/16);dim3 threads(16, 16);fp8_matmul<<<blocks, threads>>>(d_a, d_b, d_c, m, n, k);cudaDeviceSynchronize();cudaFree(d_a);cudaFree(d_b);cudaFree(d_c);return 0;
}

4. 其他重要新功能

4.1 动态并行增强

Blackwell 改进了动态并行能力，允许更灵活的内核启动模式。

// 动态并行示例
__global__ void dynamic_parallel_kernel(int depth, int max_depth) {if (depth >= max_depth) return;printf("Depth %d, thread %d\n", depth, threadIdx.x);if (threadIdx.x == 0 && depth < max_depth) {dynamic_parallel_kernel<<<1, 32>>>(depth+1, max_depth);}// 不需要显式同步，设备运行时自动处理
}int main() {dynamic_parallel_kernel<<<1, 32>>>(0, 3);cudaDeviceSynchronize();return 0;
}

4.2 增强的 Tensor Core

Blackwell 引入了新一代 Tensor Core，支持更广泛的精度和操作。

// Tensor Core 示例 (使用 WMMA API)
#include <cuda.h>
#include <cuda_runtime.h>
#include <cuda_fp16.h>
#include <mma.h>
#include <stdio.h>using namespace nvcuda;__global__ void wmma_example(half *a, half *b, float *c, int m, int n, int k) {// 声明矩阵分片wmma::fragment<wmma::matrix_a, 16, 16, 16, half, wmma::row_major> a_frag;wmma::fragment<wmma::matrix_b, 16, 16, 16, half, wmma::col_major> b_frag;wmma::fragment<wmma::accumulator, 16, 16, 16, float> c_frag;// 初始化累加器wmma::fill_fragment(c_frag, 0.0f);// 加载和计算for (int i = 0; i < k; i += 16) {wmma::load_matrix_sync(a_frag, a + blockIdx.y * m * k + i, k);wmma::load_matrix_sync(b_frag, b + i * n + blockIdx.x * 16, n);wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);}// 存储结果wmma::store_matrix_sync(c + blockIdx.y * m * n + blockIdx.x * 16, c_frag, n, wmma::mem_row_major);
}int main() {const int m = 64, n = 64, k = 64;half *d_a, *d_b;float *d_c;cudaMalloc(&d_a, m*k*sizeof(half));cudaMalloc(&d_b, k*n*sizeof(half));cudaMalloc(&d_c, m*n*sizeof(float));// 初始化代码省略...dim3 gridDim(n/16, m/16);dim3 blockDim(32, 1);wmma_example<<<gridDim, blockDim>>>(d_a, d_b, d_c, m, n, k);cudaDeviceSynchronize();cudaFree(d_a);cudaFree(d_b);cudaFree(d_c);return 0;
}