Merge pull request #603 from edgargabriel/topic/float16_unit_tests
introduce support for ncclFloat16/half in UT
This commit is contained in:
Edgar Gabriel
GitHub
@@ -13,7 +13,7 @@ namespace RcclUnitTesting
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollAllReduce};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclFloat32, ncclFloat64, ncclBfloat16};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclFloat16, ncclFloat32, ncclFloat64, ncclBfloat16};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum, ncclMin};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {1048576, 53327, 1024};
|
||||
|
||||
@@ -47,7 +47,7 @@ namespace RcclUnitTesting
|
||||
{
|
||||
TestBed testBed;
|
||||
// Configuration
|
||||
std::vector<ncclDataType_t> const& dataTypes = {ncclInt32, ncclFloat64};
|
||||
std::vector<ncclDataType_t> const& dataTypes = {ncclInt32, ncclFloat64, ncclFloat16};
|
||||
std::vector<int> const numElements = {1048576, 53327, 1024};
|
||||
bool const inPlace = false;
|
||||
bool const useManagedMem = false;
|
||||
|
||||
@@ -13,7 +13,7 @@ namespace RcclUnitTesting
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollGather};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclInt8, ncclInt32, ncclInt64};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclInt8, ncclInt32, ncclInt64, ncclFloat16};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {1048576, 53327, 1024};
|
||||
|
||||
@@ -13,10 +13,10 @@ namespace RcclUnitTesting
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollReduceScatter};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclInt8, ncclInt32, ncclInt64};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclInt8, ncclInt32, ncclInt64, ncclFloat16};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclSum, ncclProd};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {1048576, 53327, 1024};
|
||||
std::vector<int> const numElements = {1048576, 53327, 5461, 1024};
|
||||
std::vector<bool> const inPlaceList = {true};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
|
||||
|
||||
@@ -13,10 +13,10 @@ namespace RcclUnitTesting
|
||||
|
||||
// Configuration
|
||||
std::vector<ncclFunc_t> const funcTypes = {ncclCollReduceScatter};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclFloat32, ncclFloat64, ncclBfloat16};
|
||||
std::vector<ncclDataType_t> const dataTypes = {ncclFloat16, ncclFloat32, ncclFloat64, ncclBfloat16};
|
||||
std::vector<ncclRedOp_t> const redOps = {ncclMin, ncclMax, ncclAvg};
|
||||
std::vector<int> const roots = {0};
|
||||
std::vector<int> const numElements = {1048576, 53327, 1024};
|
||||
std::vector<int> const numElements = {1048576, 53327, 5461, 1024};
|
||||
std::vector<bool> const inPlaceList = {false};
|
||||
std::vector<bool> const managedMemList = {false};
|
||||
|
||||
|
||||
@@ -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);
|
||||
|
||||
@@ -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<float>(valueF)); break;
|
||||
case ncclFloat32: F4[idx] = valueF; break;
|
||||
case ncclFloat64: F8[idx] = valueF; break;
|
||||
case ncclBfloat16: B2[idx] = rccl_bfloat16(static_cast<float>(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;
|
||||
|
||||
@@ -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
|
||||
|
||||
مرجع در شماره جدید
Block a user