NCCL 2.27.7-1

Prevent initialization failures in certain configurations when attempting
to load fp8-specific symmetric multicast kernels on GPUs older than
Blackwell.
This commit is contained in:
Kamil Iskra
2025-07-24 10:39:53 -07:00
parent 0d1ece2b43
commit 593de54e52
5 changed files with 382 additions and 7 deletions
+1 -4
View File
@@ -38,12 +38,9 @@ ncclResult_t ncclInitKernelsForDevice(int cudaArch, int maxSharedMem, size_t* ma
if (fn == nullptr) continue;
cudaError_t errcode = cudaFuncGetAttributes(&attr, fn);
if (errcode == cudaErrorNoKernelImageForDevice) continue;
CUDACHECKGOTO(errcode, result, ignore0);
if (errcode != cudaSuccess) continue; // Silently ignore failures
if (maxStackSize) {
if (attr.localSizeBytes > *maxStackSize) *maxStackSize = attr.localSizeBytes;
ignore0:;
}
if (carveout) {
CUDACHECKGOTO(cudaFuncSetAttribute(fn,