Skip to content

Tune the nThreadsPerBlock for FP8 and Half datatype on MI300 #694

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
wants to merge 4 commits into
base: main
Choose a base branch
from
+25 −0
Open

Tune the nThreadsPerBlock for FP8 and Half datatype on MI300 #694

Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
c409d0a
Tune the nThreadsPerBlock for 32KB to 256KB with FP8 and Half datatype
seagater Nov 20, 2025
ecefa62
Improve performance of 32KB and 64KB further for FP8
seagater Nov 21, 2025
afc8801
Merge branch 'main' of https://github.com/microsoft/mscclpp into qing…
seagater Nov 21, 2025
cceb271
Use else if for condition checks
seagater Nov 21, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
25 changes: 25 additions & 0 deletions apps/nccl/src/allreduce.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -48,6 +48,31 @@ struct AllpairAdapter {
nBlocks = (worldSize - 1) * 8;
nThreadsPerBlock = (nelems <= 76800) ? 512 : 1024;
deviceFlag = deviceFlag56;

#if defined(__HIP_PLATFORM_AMD__)
size_t sizeBytes = sizeof(T) * nelems;
if constexpr (std::is_same_v<T, __half>) {
// Half-specific tuning for 32KB-256KB range
if (sizeBytes == (32 << 10)) {
nThreadsPerBlock = 64;
} else if (sizeBytes >= (64 << 10) && sizeBytes <= (256 << 10)) {
nThreadsPerBlock = 128;
}
}

#if defined(__FP8_TYPES_EXIST__)
// FP8-specific tuning for 32KB-256KB range
if constexpr (std::is_same_v<T, __fp8_e4m3> || std::is_same_v<T, __fp8_e5m2>) {
if (sizeBytes == (32 << 10)) {
nThreadsPerBlock = 64;
} else if (sizeBytes == (64 << 10)) {
nThreadsPerBlock = 128;
} else if (sizeBytes >= (128 << 10) && sizeBytes <= (256 << 10)) {
nThreadsPerBlock = 256;
}
}
#endif
#endif
}
#if defined(ENABLE_NPKIT)
size_t NpkitSharedMemSize = NPKIT_SHM_NUM_EVENTS * sizeof(NpKitEvent);
Expand Down