TensorCore是NVIDIA GPU中专用于加速矩阵乘法与张量运算的硬件单元，自Volta架构起存在，需通过cuBLASLt或WMMA API等CUDA生态接口调用，纯C++无法直接触发。
TensorCore 是什么，为什么不能直接用 C++ 调用
TensorCore 不是 CPU 指令集，也不是 C++ 标准库或 STL 提供的接口；它是 NVIDIA GPU 上的专用硬件单元，只在特定架构（Volta 及之后，如 Turing、Ampere、Hopper）中存在，且必须通过 CUDA 生态调用。你写纯 C++ 代码（哪怕开了 -O3 或 -ffast-math）完全无法触发 TensorCore——编译器根本看不到它。
真正能调度 TensorCore 的路径只有两条：cuBLASLt（推荐）或 WMMA API（底层可控但复杂）。前者封装了 GEMM 的 TensorCore 启用逻辑，后者需要手写 warp-level 矩阵分块、load/store、mma.sync 指令序列。
用 cuBLASLt 做半精度 GEMM：关键参数和易错点
cuBLASLt 是目前最实用的选择，它会在运行时自动判断是否启用 TensorCore（比如当 computeType == CUBLASLT_COMPUTE_16F 且矩阵尺寸满足 tile 对齐要求时）。但默认行为不保证用 TensorCore——你得显式设置 heuristics 并检查结果。
- 输入矩阵必须是
half（即__half），不是float16_t或自定义 struct；需用cudaMalloc分配显存，并用cudaMemcpy传输 - 调用前必须设置
cublasLtMatmulHeuristicResult_t：先调cublasLtMatmulPreferenceSetAttribute开启CUBLASLT_MATMUL_PREF_REDUCTION_SCHEME，再用cublasLtMatmulHeuristic获取候选算法——其中heuristicResult.algoId >= 0且heuristicResult.workspaceSize > 0才表示 TensorCore 路径被选中 - 常见错误：
CUBLAS_STATUS_NOT_SUPPORTED多因 m/n/k 未对齐（Ampere 要求 m/n/k 都是 8 的倍数，Hopper 要求 16）；CUBLAS_STATUS_INVALID_VALUE常因ldA/ldB/ldC（leading dimension）没设对（例如ldA = m但实际分配的是按行优先的m x k，则应为ldA = k）
WMMA API 手动启用 TensorCore：什么时候值得折腾
除非你要做 kernel fusion（比如 GEMM + bias + gelu 一体化）、定制量化策略，或者调试 TensorCore 利用率，否则别碰 WMMA。它暴露的是 warp-level 接口，每个 wmma::fragments 必须严格匹配 shape（如 wmma::fragment<:matrix_a wmma::row_major wmma::half></:matrix_a>），且 load/store 必须用 wmma::load_matrix_sync，计算必须用 wmma::mma_sync。
立即学习“C++免费学习笔记（深入）”；
典型陷阱：
- host 端数据类型必须用
__half，且内存 layout 必须是 column-major（WMMA 默认按列加载 A/B），否则load_matrix_sync会读错数据 - 所有 fragment 操作必须在同一个 warp 内完成，不能跨 block 或依赖 global sync；
__syncthreads()对 fragment 无效，要用wmma::fill_fragment初始化 - 没有自动 padding：若 m=31，则最后一个 warp 的 fragment 会越界——你得手动 zero-pad 到 16/32 倍数，或加条件分支跳过非法 lane
验证 TensorCore 是否真在跑：别信文档，看指标
即使 cuBLASLt 返回成功，也不代表 TensorCore 被用上。最可靠的方式是用 nsys profile 抓 trace：
- 过滤
volta_sgemm_32x32或ampere_hgemm_16816这类 kernel 名——带hgemm且含16816（表示 16×8×16 tile）的才是 TensorCore 路径 - 看
SOL% (Speed of Light)：TensorCore GEMM 通常 SOL > 70%，而模拟 FP16 的 FP32 GEMM 一般 - 用
nvidia-smi dmon -s u观察sm__sass_thread_inst_executed_op_tensor计数器是否显著增长
一个常被忽略的细节：TensorCore 加速效果高度依赖访存带宽。如果你的 half 数据反复从 global memory 搬入搬出（没用 shared memory 做 tiling），实际加速比可能不到 2×，甚至不如优化过的 cublasHgemm。
