深入解析NVIDIA 3090 FP64算力：架构优势与科学计算实战

1次阅读

共计 2271 个字符，预计需要花费 6 分钟才能阅读完成。

NVIDIA RTX 3090 采用的 GA102 芯片在 FP64（双精度浮点）设计上与上一代 Turing 架构有显著差异。通过 SM（Streaming Multiprocessor）结构对比可以看到 Ampere 的精妙取舍：

1:64 的 FP32/FP64 比例：每个 SM 内仅有 1 个 FP64 计算单元对应 64 个 FP32 单元，这种设计直接反映在官方 556 TFLOPS FP32 vs 0.87 TFLOPS FP64 的算力数据上
双功能 ALU 设计：FP64 单元与 INT32 单元共享硬件资源，当执行 FP64 指令时会暂时挂起 INT32 运算能力
显存子系统优化：配备 936GB/ s 带宽的 GDDR6X 显存和 ECC 纠错机制，缓解 FP64 计算对数据搬运的高需求

（示意图：红框标注 FP64 单元在 SM 中的位置）

以下是通过共享内存优化 bank conflict 的 FP64 GEMM 核心代码（CUDA 12.2 环境）：

__global__ void fp64MatMul(const double* A, const double* B, double* C, 
                          int M, int N, int K) {__shared__ double sA[TILE][TILE];  // 使用方阵分块减少 bank 冲突
    __shared__ double sB[TILE][TILE+1]; // 添加 padding 避免 bank conflict

    int tx = threadIdx.x, ty = threadIdx.y;
    int row = blockIdx.y * TILE + ty;
    int col = blockIdx.x * TILE + tx;
    double sum = 0.0;

    for (int i = 0; i < K; i += TILE) {
        // 合作式加载到共享内存
        sA[ty][tx] = (row < M && i+tx < K) ? A[row*K + i+tx] : 0.0;
        sB[ty][tx] = (i+ty < K && col < N) ? B[(i+ty)*N + col] : 0.0;
        __syncthreads();

        // 计算 tile 内乘积
        for (int k = 0; k < TILE; k++)
            sum += sA[ty][k] * sB[k][tx];
        __syncthreads();}

    if (row < M && col < N)
        C[row*N + col] = sum;
}

关键优化点说明：

TILE_SIZE 建议设为 32 的倍数（如 32/64），与 warp 调度器宽度匹配
sB 矩阵添加 1 列 padding 彻底消除 bank conflict（GDDR6X 的 bank 宽度为 32 字节）
使用 __syncthreads() 确保正确的内存可见性

虽然 3090 的 Tensor Core 主要面向 FP16/FP32，但通过 cuBLASLt 仍可实现 FP64 加速：

cublasLtMatmulDesc_t operationDesc;
cublasLtMatrixLayout_t Adesc, Bdesc, Cdesc;

// 创建 TF32 计算描述符
cublasLtMatmulDescCreate(&operationDesc, CUBLAS_COMPUTE_32F_FAST_TF32, CUDA_R_32F);

// 显式设置数学模式
cublasLtMatmulDescSetAttribute(operationDesc, CUBLASLT_MATMUL_DESC_MATH_MODE, 
                              &mathMode, sizeof(mathMode));

// 执行混合精度矩阵乘
cublasLtMatmul(ltHandle, operationDesc,
               &alpha, A, Adesc, B, Bdesc,
               &beta, C, Cdesc, C, Cdesc,
               &algo, workspace, workspaceSize, stream);

实现原理：