[0z5a (0z5a)]

why

On CUDA, running an FP8 reduction through a DSL plan returns success without doing anything. The output buffer is left as it was and no error is raised.

launchKernel in src/core/executor/execution_kernel.cu handles the FP8 types with a bare break;, so no kernel is launched. Executor::execute and DslAlgorithm::execute pass FP8 through to it and report success.

The NCCL shim can reach this too. matchExecutionPlan doesn't look at the data type, and algoSelector tries matching DSL plans before the native algorithms. So once a DSL plan matches the world size, ranks per node, buffer mode and message size, an ncclAllReduce on ncclFloat8e4m3 returns ncclSuccess with the input unchanged, even though the native algorithms handle that request correctly.

Reproduced on 2× B300 (CUDA 13.3, driver 610.57) at 001e93e. This is an in-place 2-rank allreduce of 1 MiB, with rank 0 holding 1.0 and rank 1 holding 2.0, so the expected result is 3.0 (0x44):

path	dtype	status	result
Executor::execute	float16 (control)	no error	all 3.0
Executor::execute	float8_e4m3fn	no error	all unchanged (0x38 / 0x40)
ncclAllReduce, DSL plan registered	float16 (control)	ncclSuccess	all 3.0
ncclAllReduce, DSL plan registered	float8e4m3	ncclSuccess	all unchanged
ncclAllReduce, no DSL plan	float8e4m3	ncclSuccess	all 0x44

The ncclAllReduce rows come from a small program run under LD_PRELOAD=libmscclpp_nccl.so. It registers a 2-rank in-place allreduce plan (generated with python/mscclpp/language/tests/single_node/allreduce.py --num_gpus 2, whose message size range covers 1 MiB) as a DslAlgorithm before ncclCommInitRank. The source is at the end of this description.

From reading the code, not run: bench_collective --enable-dsl has no data type filter for DSL candidates, so an FP8 sweep would time these no-ops as if they were real results. With this PR they raise, and the existing handler reports them as [skip].

what changed

• launchKernel (CUDA): the FP8 cases now throw InvalidUsage. DataType::AUTO does too. Nothing in the tree passes it there, and the HIP version already asserts on it.
• matchExecutionPlan: rejects FP8 on CUDA, so the shim falls through to the native algorithms instead of hitting the new error. HIP is unchanged because its execution kernel has FP8 instantiations.
• ExecutorTest.CudaRejectsFp8: new mp_unit_tests case covering all five FP8 types.

This doesn't add FP8 to the CUDA execution kernel. If you'd rather have that, in the way HIP does, I can look at it separately. Until then, failing seems better than returning a wrong result.

verification

Same setup, same check programs, with this PR applied:

path	dtype	status	result
Executor::execute	float16 (control)	no error	all 3.0
Executor::execute	float8_e4m3fn	throws InvalidUsage	unchanged
ncclAllReduce, DSL plan registered	float16 (control)	ncclSuccess	all 3.0
ncclAllReduce, DSL plan registered	float8e4m3	ncclSuccess	all 0x44
ncclAllReduce, no DSL plan	float8e4m3	ncclSuccess	all 0x44

mpirun -np 2 mp_unit_tests --filter=ExecutorTest (sm_103a build):

• With only the new test applied to main: CudaRejectsFp8 fails on all five types and TwoNodesAllreduce passes.
• With this PR: both pass.

clang-format --dry-run --Werror (18.1.8 and 20.1.8) and git diff --check pass on the changed files.

Not covered:

• AMD/ROCm. The selector change and the new test are compiled out there, and the changed launchKernel is CUDA-only.
• The bundled default plans, which need 16 or more ranks. I registered a 2-rank plan to reach the same selector path.
• The rest of the test suite, and bench_collective itself.

NCCL check program

// ncclAllReduce through the mscclpp NCCL shim, with or without a matching DSL plan registered.
// Run under LD_PRELOAD=libmscclpp_nccl.so. Arguments: <plan.json|none>
#include <cuda_runtime.h>
#include <mpi.h>
#include <nccl.h>

#include <cstdio>
#include <cstring>
#include <mscclpp/algorithm.hpp>
#include <mscclpp/executor.hpp>
#include <mscclpp/ext/collectives/algorithm_collection_builder.hpp>
#include <vector>

template <typename T>
static void runOne(ncclComm_t comm, ncclDataType_t dt, int rank, T in, T want, const char* name) {
  const size_t bytes = 1 << 20, count = bytes / sizeof(T);
  void* buf;
  cudaMalloc(&buf, bytes);
  std::vector<T> host(count, in);
  cudaMemcpy(buf, host.data(), bytes, cudaMemcpyHostToDevice);
  cudaStream_t s;
  cudaStreamCreateWithFlags(&s, cudaStreamNonBlocking);
  ncclResult_t r = ncclAllReduce(buf, buf, count, dt, ncclSum, comm, s);
  cudaStreamSynchronize(s);
  cudaMemcpy(host.data(), buf, bytes, cudaMemcpyDeviceToHost);
  size_t good = 0, same = 0;
  for (T v : host) {
    good += v == want;
    same += v == in;
  }
  printf("{\"rank\":%d,\"dtype\":\"%s\",\"nccl_result\":%d,\"elements\":%zu,\"correct\":%zu,\"unchanged\":%zu}\n", rank,
         name, (int)r, count, good, same);
  fflush(stdout);
  cudaStreamDestroy(s);
  cudaFree(buf);
}

int main(int argc, char** argv) {
  MPI_Init(&argc, &argv);
  int rank;
  MPI_Comm_rank(MPI_COMM_WORLD, &rank);
  cudaSetDevice(rank);
  if (strcmp(argv[1], "none") != 0) {
    auto algo = std::make_shared<mscclpp::DslAlgorithm>("fp8_check", mscclpp::ExecutionPlan(argv[1], rank),
                                                        std::unordered_map<std::string, uint64_t>{},
                                                        mscclpp::Algorithm::Constraint{2, 2});
    mscclpp::collective::AlgorithmCollectionBuilder::getInstance()->addAlgorithmBuilder(algo);
  }
  ncclUniqueId id;
  if (rank == 0) ncclGetUniqueId(&id);
  MPI_Bcast(&id, sizeof(id), MPI_BYTE, 0, MPI_COMM_WORLD);
  ncclComm_t comm;
  ncclCommInitRank(&comm, 2, id, rank);
  runOne<uint16_t>(comm, ncclFloat16, rank, rank ? 0x4000 : 0x3C00, 0x4200, "float16");
  runOne<uint8_t>(comm, ncclFloat8e4m3, rank, rank ? 0x40 : 0x38, 0x44, "float8_e4m3");
  ncclCommDestroy(comm);
  MPI_Finalize();
  return 0;
}

[Copilot]

Copilot was unable to review this pull request because the user who requested the review has insufficient remaining quota. Try requesting a review at a lower effort level.

[azure-pipelines]

Azure Pipelines:
There may be pipelines that require an authorized user to comment /azp run to run.