From b5306a69c5c3821bb5cc44cbb50adfa76059fab7 Mon Sep 17 00:00:00 2001 From: David Addison Date: Sat, 1 Feb 2025 16:10:41 -0800 Subject: [PATCH] Also check NCCL runtime version version for fp8/bf16 support --- src/common.cu | 30 ++++++++++++++++++------------ 1 file changed, 18 insertions(+), 12 deletions(-) diff --git a/src/common.cu b/src/common.cu index 9b613e1..8819baa 100644 --- a/src/common.cu +++ b/src/common.cu @@ -956,11 +956,15 @@ int main(int argc, char* argv[]) { if (NCCL_VERSION_CODE >= NCCL_VERSION(2,11,0) && test_ncclVersion >= NCCL_VERSION(2,11,0)) { test_opnum++; // PreMulSum } - #if HAVE_BF16 + #if defined(__CUDA_BF16_TYPES_EXIST__) + if (NCCL_VERSION_CODE >= NCCL_VERSION(2,10,0) && test_ncclVersion >= NCCL_VERSION(2,10,0)) { test_typenum++; // bfloat16 + } #endif - #if HAVE_FP8 + #if defined(__CUDA_FP8_TYPES_EXIST__) + if (NCCL_VERSION_CODE >= NCCL_VERSION(2,24,0) && test_ncclVersion >= NCCL_VERSION(2,24,0)) { test_typenum += 2; // fp8 e4m3,e5m2 + } #endif #endif @@ -1275,18 +1279,20 @@ testResult_t run() { #ifdef MPI_SUPPORT MPI_Allreduce(MPI_IN_PLACE, &minCudaArch, 1, MPI_INT, MPI_MIN, MPI_COMM_WORLD); #endif -#if HAVE_FP8 - if (minCudaArch < 900) { // Filter out fp8 on pre-Hopper hardware - int n = 0; - for (int i=0; i < test_typenum; i++) { - if (!(test_types[i] == ncclFloat8e4m3 || test_types[i] == ncclFloat8e5m2)) { - test_types[n] = test_types[i]; - test_typenames[n] = test_typenames[i]; - n += 1; +#if defined(__CUDA_FP8_TYPES_EXIST__) + if (NCCL_VERSION_CODE >= NCCL_VERSION(2,24,0) && test_ncclVersion >= NCCL_VERSION(2,24,0)) { + if (minCudaArch < 900) { // Filter out fp8 on pre-Hopper hardware + int n = 0; + for (int i=0; i < test_typenum; i++) { + if (!(test_types[i] == ncclFloat8e4m3 || test_types[i] == ncclFloat8e5m2)) { + test_types[n] = test_types[i]; + test_typenames[n] = test_typenames[i]; + n += 1; + } } + test_typenum = n; } - test_typenum = n; - }; + } #endif //if parallel init is not selected, use main thread to initialize NCCL