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

Switching to old version of rccl_float8 for ROCm versions earlier than 6.3 for backward compatibility. (#128)

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

[ROCm/rccl-tests commit: 9076091602]
This commit is contained in:
mberenjk
2025-05-16 09:14:46 -05:00
committed by GitHub
parent fa5259894c
commit db5ab33461
2 changed files with 8 additions and 11 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rccl-tests/src/rccl_float8.h
+1 -1
View File
@@ -40,7 +40,7 @@ typedef struct
} rccl_bfloat8; } rccl_bfloat8;
// __cplusplus < 201103L || (!defined(__HIP_PLATFORM_AMD__) && !defined(__HIPCC__)) // __cplusplus < 201103L || (!defined(__HIP_PLATFORM_AMD__) && !defined(__HIPCC__))
#elif HIP_VERSION >= 60200000 #elif HIP_VERSION >= 60300000
#include <hip/hip_fp8.h> #include <hip/hip_fp8.h>
projects/rccl-tests/verifiable/verifiable.cu
+7 -10
View File
@@ -392,7 +392,7 @@ struct FloatLayout<hip_bfloat16> {
}; };
#endif #endif
#if RCCL_FLOAT8 == 1 #if RCCL_FLOAT8 == 1
#if __HIP_DEVICE_COMPILE__ #if __HIP_DEVICE_COMPILE__ || HIP_VERSION < 60300000
template<> template<>
struct FloatLayout<rccl_float8> { struct FloatLayout<rccl_float8> {
static constexpr bool is_floating_point = true; static constexpr bool is_floating_point = true;
@@ -993,11 +993,10 @@ cudaError_t prepareInput1(
#if HAVE_ncclBfloat16 #if HAVE_ncclBfloat16
case ncclBfloat16: fn = (void const*)&prepareInput2<hip_bfloat16, ReduceOp>; break; case ncclBfloat16: fn = (void const*)&prepareInput2<hip_bfloat16, ReduceOp>; break;
#endif #endif
#if HAVE_ncclfp8_DEVICE #if HAVE_ncclfp8_DEVICE || HIP_VERSION < 60300000
case ncclFloat8e4m3: fn = (void const*)&prepareInput2<rccl_float8, ReduceOp>; break; case ncclFloat8e4m3: fn = (void const*)&prepareInput2<rccl_float8, ReduceOp>; break;
case ncclFloat8e5m2: fn = (void const*)&prepareInput2<rccl_bfloat8, ReduceOp>; break; case ncclFloat8e5m2: fn = (void const*)&prepareInput2<rccl_bfloat8, ReduceOp>; break;
#endif #elif HAVE_ncclfp8_HOST
#if HAVE_ncclfp8_HOST
case ncclFloat8e4m3: if (rccl_float8_useFnuz) { fn = (void const*)&prepareInput2<__hip_fp8_e4m3_fnuz, ReduceOp>; break;} case ncclFloat8e4m3: if (rccl_float8_useFnuz) { fn = (void const*)&prepareInput2<__hip_fp8_e4m3_fnuz, ReduceOp>; break;}
else { fn = (void const*)&prepareInput2<__hip_fp8_e4m3, ReduceOp>; break;} else { fn = (void const*)&prepareInput2<__hip_fp8_e4m3, ReduceOp>; break;}
case ncclFloat8e5m2: if (rccl_float8_useFnuz) { fn = (void const*)&prepareInput2<__hip_fp8_e5m2_fnuz, ReduceOp>; break;} case ncclFloat8e5m2: if (rccl_float8_useFnuz) { fn = (void const*)&prepareInput2<__hip_fp8_e5m2_fnuz, ReduceOp>; break;}
@@ -1084,11 +1083,10 @@ cudaError_t prepareExpected1(
#if HAVE_ncclBfloat16 #if HAVE_ncclBfloat16
case ncclBfloat16: fn = (void const*)&prepareExpected2<hip_bfloat16, ReduceOp>; break; case ncclBfloat16: fn = (void const*)&prepareExpected2<hip_bfloat16, ReduceOp>; break;
#endif #endif
#if HAVE_ncclfp8_DEVICE #if HAVE_ncclfp8_DEVICE || HIP_VERSION < 60300000 //for backward compatibility
case ncclFloat8e4m3: fn = (void const*)&prepareExpected2<rccl_float8, ReduceOp>; break; case ncclFloat8e4m3: fn = (void const*)&prepareExpected2<rccl_float8, ReduceOp>; break;
case ncclFloat8e5m2: fn = (void const*)&prepareExpected2<rccl_bfloat8, ReduceOp>; break; case ncclFloat8e5m2: fn = (void const*)&prepareExpected2<rccl_bfloat8, ReduceOp>; break;
#endif #elif HAVE_ncclfp8_HOST
#if HAVE_ncclfp8_HOST
case ncclFloat8e4m3: if (rccl_float8_useFnuz) { fn = (void const*)&prepareExpected2<__hip_fp8_e4m3_fnuz, ReduceOp>; break; } case ncclFloat8e4m3: if (rccl_float8_useFnuz) { fn = (void const*)&prepareExpected2<__hip_fp8_e4m3_fnuz, ReduceOp>; break; }
else { fn = (void const*)&prepareExpected2<__hip_fp8_e4m3, ReduceOp>; break; } else { fn = (void const*)&prepareExpected2<__hip_fp8_e4m3, ReduceOp>; break; }
case ncclFloat8e5m2: if (rccl_float8_useFnuz) { fn = (void const*)&prepareExpected2<__hip_fp8_e5m2_fnuz, ReduceOp>; break; } case ncclFloat8e5m2: if (rccl_float8_useFnuz) { fn = (void const*)&prepareExpected2<__hip_fp8_e5m2_fnuz, ReduceOp>; break; }
@@ -1323,11 +1321,10 @@ hipError_t ncclVerifiableVerify(
#if HAVE_ncclBfloat16 #if HAVE_ncclBfloat16
case ncclBfloat16: CASE_TY(hip_bfloat16, uint16_t) case ncclBfloat16: CASE_TY(hip_bfloat16, uint16_t)
#endif #endif
#if HAVE_ncclfp8_DEVICE #if HAVE_ncclfp8_DEVICE || HIP_VERSION < 60300000
case ncclFloat8e4m3: CASE_TY(rccl_float8, uint8_t) case ncclFloat8e4m3: CASE_TY(rccl_float8, uint8_t)
case ncclFloat8e5m2: CASE_TY(rccl_bfloat8, uint8_t) case ncclFloat8e5m2: CASE_TY(rccl_bfloat8, uint8_t)
#endif #elif HAVE_ncclfp8_HOST
#if HAVE_ncclfp8_HOST
case ncclFloat8e4m3: if (rccl_float8_useFnuz) { CASE_TY(__hip_fp8_e4m3_fnuz, uint8_t);} case ncclFloat8e4m3: if (rccl_float8_useFnuz) { CASE_TY(__hip_fp8_e4m3_fnuz, uint8_t);}
else { CASE_TY(__hip_fp8_e4m3, uint8_t);} else { CASE_TY(__hip_fp8_e4m3, uint8_t);}
case ncclFloat8e5m2: if (rccl_float8_useFnuz) { CASE_TY(__hip_fp8_e5m2_fnuz, uint8_t);} case ncclFloat8e5m2: if (rccl_float8_useFnuz) { CASE_TY(__hip_fp8_e5m2_fnuz, uint8_t);}