From f6e00dec139898865b8879efaebb4c375b58b72e Mon Sep 17 00:00:00 2001 From: Edgar Gabriel Date: Tue, 23 Aug 2022 19:08:38 +0000 Subject: [PATCH] introduce support for ncclFloat16/half in UT --- test/AllReduce_OutOfPlace.cpp | 2 +- test/AllToAllv_OutOfPlace.cpp | 2 +- test/Gather_InPlace.cpp | 2 +- test/ReduceScatter_InPlace.cpp | 4 ++-- test/ReduceScatter_OutOfPlace.cpp | 4 ++-- test/common/EnvVars.cpp | 3 +-- test/common/PtrUnion.cpp | 9 +++++++++ test/common/PtrUnion.hpp | 2 ++ 8 files changed, 19 insertions(+), 9 deletions(-) diff --git a/test/AllReduce_OutOfPlace.cpp b/test/AllReduce_OutOfPlace.cpp index 660d5eb1b8..6fc8aef9f4 100644 --- a/test/AllReduce_OutOfPlace.cpp +++ b/test/AllReduce_OutOfPlace.cpp @@ -13,7 +13,7 @@ namespace RcclUnitTesting // Configuration std::vector const funcTypes = {ncclCollAllReduce}; - std::vector const dataTypes = {ncclFloat32, ncclFloat64, ncclBfloat16}; + std::vector const dataTypes = {ncclFloat16, ncclFloat32, ncclFloat64, ncclBfloat16}; std::vector const redOps = {ncclSum, ncclMin}; std::vector const roots = {0}; std::vector const numElements = {1048576, 53327, 1024}; diff --git a/test/AllToAllv_OutOfPlace.cpp b/test/AllToAllv_OutOfPlace.cpp index f072558b30..e15cafdb3a 100644 --- a/test/AllToAllv_OutOfPlace.cpp +++ b/test/AllToAllv_OutOfPlace.cpp @@ -47,7 +47,7 @@ namespace RcclUnitTesting { TestBed testBed; // Configuration - std::vector const& dataTypes = {ncclInt32, ncclFloat64}; + std::vector const& dataTypes = {ncclInt32, ncclFloat64, ncclFloat16}; std::vector const numElements = {1048576, 53327, 1024}; bool const inPlace = false; bool const useManagedMem = false; diff --git a/test/Gather_InPlace.cpp b/test/Gather_InPlace.cpp index d9dec85ce1..6fe6f0e5a8 100644 --- a/test/Gather_InPlace.cpp +++ b/test/Gather_InPlace.cpp @@ -13,7 +13,7 @@ namespace RcclUnitTesting // Configuration std::vector const funcTypes = {ncclCollGather}; - std::vector const dataTypes = {ncclInt8, ncclInt32, ncclInt64}; + std::vector const dataTypes = {ncclInt8, ncclInt32, ncclInt64, ncclFloat16}; std::vector const redOps = {ncclSum}; std::vector const roots = {0}; std::vector const numElements = {1048576, 53327, 1024}; diff --git a/test/ReduceScatter_InPlace.cpp b/test/ReduceScatter_InPlace.cpp index 96cc283f8a..53f200a2ec 100644 --- a/test/ReduceScatter_InPlace.cpp +++ b/test/ReduceScatter_InPlace.cpp @@ -13,10 +13,10 @@ namespace RcclUnitTesting // Configuration std::vector const funcTypes = {ncclCollReduceScatter}; - std::vector const dataTypes = {ncclInt8, ncclInt32, ncclInt64}; + std::vector const dataTypes = {ncclInt8, ncclInt32, ncclInt64, ncclFloat16}; std::vector const redOps = {ncclSum, ncclProd}; std::vector const roots = {0}; - std::vector const numElements = {1048576, 53327, 1024}; + std::vector const numElements = {1048576, 53327, 5461, 1024}; std::vector const inPlaceList = {true}; std::vector const managedMemList = {false}; diff --git a/test/ReduceScatter_OutOfPlace.cpp b/test/ReduceScatter_OutOfPlace.cpp index 09d5868a49..c4f7ee6b99 100644 --- a/test/ReduceScatter_OutOfPlace.cpp +++ b/test/ReduceScatter_OutOfPlace.cpp @@ -13,10 +13,10 @@ namespace RcclUnitTesting // Configuration std::vector const funcTypes = {ncclCollReduceScatter}; - std::vector const dataTypes = {ncclFloat32, ncclFloat64, ncclBfloat16}; + std::vector const dataTypes = {ncclFloat16, ncclFloat32, ncclFloat64, ncclBfloat16}; std::vector const redOps = {ncclMin, ncclMax, ncclAvg}; std::vector const roots = {0}; - std::vector const numElements = {1048576, 53327, 1024}; + std::vector const numElements = {1048576, 53327, 5461, 1024}; std::vector const inPlaceList = {false}; std::vector const managedMemList = {false}; diff --git a/test/common/EnvVars.cpp b/test/common/EnvVars.cpp index a2e9fa8031..3e557cec77 100644 --- a/test/common/EnvVars.cpp +++ b/test/common/EnvVars.cpp @@ -95,8 +95,7 @@ namespace RcclUnitTesting dataTypes.push_back(ncclUint32); dataTypes.push_back(ncclInt64); dataTypes.push_back(ncclUint64); - // Half-precision floats disabled due to lack of host-side support - // dataTypes.push_back(ncclFloat16); + dataTypes.push_back(ncclFloat16); dataTypes.push_back(ncclFloat32); dataTypes.push_back(ncclFloat64); dataTypes.push_back(ncclBfloat16); diff --git a/test/common/PtrUnion.cpp b/test/common/PtrUnion.cpp index 82eec6a524..708b86d606 100644 --- a/test/common/PtrUnion.cpp +++ b/test/common/PtrUnion.cpp @@ -153,6 +153,7 @@ namespace RcclUnitTesting case ncclUint32: U4[idx] = valueI; break; case ncclInt64: I8[idx] = valueI; break; case ncclUint64: U8[idx] = valueI; break; + case ncclFloat16: F2[idx] = __float2half(static_cast(valueF)); break; case ncclFloat32: F4[idx] = valueF; break; case ncclFloat64: F8[idx] = valueF; break; case ncclBfloat16: B2[idx] = rccl_bfloat16(static_cast(valueF)); break; @@ -173,6 +174,7 @@ namespace RcclUnitTesting case ncclUint32: valueI = U4[idx]; break; case ncclInt64: valueI = I8[idx]; break; case ncclUint64: valueI = U8[idx]; break; + case ncclFloat16: valueF = __half2float(F2[idx]); break; case ncclFloat32: valueF = F4[idx]; break; case ncclFloat64: valueF = F8[idx]; break; case ncclBfloat16: valueF = B2[idx]; break; @@ -202,6 +204,7 @@ namespace RcclUnitTesting case ncclUint32: U4[idx] *= scalarsPerRank.U4[rank]; break; case ncclInt64: I8[idx] *= scalarsPerRank.I8[rank]; break; case ncclUint64: U8[idx] *= scalarsPerRank.U8[rank]; break; + case ncclFloat16: F2[idx] = __float2half(__half2float(F2[idx]) * __half2float(scalarsPerRank.F2[rank])); break; case ncclFloat32: F4[idx] *= scalarsPerRank.F4[rank]; break; case ncclFloat64: F8[idx] *= scalarsPerRank.F8[rank]; break; case ncclBfloat16: B2[idx] *= scalarsPerRank.B2[rank]; break; @@ -234,6 +237,7 @@ namespace RcclUnitTesting case ncclUint32: U4[idx] = ReduceOp(op, U4[idx], inputCpu.U4[idx]); break; case ncclInt64: I8[idx] = ReduceOp(op, I8[idx], inputCpu.I8[idx]); break; case ncclUint64: U8[idx] = ReduceOp(op, U8[idx], inputCpu.U8[idx]); break; + case ncclFloat16: F2[idx] = __float2half(ReduceOp(op, __half2float(F2[idx]), __half2float(inputCpu.F2[idx]))); break; case ncclFloat32: F4[idx] = ReduceOp(op, F4[idx], inputCpu.F4[idx]); break; case ncclFloat64: F8[idx] = ReduceOp(op, F8[idx], inputCpu.F8[idx]); break; case ncclBfloat16: B2[idx] = ReduceOp(op, B2[idx], inputCpu.B2[idx]); break; @@ -260,6 +264,7 @@ namespace RcclUnitTesting case ncclUint32: U4[idx] /= divisor; break; case ncclInt64: I8[idx] /= divisor; break; case ncclUint64: U8[idx] /= divisor; break; + case ncclFloat16: F2[idx] = __float2half(__half2float(F2[idx])/divisor); break; case ncclFloat32: F4[idx] /= divisor; break; case ncclFloat64: F8[idx] /= divisor; break; case ncclBfloat16: B2[idx] = (rccl_bfloat16((float)(B2[idx]) / divisor)); break; @@ -289,6 +294,7 @@ namespace RcclUnitTesting case ncclUint32: isMatch = (U4[idx] == expected.U4[idx]); break; case ncclInt64: isMatch = (I8[idx] == expected.I8[idx]); break; case ncclUint64: isMatch = (U8[idx] == expected.U8[idx]); break; + case ncclFloat16: isMatch = (fabs(__half2float(F2[idx]) - __half2float(expected.F2[idx])) < 9e-2); break; case ncclFloat32: isMatch = (fabs(F4[idx] - expected.F4[idx]) < 1e-5); break; case ncclFloat64: isMatch = (fabs(F8[idx] - expected.F8[idx]) < 1e-12); break; case ncclBfloat16: isMatch = (fabs((float)B2[idx] - (float)expected.B2[idx]) < 9e-2); break; @@ -315,6 +321,8 @@ namespace RcclUnitTesting ERROR("Expected output: %ld. Actual output: %ld at index %lu\n", expected.I8[idx], I8[idx], idx); break; case ncclUint64: ERROR("Expected output: %lu. Actual output: %lu at index %lu\n", expected.U8[idx], U8[idx], idx); break; + case ncclFloat16: + ERROR("Expected output: %f. Actual output: %f at index %lu\n", __half2float(expected.F2[idx]), __half2float(F2[idx]), idx); break; case ncclFloat32: ERROR("Expected output: %f. Actual output: %f at index %lu\n", expected.F4[idx], F4[idx], idx); break; case ncclFloat64: @@ -343,6 +351,7 @@ namespace RcclUnitTesting case ncclUint32: ss << U4[i]; break; case ncclInt64: ss << I8[i]; break; case ncclUint64: ss << U8[i]; break; + case ncclFloat16: ss << __half2float(F2[i]); break; case ncclFloat32: ss << F4[i]; break; case ncclFloat64: ss << F8[i]; break; case ncclBfloat16: ss << (float)B2[i]; break; diff --git a/test/common/PtrUnion.hpp b/test/common/PtrUnion.hpp index 67323e93d5..79c10ed979 100644 --- a/test/common/PtrUnion.hpp +++ b/test/common/PtrUnion.hpp @@ -2,6 +2,7 @@ #include "ErrCode.hpp" #include "rccl/rccl.h" #include "rccl_bfloat16.h" +#include "hip/hip_fp16.h" namespace RcclUnitTesting { @@ -35,6 +36,7 @@ namespace RcclUnitTesting uint32_t* U4; // ncclUint32 int64_t* I8; // ncclInt64 uint64_t* U8; // ncclUint64 + __half* F2; // ncclFloat16 float* F4; // ncclFloat32 double* F8; // ncclFloat64 rccl_bfloat16* B2; // ncclBfloat16