From 86c1297a95e2a0d1e09a96c568ecda42eb4ae5be Mon Sep 17 00:00:00 2001 From: Katie Date: Mon, 8 Dec 2025 19:12:28 +0000 Subject: [PATCH] device implementations (-D > 0) require device API support from the communicator. Return invalidUsage if not supported --- src/all_reduce.cu | 2 +- src/alltoall.cu | 2 +- src/common.cu | 12 ++++++++++-- 3 files changed, 12 insertions(+), 4 deletions(-) diff --git a/src/all_reduce.cu b/src/all_reduce.cu index 260e184..ee23426 100644 --- a/src/all_reduce.cu +++ b/src/all_reduce.cu @@ -78,7 +78,7 @@ testResult_t AllReduceGetDevCommRequirements(int deviceImpl, ncclDevCommRequirem case 4: // allReduceMultimemVectorizedKernel if (!commProperties->multimemSupport) { fprintf(stderr, "This test requires multimem support, but multimem support is not enabled for this communicator.\n"); - return testInternalError; + return testInvalidUsage; } reqs->lsaMultimem = true; reqs->lsaBarrierCount = deviceCtaCount; diff --git a/src/alltoall.cu b/src/alltoall.cu index f354e26..c07c82e 100644 --- a/src/alltoall.cu +++ b/src/alltoall.cu @@ -65,7 +65,7 @@ testResult_t AlltoAllGetDevCommRequirements(int deviceImpl, ncclDevCommRequireme case 4: // HybridAlltoAllKernel (LSA+GIN) if (commProperties->ginType == NCCL_GIN_TYPE_NONE) { fprintf(stderr, "This test requires GIN support, but GIN support is not enabled for this communicator.\n"); - return testInternalError; + return testInvalidUsage; } reqs->barrierCount = deviceCtaCount; reqs->ginSignalCount = deviceCtaCount; diff --git a/src/common.cu b/src/common.cu index 4da475f..d9a04be 100644 --- a/src/common.cu +++ b/src/common.cu @@ -827,8 +827,12 @@ testResult_t threadInit(struct threadArgs* args) { fprintf(stderr, "Device implementation %d is not supported by this test\n", deviceImpl); return testNotImplemented; } - ncclCommProperties commProperties = NCCL_COMM_PROPERTIES_INITIALIZER; + ncclCommProperties_t commProperties = NCCL_COMM_PROPERTIES_INITIALIZER; NCCLCHECK(ncclCommQueryProperties(args->comms[0], &commProperties)); + if (!commProperties.deviceApiSupport) { + fprintf(stderr, "Device API is not supported by this communicator.\n"); + return testInvalidUsage; + } TESTCHECK(ncclTestEngine.getDevCommRequirements(deviceImpl, &reqs, &commProperties)); #else if (test_ncclVersion >= NCCL_VERSION(2,29,0)) { @@ -1422,8 +1426,12 @@ testResult_t run() { fprintf(stderr, "Device implementation %d is not supported by this test\n", deviceImpl); return testNotImplemented; } - ncclCommProperties commProperties = NCCL_COMM_PROPERTIES_INITIALIZER; + ncclCommProperties_t commProperties = NCCL_COMM_PROPERTIES_INITIALIZER; NCCLCHECK(ncclCommQueryProperties(comms[0], &commProperties)); + if (!commProperties.deviceApiSupport) { + fprintf(stderr, "Device API is not supported by this communicator.\n"); + return testInvalidUsage; + } TESTCHECK(ncclTestEngine.getDevCommRequirements(deviceImpl, &reqs, &commProperties)); #else if (test_ncclVersion >= NCCL_VERSION(2,29,0)) {