Skip to content

06_device_attributes.md

Latest commit

 

History

History
250 lines (185 loc) · 9.72 KB

06_device_attributes.md

File metadata and controls

250 lines (185 loc) · 9.72 KB
Raw

GPU 原子操作与 PCIe 能力查询

基于 RTX 5090 (Blackwell, CC 12.0) 和 A100-SXM4-80GB (Ampere, CC 8.0) 双平台实测。cudaDeviceGetAttribute 可查询 100+ 种设备属性，本文聚焦 PCIe 原子操作、Host Native Atomic 等底层硬件能力的查询方法。

验证代码：本文所有实测数据来自配套工具：06_device_attributes.cu。在任意 GPU 上编译运行即可查看完整的硬件能力清单。

1. cudaDeviceGetAttribute 概述

cudaDeviceProp 结构体只暴露了最常用的属性（SM 数量、显存大小、时钟频率等）。cudaDeviceGetAttribute 能查询更底层的硬件能力：能力类属性多为 0/1 开关，规格类属性是数值，它们直接决定某些 CUDA API 是否可用。

CUDA 高级功能在不同 GPU 上的支持状态不同：

  • GPU atomic 直写 Host 内存：依赖 HostNativeAtomicSupported——文档定义是「设备与主机之间的链路支持原生原子操作」，为 0 时不能用 GPU atomic 直接操作 Host 内存
  • Unified Memory：ConcurrentManagedAccess 决定 Host 和 GPU 能否同时访问同一块 managed memory
  • Cooperative Launch：CooperativeLaunch 决定能否使用 CUDA Cooperative Groups

正确使用模式是运行时能力检查 + 优雅降级：每次运行前查询，并准备好降级路径，不靠文档假设。

int value;
cudaError_t err = cudaDeviceGetAttribute(&value, cudaDevAttrXxx, device_id);

2. PCIe 原子操作能力

2.1 什么是 PCIe Atomic

PCIe 原子操作允许 GPU 通过 PCIe 总线对 Host 内存执行单事务原子读写（无需软件加锁）。RDMA 和 GPUDirect 场景依赖这条路径。

2.2 查询原子操作支持

cat > check_atomic.cu << 'EOF'
#include <cuda_runtime.h>
#include <stdio.h>

#define CHECK_ATTR(attr, id) do {                              \
    int v;                                                      \
    cudaDeviceGetAttribute(&v, attr, id);                      \
    printf("  %-45s = %d%s\n", #attr, v, v ? " ✓" : "");        \
} while(0)

int main() {
    int count;
    cudaGetDeviceCount(&count);

    for (int i = 0; i < count; i++) {
        cudaDeviceProp prop;
        cudaGetDeviceProperties(&prop, i);
        printf("GPU %d: %s (CC %d.%d)\n", i, prop.name,
               prop.major, prop.minor);

        printf("\n--- PCIe & Atomic ---\n");
        CHECK_ATTR(cudaDevAttrHostNativeAtomicSupported, i);
        CHECK_ATTR(cudaDevAttrCanUseHostPointerForRegisteredMem, i);

        printf("\n--- Managed Memory ---\n");
        CHECK_ATTR(cudaDevAttrConcurrentManagedAccess, i);
        CHECK_ATTR(cudaDevAttrPageableMemoryAccess, i);
        CHECK_ATTR(cudaDevAttrPageableMemoryAccessUsesHostPageTables, i);
        CHECK_ATTR(cudaDevAttrDirectManagedMemAccessFromHost, i);

        printf("\n--- Launch Capability ---\n");
        CHECK_ATTR(cudaDevAttrCooperativeLaunch, i);
        CHECK_ATTR(cudaDevAttrCooperativeMultiDeviceLaunch, i);

        printf("\n--- SM Resources ---\n");
        CHECK_ATTR(cudaDevAttrMaxRegistersPerMultiprocessor, i);
        CHECK_ATTR(cudaDevAttrMaxBlocksPerMultiprocessor, i);
        CHECK_ATTR(cudaDevAttrMaxThreadsPerMultiProcessor, i);
        CHECK_ATTR(cudaDevAttrAsyncEngineCount, i);

        printf("\n--- Precision ---\n");
        CHECK_ATTR(cudaDevAttrSingleToDoublePrecisionPerfRatio, i);
    }
    return 0;
}
EOF

nvcc -o check_atomic check_atomic.cu && ./check_atomic

RTX 5090 实际输出：

GPU 0: NVIDIA GeForce RTX 5090 (CC 12.0)

--- PCIe & Atomic ---
  cudaDevAttrHostNativeAtomicSupported     = 0
  cudaDevAttrCanUseHostPointerForRegisteredMem = 1 ✓

--- Managed Memory ---
  cudaDevAttrConcurrentManagedAccess       = 1 ✓
  cudaDevAttrPageableMemoryAccess          = 0
  cudaDevAttrPageableMemoryAccessUsesHostPageTables = 0
  cudaDevAttrDirectManagedMemAccessFromHost = 0

--- Launch Capability ---
  cudaDevAttrCooperativeLaunch             = 1 ✓
  cudaDevAttrCooperativeMultiDeviceLaunch  = 1 ✓

--- SM Resources ---
  cudaDevAttrMaxRegistersPerMultiprocessor = 65536
  cudaDevAttrMaxBlocksPerMultiprocessor    = 24
  cudaDevAttrMaxThreadsPerMultiProcessor   = 1536
  cudaDevAttrAsyncEngineCount              = 2

--- Precision ---
  cudaDevAttrSingleToDoublePrecisionPerfRatio = 64

A100-SXM4-80GB 实际输出（2026-05 实测）：

GPU 0: NVIDIA A100-SXM4-80GB (CC 8.0)

--- PCIe & Atomic ---
  cudaDevAttrHostNativeAtomicSupported          = 0
  cudaDevAttrCanUseHostPointerForRegisteredMem  = 1 ✓

--- Managed Memory ---
  cudaDevAttrConcurrentManagedAccess            = 1 ✓
  cudaDevAttrPageableMemoryAccess               = 0
  cudaDevAttrPageableMemoryAccessUsesHostPageTables = 0
  cudaDevAttrDirectManagedMemAccessFromHost     = 0

--- Launch Capability ---
  cudaDevAttrCooperativeLaunch                  = 1 ✓

--- SM Resources ---
  cudaDevAttrMaxRegistersPerMultiprocessor      = 65536
  cudaDevAttrMaxBlocksPerMultiprocessor         = 32
  cudaDevAttrMaxThreadsPerMultiProcessor        = 2048
  cudaDevAttrAsyncEngineCount                   = 3

--- Precision ---
  cudaDevAttrSingleToDoublePrecisionPerfRatio   = 2

2.3 关键解读

属性 RTX 5090 A100 数据中心 GPU 影响
HostNativeAtomicSupported 0 ❌ 0 ❌ 1 ✅ GPU atomic 无法直接操作 Host 内存
CanUseHostPointerForRegisteredMem 1 ✅ 1 ✅ 1 ✅ 可以使用 registered host memory
PageableMemoryAccess 0 ❌ 0 ❌ 1 ✅ (H100+) 为 0 时设备不能直接访问 pageable 内存
DirectManagedMemAccessFromHost 0 ❌ 0 ❌ 1 ✅ (GH200) 为 0 时 Host 访问 managed 页需迁移
ConcurrentManagedAccess 1 ✅ 1 ✅ 1 ✅ 支持 Host+GPU 同时访问 managed memory
CooperativeMultiDeviceLaunch 1 ✅ 未采样 — 配套脚本会打印，表中以待运行实测为准
SingleToDoublePrecisionPerfRatio 64:1 2:1 2:1 (A100/H100) 消费级 GPU 双精度性能严重受限

注意：A100 上 HostNativeAtomicSupported 和 PageableMemoryAccess 均为 0（与某些资料中「数据中心 GPU = 1」的预期不一致）。这说明这两个能力并非由「数据中心 vs 消费级」决定，而是特定型号的硬件特性。代码实践中始终用 cudaDeviceGetAttribute 做运行时能力检查，不查表假设。

3. nvidia-smi 侧确认 PCIe Atomic

nvidia-smi -q 的 PCI 部分会显示 atomic 能力：

nvidia-smi -q | grep -A2 "Atomic"

RTX 5090 输出：

Atomic Caps Outbound    : N/A
Atomic Caps Inbound     : FETCHADD_32 FETCHADD_64 SWAP_32 SWAP_64 CAS_32 CAS_64
  • Inbound：GPU 接受来自 PCIe 的 atomic 请求，此处支持 FetchAdd/Swap/CAS 32/64
  • Outbound：GPU 向 Host 发起 atomic 请求，此处 N/A（消费级 GPU 不支持）

4. 常用属性速查表

4.1 内存相关

属性 说明 RTX 5090
cudaDevAttrMaxSharedMemoryPerBlockOptin 可选最大 shared memory 待实测（运行配套脚本可得）
cudaDevAttrMemoryPoolsSupported 内存池支持 支持
cudaDevAttrUnifiedAddressing 统一寻址 支持

4.2 执行模型

属性 说明 RTX 5090 A100
cudaDevAttrMaxRegistersPerMultiprocessor SM 寄存器总数 65536 65536
cudaDevAttrMaxBlocksPerMultiprocessor SM 最大 block 数 24 32
cudaDevAttrMaxThreadsPerMultiProcessor SM 最大线程数 1536 2048
cudaDevAttrWarpSize (from prop) Warp 大小 32 32

4.3 流与并发

属性 RTX 5090 A100 说明
cudaDevAttrAsyncEngineCount 2 3 异步引擎数量（通常即 copy engine）
cudaDevAttrConcurrentKernels (prop) 1 1 支持 concurrent kernel 执行
cudaDevAttrMaxSurface1DLayeredLayers 待实测 待实测 1D 分层 surface 的层数上限

5. 编程实践

5.1 能力检查模式

// 运行时检查能力，优雅降级
int host_atomic;
cudaDeviceGetAttribute(&host_atomic,
    cudaDevAttrHostNativeAtomicSupported, 0);

if (host_atomic) {
    // 使用 GPU atomic 直接操作 Host 内存
    launch_kernel_with_host_atomic<<<...>>>();
} else {
    // 回退方案：先 D2H → CPU atomic → H2D
    launch_kernel_without_host_atomic<<<...>>>();
    cudaMemcpy(...);
    cpu_atomic_op(...);
    cudaMemcpy(...);
}

5.2 设备选择

// 选择支持特定能力的 GPU
for (int i = 0; i < count; i++) {
    int v;
    cudaDeviceGetAttribute(&v, cudaDevAttrCooperativeLaunch, i);
    if (v) {
        cudaSetDevice(i);
        break;
    }
}

参考