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)) {