obscura/nccl-tests
1
0
Fork 0
mirror of https://github.com/NVIDIA/nccl-tests.git synced 2026-09-18 08:10:13 +00:00
Code Issues Packages Projects Releases Wiki Activity

Pass COMM to getDevCommRequirements function

Browse Source 
Pass the target ncclComm_t to the getDevCommmRequirements function and
use that to retrieve properties.

This change is required to add a condition in getDevCommRequirements
that validates whether a test can be run on a cluster given its configuration.

Signed-off-by: Ahsan Pervaiz <[email protected]>
This commit is contained in:
Ahsan Pervaiz
2026-01-12 16:32:14 -08:00
parent ca7625d41a
commit 1d317d53ba
4 changed files with 19 additions and 21 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/all_reduce.cu
+8 -3
View File
@@ -66,8 +66,13 @@ void AllReduceGetBw(size_t count, int typesize, double sec, double* algBw, doubl
#if NCCL_VERSION_CODE >= NCCL_VERSION(2,29,0)
// set devComm reqs for allreduce device kernels
testResult_t AllReduceGetDevCommRequirements(int deviceImpl, ncclDevCommRequirements* reqs, ncclCommProperties_t* commProperties) {
if (!reqs || !commProperties) return testInternalError;
testResult_t AllReduceGetDevCommRequirements(int deviceImpl, ncclDevCommRequirements* reqs, ncclComm_t comm) {
if (!reqs || !comm) return testInternalError;
ncclCommProperties_t commProperties = NCCL_COMM_PROPERTIES_INITIALIZER;
if (ncclCommQueryProperties(comm, &commProperties) != ncclSuccess) {
return testNcclError;
}
switch(deviceImpl) {
case 1: // allReduceLsaKernel
@@ -76,7 +81,7 @@ testResult_t AllReduceGetDevCommRequirements(int deviceImpl, ncclDevCommRequirem
return testSuccess;
case 3: // allReduceMultimemKernel
case 4: // allReduceMultimemVectorizedKernel
if (!commProperties->multimemSupport) {
if (!commProperties.multimemSupport) {
fprintf(stderr, "This test requires multimem support, but multimem support is not enabled for this communicator.\n");
return testInvalidUsage;
}
src/alltoall.cu
+8 -3
View File
@@ -53,8 +53,13 @@ void AlltoAllGetBw(size_t count, int typesize, double sec, double* algBw, double
#if NCCL_VERSION_CODE >= NCCL_VERSION(2,29,0)
// set devComm reqs for alltoall device kernels
testResult_t AlltoAllGetDevCommRequirements(int deviceImpl, ncclDevCommRequirements* reqs, ncclCommProperties_t* commProperties) {
if (!reqs || !commProperties) return testInternalError;
testResult_t AlltoAllGetDevCommRequirements(int deviceImpl, ncclDevCommRequirements* reqs, ncclComm_t comm) {
if (!reqs || !comm) return testInternalError;
ncclCommProperties_t commProperties = NCCL_COMM_PROPERTIES_INITIALIZER;
if (ncclCommQueryProperties(comm, &commProperties) != ncclSuccess) {
return testNcclError;
}
switch(deviceImpl) {
case 1: // NvlAlltoAllKernel
@@ -63,7 +68,7 @@ testResult_t AlltoAllGetDevCommRequirements(int deviceImpl, ncclDevCommRequireme
return testSuccess;
case 3: // GinAlltoAllKernel
case 4: // HybridAlltoAllKernel (LSA+GIN)
if (commProperties->ginType == NCCL_GIN_TYPE_NONE) {
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 testInvalidUsage;
}
src/common.cu
+2 -14
View File
@@ -834,13 +834,7 @@ testResult_t threadInit(struct threadArgs* args) {
fprintf(stderr, "Device implementation %d is not supported by this test\n", deviceImpl);
return testNotImplemented;
}
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));
TESTCHECK(ncclTestEngine.getDevCommRequirements(deviceImpl, &reqs, args->comms[0]));
#else
if (test_ncclVersion >= NCCL_VERSION(2,29,0)) {
fprintf(stderr, "Incompatible NCCL versions. nccl-tests was compiled with NCCL 2.28, but is running with NCCL %d. "
@@ -1414,13 +1408,7 @@ testResult_t run() {
fprintf(stderr, "Device implementation %d is not supported by this test\n", deviceImpl);
return testNotImplemented;
}
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));
TESTCHECK(ncclTestEngine.getDevCommRequirements(deviceImpl, &reqs, comms[0]));
#else
if (test_ncclVersion >= NCCL_VERSION(2,29,0)) {
fprintf(stderr, "Incompatible NCCL versions. nccl-tests was compiled with NCCL 2.28, but is running with NCCL %d. "
src/common.h
+1 -1
View File
@@ -114,7 +114,7 @@ struct testEngine {
const char* typeName, ncclRedOp_t op, const char* opName);
#if NCCL_VERSION_CODE >= NCCL_VERSION(2,29,0)
testResult_t (*getDevCommRequirements)(int deviceImpl, ncclDevCommRequirements* reqs, ncclCommProperties_t* commProperties);
testResult_t (*getDevCommRequirements)(int deviceImpl, ncclDevCommRequirements* reqs, ncclComm_t comm);
#elif NCCL_VERSION_CODE >= NCCL_VERSION(2,28,0)
bool (*getDevCommRequirements)(int deviceImpl, ncclDevCommRequirements* reqs);
#endif