feat(rccl): add gfx1150 support

This commit is contained in:
2026-03-15 10:59:18 +01:00
bovenliggende a2686c9f41
commit 0c3405a6cc
5 gewijzigde bestanden met toevoegingen van 5 en 4 verwijderingen
+1
Bestand weergeven
@@ -62,6 +62,7 @@ set(DEFAULT_GPUS
gfx1102
gfx1200
gfx1201
gfx1150
gfx1151)
# Load CMake modules
+1 -1
Bestand weergeven
@@ -26,7 +26,7 @@
{ __atomic_store_n((DST), (SRC), __ATOMIC_SEQ_CST); }
#endif
#if defined(__gfx1100__) || defined(__gfx1101__) || defined(__gfx1102__) || defined(__gfx1151__) || defined(__gfx1200__) || defined(__gfx1201__)
#if defined(__gfx1100__) || defined(__gfx1101__) || defined(__gfx1102__) || defined(__gfx1150__) || defined(__gfx1151__) || defined(__gfx1200__) || defined(__gfx1201__)
#define __trace_hwreg() \
collTrace->data_0 = 0;
#else
+1 -1
Bestand weergeven
@@ -1019,7 +1019,7 @@ ncclResult_t ncclTopoGetAlgoTime(struct ncclComm* comm, int coll, int algorithm,
int rcclGetTuningIndexForArch(const char* gfxarch) {
static const std::vector<std::pair<std::string, int>> tuningIndexMap = {
{"gfx906", 0}, {"gfx908", 0}, {"gfx90a", 0}, {"gfx942", 5},
{"gfx950", 6}, {"gfx1030", 0}, {"gfx1100", 0}, {"gfx1102", 0}, {"gfx1151", 0},
{"gfx950", 6}, {"gfx1030", 0}, {"gfx1100", 0}, {"gfx1102", 0}, {"gfx1150", 0}, {"gfx1151", 0},
{"gfx1200", 7}, {"gfx1201", 7}
};
if (gfxarch == nullptr) return 0;
@@ -41,7 +41,7 @@ typedef struct
} rccl_bfloat8;
// __cplusplus < 201103L || (!defined(__HIP_PLATFORM_AMD__) && !defined(__HIPCC__))
#elif HIP_VERSION >= 60300000 && !(defined(__gfx1100__) || defined(__gfx1101__) || defined(__gfx1102__) || defined(__gfx1151__) || defined(__gfx1030__))
#elif HIP_VERSION >= 60300000 && !(defined(__gfx1100__) || defined(__gfx1101__) || defined(__gfx1102__) || defined(__gfx1150__) || defined(__gfx1151__) || defined(__gfx1030__))
#include <hip/hip_fp8.h>
@@ -43,7 +43,7 @@ THE SOFTWARE.
#endif
// Macro for collecting HW_REG_HW_ID
#if defined(__gfx1100__) || defined(__gfx1101__) || defined(__gfx1102__) || defined(__gfx1151__) || defined(__NVCC__)
#if defined(__gfx1100__) || defined(__gfx1101__) || defined(__gfx1102__) || defined(__gfx1150__) || || defined(__gfx1151__) || defined(__NVCC__)
#define GetHwId(val) \
val = 0
#else