AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Switched to using the hip_fp8 header instead of rccl_float8, resolving compatibility issues.(#109)

Browse Source 
* addressing hip_fp8 support compatibility issue

* skipping mulsum and avg test for fp8, using hip_fp8 for product

* syncing with nccl-tests

removing the fp8 filter for pre-hopper gpus and resolving the merge conflict

---------

Co-authored-by: Marzieh Berenjkoub <mberenjk@amd.com>

[ROCm/rccl-tests commit: 4b2b635766]
This commit is contained in:
mberenjk
2025-05-14 15:30:07 -05:00
committed by GitHub
parent fe47d3dd77
commit ed6ebb12a7
8 changed files with 385 additions and 193 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rccl-tests/src/all_reduce.cu
+2 -3
View File
@@ -65,7 +65,6 @@ testResult_t AllReduceRunTest(struct threadArgs* args, int root, ncclDataType_t
ncclRedOp_t *run_ops;
const char **run_typenames, **run_opnames;
int type_count, op_count;
if ((int)type != -1) {
type_count = 1;
run_types = &type;
@@ -89,8 +88,8 @@ testResult_t AllReduceRunTest(struct threadArgs* args, int root, ncclDataType_t
for (int i=0; i<type_count; i++) {
for (int j=0; j<op_count; j++) {
#if defined(RCCL_FLOAT8)
if((run_types[i] == ncclFp8E4M3 || run_types[i] == ncclFp8E5M2) && run_ops[j] == ncclProd)
continue;
if((run_types[i] == ncclFloat8e4m3 || run_types[i] == ncclFloat8e5m2) && (run_ops[j] == ncclProd || run_ops[j] == ncclAvg || strcmp(run_opnames[j],"mulsum") == 0))
continue;
#endif
TESTCHECK(TimeTest(args, run_types[i], run_typenames[i], run_ops[j], run_opnames[j], -1));
}
projects/rccl-tests/src/common.cu
+32 -8
View File
@@ -31,6 +31,13 @@ int test_ncclVersion = 0; // init'd with ncclGetVersion()
int32_t gpu_block3;
size_t cache_bytes = 192 * 1024 * 1024; // Use 192MB
// RCCL_FLOAT8 support
bool rccl_float8_useFnuz = false;
bool IsArchMatch(char const* arch, char const* target) {
// helper function to reduce clutter in code elsewhere. Returns true on match.
return (strncmp(arch, target, strlen(target)) == 0);
}
#if NCCL_MAJOR >= 2
ncclDataType_t test_types[ncclNumTypes] = {
ncclInt8, ncclUint8, ncclInt32, ncclUint32, ncclInt64, ncclUint64, ncclHalf, ncclFloat, ncclDouble
@@ -38,7 +45,7 @@ size_t cache_bytes = 192 * 1024 * 1024; // Use 192MB
, ncclBfloat16
#endif
#if RCCL_FLOAT8 == 1
, ncclFp8E4M3, ncclFp8E5M2
, ncclFloat8e4m3, ncclFloat8e5m2
#endif
};
const char *test_typenames[ncclNumTypes] = {
@@ -196,6 +203,7 @@ void Reporter::addResult(int gpusPerRank, int ranksPerNode, int totalRanks, size
}
bool Reporter::isMainThread() { return is_main_thread == 1; }
static int minCudaArch = 1<<30;
#define NUM_BLOCKS 32
@@ -304,18 +312,18 @@ static bool minReqVersion(int rmajor, int rminor, int rpatch)
}
testResult_t CheckDelta(void* results, void* expected, size_t count, size_t offset, ncclDataType_t type, ncclRedOp_t op, uint64_t seed, int nranks, int64_t *wrongEltN) {
ncclVerifiableVerify(results, expected, count, (int)type, (int)op, nranks, seed, offset, wrongEltN, cudaStreamDefault);
CUDACHECK(ncclVerifiableVerify(results, expected, count, (int)type, (int)op, nranks, seed, offset, wrongEltN, cudaStreamDefault));
CUDACHECK(cudaDeviceSynchronize());
return testSuccess;
}
testResult_t InitDataReduce(void* data, const size_t count, const size_t offset, ncclDataType_t type, ncclRedOp_t op, uint64_t seed, int nranks) {
ncclVerifiablePrepareExpected(data, count, (int)type, (int)op, nranks, seed, offset, cudaStreamDefault);
CUDACHECK(ncclVerifiablePrepareExpected(data, count, (int)type, (int)op, nranks, seed, offset, cudaStreamDefault));
return testSuccess;
}
testResult_t InitData(void* data, const size_t count, size_t offset, ncclDataType_t type, ncclRedOp_t op, uint64_t seed, int nranks, int rank) {
ncclVerifiablePrepareInput(data, count, (int)type, (int)op, nranks, rank, seed, offset, cudaStreamDefault);
CUDACHECK(ncclVerifiablePrepareInput(data, count, (int)type, (int)op, nranks, rank, seed, offset, cudaStreamDefault));
return testSuccess;
}
@@ -563,8 +571,8 @@ testResult_t startColl(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t
case ncclBfloat16: bf16 = ncclVerifiablePremulScalar<hip_bfloat16>(rank); break;
#endif
#if defined(RCCL_FLOAT8)
case ncclFp8E4M3: fp8_e4m3 = ncclVerifiablePremulScalar<rccl_float8>(rank); break;
case ncclFp8E5M2: fp8_e5m2 = ncclVerifiablePremulScalar<rccl_bfloat8>(rank); break;
case ncclFloat8e4m3: fp8_e4m3 = ncclVerifiablePremulScalar<rccl_float8>(rank); break;
case ncclFloat8e5m2 : fp8_e5m2 = ncclVerifiablePremulScalar<rccl_bfloat8>(rank); break;
#endif
case ncclNumTypes: break;
}
@@ -1330,6 +1338,13 @@ testResult_t run() {
char hostname[1024];
getHostName(hostname, 1024);
hipDeviceProp_t devProp;
CUDACHECK(hipGetDeviceProperties(&devProp, 0));
if (IsArchMatch(devProp.gcnArchName, "gfx942")) {
PRINT("On gfx942 architecture, using FNUZ FP8 types");
rccl_float8_useFnuz = true;
}
#ifdef MPI_SUPPORT
MPI_Comm_size(MPI_COMM_WORLD, &totalProcs);
MPI_Comm_rank(MPI_COMM_WORLD, &proc);
@@ -1456,12 +1471,21 @@ testResult_t run() {
gpus[i] = ((gpu0 != -1 ? gpu0 : localRank*nThreads*nGpus) + i)%numDevices;
CUDACHECK(cudaSetDevice(gpus[i]));
TESTCHECK(AllocateBuffs(sendbuffs.data()+i, sendBytes, recvbuffs.data()+i, recvBytes, expected.data()+i, (size_t)maxBytes));
if (streamnull)
if (streamnull) {
streams[i] = NULL;
else
}
else {
CUDACHECK(cudaStreamCreateWithFlags(streams.data()+i, cudaStreamNonBlocking));
}
int archMajor, archMinor;
CUDACHECK(cudaDeviceGetAttribute(&archMajor, cudaDevAttrComputeCapabilityMajor, gpus[i]));
CUDACHECK(cudaDeviceGetAttribute(&archMinor, cudaDevAttrComputeCapabilityMinor, gpus[i]));
minCudaArch = std::min(minCudaArch, 100*archMajor + 10*archMinor);
}
#ifdef MPI_SUPPORT
MPI_Allreduce(MPI_IN_PLACE, &minCudaArch, 1, MPI_INT, MPI_MIN, MPI_COMM_WORLD);
#endif
//if parallel init is not selected, use main thread to initialize NCCL
ncclComm_t* comms = (ncclComm_t*)malloc(sizeof(ncclComm_t)*nThreads*nGpus);
#if NCCL_VERSION_CODE >= NCCL_VERSION(2,19,0)
projects/rccl-tests/src/common.h
+2 -2
View File
@@ -258,8 +258,8 @@ static size_t wordSize(ncclDataType_t type) {
//case ncclInt8:
case ncclUint8:
#if NCCL_MAJOR >= 2 && RCCL_FLOAT8 == 1
case ncclFp8E4M3:
case ncclFp8E5M2:
case ncclFloat8e4m3:
case ncclFloat8e5m2:
#endif
#endif
return 1;
projects/rccl-tests/src/rccl_float8.h
+73 -19
View File
@@ -24,8 +24,9 @@
#define ROCBLAS_FLOAT8_H
#include <stdint.h>
#include <hip/hip_version.h>
#if __cplusplus < 201103L || (!defined(__HCC__) && !defined(__HIPCC__))
#if __cplusplus < 201103L || (!defined(__HIP_PLATFORM_AMD__) && !defined(__HIPCC__))
/*! \brief Struct to represent a 8 bit floating-point number. */
typedef struct
@@ -38,7 +39,60 @@ typedef struct
uint8_t data;
} rccl_bfloat8;
#else // __cplusplus < 201103L || (!defined(__HCC__) && !defined(__HIPCC__))
// __cplusplus < 201103L || (!defined(__HIP_PLATFORM_AMD__) && !defined(__HIPCC__))
#elif HIP_VERSION >= 60200000
#include <hip/hip_fp8.h>
#if __HIP_DEVICE_COMPILE__ && (defined(__gfx950__) || defined(__gfx1200__) || defined(__gfx1201__) || (defined(__gfx1100__) || defined(__gfx1101__)))//HIP_FP8_TYPE_OCP is enabled.
typedef __hip_fp8_e4m3 rccl_float8;
typedef __hip_fp8_e5m2 rccl_bfloat8;
#elif __HIP_DEVICE_COMPILE__ && (defined(__gfx942__))
typedef __hip_fp8_e4m3_fnuz rccl_float8;
typedef __hip_fp8_e5m2_fnuz rccl_bfloat8;
#else
typedef __hip_fp8_e4m3 rccl_float8;
typedef __hip_fp8_e5m2 rccl_bfloat8;
#endif
#if __HIP_DEVICE_COMPILE__
inline std::ostream& operator<<(std::ostream& os, const rccl_float8& f8)
{
return os << float(f8);
}
inline std::ostream& operator<<(std::ostream& os, const rccl_bfloat8& bf8)
{
return os << float(bf8);
}
#else
inline std::ostream& operator<<(std::ostream& os, const __hip_fp8_e4m3& f8)
{
return os << float(f8);
}
inline std::ostream& operator<<(std::ostream& os, const __hip_fp8_e5m2& bf8)
{
return os << float(bf8);
}
//adding support for those operators on the host side
inline std::ostream& operator<<(std::ostream& os, const __hip_fp8_e4m3_fnuz& f8)
{
return os << float(f8);
}
inline std::ostream& operator<<(std::ostream& os, const __hip_fp8_e5m2_fnuz& bf8)
{
return os << float(bf8);
}
#endif
extern bool rccl_float8_useFnuz;
// For older versions of ROCm that do not include hip_fp8.h,
// we provide a local version of the header file as a fallback.
#else
#define HIP_HOST_DEVICE __host__ __device__
#define HIP_HOST __host__
@@ -344,7 +398,7 @@ struct rccl_float8
// default constructor
HIP_HOST_DEVICE rccl_float8() = default;
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
#if defined(__gfx942__) || defined(__gfx950__)
// device specific optimized F8 down-conversion code
template <bool stochastic_rounding = false>
@@ -381,10 +435,10 @@ struct rccl_float8
return i8data;
}
#endif // __gfx940__
#endif // __gfx942__
// constructor from float
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
#if defined(__gfx942__) || defined(__gfx950__)
// NOTE: ON-DEVICE... always optimal bias
explicit HIP_DEVICE rccl_float8(float v,
@@ -402,7 +456,7 @@ struct rccl_float8
// Host only implementation using s/w simulation
explicit HIP_HOST
#else
// both Host and DEVICE for non-gfx940 using s/w simulation
// both Host and DEVICE for non-gfx942 using s/w simulation
explicit HIP_HOST_DEVICE
#endif
rccl_float8(float v,
@@ -446,7 +500,7 @@ struct rccl_float8
}
// convert to float
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
#if defined(__gfx942__) || defined(__gfx950__)
// upcast using device specific intrinsic
explicit inline HIP_DEVICE operator float() const
{
@@ -460,7 +514,7 @@ struct rccl_float8
}
explicit inline HIP_HOST operator float() const
#else // non gfx940
#else // non gfx942
explicit inline HIP_HOST_DEVICE operator float() const
#endif
{
@@ -511,7 +565,7 @@ struct rccl_bfloat8
// default constructor
HIP_HOST_DEVICE rccl_bfloat8() = default;
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
#if defined(__gfx942__) || defined(__gfx950__)
// device specific optimized F8 down-conversion code
template <bool stochastic_rounding = false>
@@ -548,10 +602,10 @@ struct rccl_bfloat8
return i8data;
}
#endif // __gfx940__
#endif // __gfx942__
// constructor from float
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
#if defined(__gfx942__) || defined(__gfx950__)
// NOTE: ON-DEVICE... always optimal bias
explicit HIP_DEVICE rccl_bfloat8(float v,
@@ -569,7 +623,7 @@ struct rccl_bfloat8
// Host only implementation using s/w simulation
explicit HIP_HOST
#else
// both Host and DEVICE for non-gfx940 using s/w simulation
// both Host and DEVICE for non-gfx942 using s/w simulation
explicit HIP_HOST_DEVICE
#endif
rccl_bfloat8(float v,
@@ -613,7 +667,7 @@ struct rccl_bfloat8
}
// convert to float
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
#if defined(__gfx942__) || defined(__gfx950__)
// upcast using device specific intrinsic
explicit inline HIP_DEVICE operator float() const
{
@@ -627,7 +681,7 @@ struct rccl_bfloat8
}
explicit inline HIP_HOST operator float() const
#else // non gfx940
#else // non gfx942
explicit inline HIP_HOST_DEVICE operator float() const
#endif
{
@@ -969,7 +1023,7 @@ inline __host__ __device__ T explicit_downcast(Ta a, uint32_t rng = 0)
return a;
}
// Use h/w intrinsic and optimized version when __gfx940__
// Use h/w intrinsic and optimized version when __gfx942__
template <
typename T,
typename Ta,
@@ -980,7 +1034,7 @@ template <
= 0>
inline __host__ __device__ T explicit_downcast(Ta a, uint32_t rng)
{
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
#if defined(__gfx942__) || defined(__gfx950__)
// NOTE: we are directly calling cast_to_f8_from_f32 instead of constructor to optimize away one runtime branch
T val;
if(std::is_same<T, rccl_float8>::value)
@@ -988,12 +1042,12 @@ inline __host__ __device__ T explicit_downcast(Ta a, uint32_t rng)
else
val.data = rccl_bfloat8::cast_to_bf8_from_f32<stochastic_rounding>(float(a), rng);
return val;
#else // non gfx940
#else // non gfx942
return T(float(a),
stochastic_rounding ? T::rocblas_hip_f8_rounding_mode::stochastic
: T::rocblas_hip_f8_rounding_mode::standard,
rng);
#endif // __gfx940__
#endif // __gfx942__
}
// NOTE NOTE: The above code is good if we don't consider HIP-GEMM code and only consider the quantization
@@ -1016,6 +1070,6 @@ inline __host__ __device__ T explicit_downcast(Ta a, uint32_t rng)
// =================================================================================================
#endif // __cplusplus < 201103L || (!defined(__HCC__) && !defined(__HIPCC__))
#endif
#endif // ROCBLAS_FLOAT8_H
projects/rccl-tests/src/reduce.cu
+2 -2
View File
@@ -96,8 +96,8 @@ testResult_t ReduceRunTest(struct threadArgs* args, int root, ncclDataType_t typ
for (int i=0; i<type_count; i++) {
for (int j=0; j<op_count; j++) {
#if defined(RCCL_FLOAT8)
if((run_types[i] == ncclFp8E4M3 || run_types[i] == ncclFp8E5M2) && run_ops[j] == ncclProd)
continue;
if((run_types[i] == ncclFloat8e4m3 || run_types[i] == ncclFloat8e5m2) && (run_ops[j] == ncclProd || run_ops[j] == ncclAvg || strcmp(run_opnames[j],"mulsum") == 0))
continue;
#endif
for (int k=begin_root; k<=end_root; k++) {
TESTCHECK(TimeTest(args, run_types[i], run_typenames[i], run_ops[j], run_opnames[j], k));
projects/rccl-tests/src/reduce_scatter.cu
+2 -2
View File
@@ -91,8 +91,8 @@ testResult_t ReduceScatterRunTest(struct threadArgs* args, int root, ncclDataTyp
for (int i=0; i<type_count; i++) {
for (int j=0; j<op_count; j++) {
#if defined(RCCL_FLOAT8)
if((run_types[i] == ncclFp8E4M3 || run_types[i] == ncclFp8E5M2) && run_ops[j] == ncclProd)
continue;
if((run_types[i] == ncclFloat8e4m3 || run_types[i] == ncclFloat8e5m2) && (run_ops[j] == ncclProd || run_ops[j] == ncclAvg || strcmp(run_opnames[j],"mulsum") == 0))
continue;
#endif
TESTCHECK(TimeTest(args, run_types[i], run_typenames[i], run_ops[j], run_opnames[j], -1));
}