mirror of
https://github.com/NVIDIA/nccl-tests.git
synced 2026-10-21 08:07:59 +00:00
GIN perf: add host memory type
Also remame the 'hBuf' with 'devBufHandle', just to avoid ambiguity, since `hostBufWindow` will be added. Signed-off-by: Bojun Zhu <[email protected]>
This commit is contained in:
Bojun Zhu
+1
-1
@@ -155,7 +155,7 @@ DEVICE_API_GIN_UTILS_DST := $(BUILDDIR)/device_api/gin/utils
|
||||
DEVICE_API_GIN_MPI_ENABLED := $(filter 1,$(MPI))
|
||||
DEVICE_API_GIN_EXTRA_INC := -I$(DEVICE_API_GIN_UTILS_SRCDIR) -I$(NCCL_HOME)/include -I$(MPI_HOME)/include \
|
||||
-I$(CUDA_INC) -I$(CUDA_INC)/cccl
|
||||
DEVICE_API_GIN_NVLDFLAGS := $(filter-out -lnccl_test_os -L$(TEST_OS_BUILDDIR),$(NVLDFLAGS))
|
||||
DEVICE_API_GIN_NVLDFLAGS := $(filter-out -lnccl_test_os -L$(TEST_OS_BUILDDIR),$(NVLDFLAGS)) -lcuda
|
||||
|
||||
# GIN throughput tests support 1024-thread CTAs, which require at most 60
|
||||
# registers per thread. Change to use these guidelines -> threads:regs
|
||||
|
||||
@@ -58,6 +58,7 @@ function(add_gin_perf_category CATEGORY_DIR)
|
||||
nccl
|
||||
MPI::MPI_CXX
|
||||
CUDA::cudart
|
||||
CUDA::cuda_driver
|
||||
)
|
||||
|
||||
foreach(RUNNER_FILE ${RUNNER_SOURCES})
|
||||
@@ -87,6 +88,7 @@ function(add_gin_perf_category CATEGORY_DIR)
|
||||
nccl
|
||||
MPI::MPI_CXX
|
||||
CUDA::cudart
|
||||
CUDA::cuda_driver
|
||||
m
|
||||
stdc++
|
||||
)
|
||||
|
||||
@@ -121,6 +121,10 @@ support are rejected rather than ignored.
|
||||
Requires `-t` to be a multiple of 32.
|
||||
* `--gin_bd` bidirectional bandwidth: both ranks send; reported metric is the
|
||||
sum of each rank's average (throughput only). Default: off.
|
||||
* `--remote_mem_type <device|host>` selects device memory (default) or CPU-backed host
|
||||
memory for the remote rank (source memory in GET or destination memory in PUT).
|
||||
* `--local_mem_type <device|host>` selects device memory (default) or CPU-backed host
|
||||
memory for the local rank (destination memory in GET or source memory in PUT).
|
||||
|
||||
### Optional optimizations
|
||||
|
||||
|
||||
@@ -6,11 +6,11 @@
|
||||
#include "gin_context.h"
|
||||
|
||||
template <bool skipCreditCheck, ncclGinResourceSharingMode rsm>
|
||||
__global__ void ginGetPingKernel(ncclDevComm comm, ncclDevResourceHandle hBuf, int iters, size_t bytes) {
|
||||
__global__ void ginGetPingKernel(ncclDevComm comm, ncclDevResourceHandle devBufHandle, int iters, size_t bytes) {
|
||||
#if __CUDA_ARCH__ >= 700
|
||||
ncclTeam team = ncclTeamWorld(comm);
|
||||
ncclGin gin(comm, 0, rsm);
|
||||
ncclSymPtr<int> sbuf = (ncclSymPtr<int>)ncclGetResourceBuffer(comm, hBuf);
|
||||
ncclSymPtr<int> sbuf = (ncclSymPtr<int>)ncclGetResourceBuffer(comm, devBufHandle);
|
||||
|
||||
constexpr uint32_t optFlags =
|
||||
skipCreditCheck ? ncclGinOptFlagsMaySkipCreditCheck : ncclGinOptFlagsDefault;
|
||||
@@ -24,11 +24,11 @@ __global__ void ginGetPingKernel(ncclDevComm comm, ncclDevResourceHandle hBuf, i
|
||||
}
|
||||
|
||||
template <ncclGinResourceSharingMode rsm>
|
||||
static void ginGetLatencyPingLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
static void ginGetLatencyPingLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, size_t numElems, int iters) {
|
||||
size_t bytes = numElems * sizeof(int);
|
||||
#define LAUNCH_PING_GET(SKIP) \
|
||||
ginGetPingKernel<SKIP, rsm><<<1, 1, 0, stream>>>(dcomm, hBuf, iters, bytes)
|
||||
ginGetPingKernel<SKIP, rsm><<<1, 1, 0, stream>>>(dcomm, devBufHandle, iters, bytes)
|
||||
|
||||
if (args->ginSkipCreditCheck) LAUNCH_PING_GET(true);
|
||||
else LAUNCH_PING_GET(false);
|
||||
@@ -37,13 +37,13 @@ static void ginGetLatencyPingLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle
|
||||
CUDACHECK_FATAL(cudaGetLastError());
|
||||
}
|
||||
|
||||
void ncclTestGinGetLatencyPingLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
void ncclTestGinGetLatencyPingLaunch(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, size_t numElems, int iters) {
|
||||
if (args->ginRsm == ncclGinRsmThread) {
|
||||
ginGetLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_THREAD>(dcomm, devBufHandle, hostBufWindow, stream, args, numElems, iters);
|
||||
ginGetLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_THREAD>(dcomm, devBufHandle, stream, args, numElems, iters);
|
||||
} else if (args->ginRsm == ncclGinRsmCta) {
|
||||
ginGetLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, hBuf, stream, args, numElems, iters);
|
||||
ginGetLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, devBufHandle, stream, args, numElems, iters);
|
||||
} else {
|
||||
ginGetLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, hBuf, stream, args, numElems, iters);
|
||||
ginGetLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, devBufHandle, stream, args, numElems, iters);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -16,6 +16,7 @@ int main(int argc, char** argv) {
|
||||
caps.allowAggregateRequests = false;
|
||||
caps.allowBidirectional = false;
|
||||
caps.allowMultiCtaThreads = false;
|
||||
caps.canUseHostMemory = false;
|
||||
ncclTestGinConfigureArgs(&args, &caps);
|
||||
|
||||
ginBenchmark_t bench = {};
|
||||
|
||||
@@ -12,23 +12,23 @@
|
||||
* ginBenchmark_t::run directly.
|
||||
*/
|
||||
|
||||
void ncclTestGinPutLatencyPingLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
void ncclTestGinPutLatencyPingLaunch(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, size_t numElems, int iters);
|
||||
void ncclTestGinPutSignalLatencyPingLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
void ncclTestGinPutSignalLatencyPingLaunch(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, size_t numElems, int iters);
|
||||
void ncclTestGinPutCounterLatencyPingLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
void ncclTestGinPutCounterLatencyPingLaunch(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, size_t numElems, int iters);
|
||||
|
||||
void ncclTestGinGetLatencyPingLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
void ncclTestGinGetLatencyPingLaunch(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, size_t numElems, int iters);
|
||||
|
||||
void ncclTestGinSignalLatencyPingLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
void ncclTestGinSignalLatencyPingLaunch(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, size_t numElems, int iters);
|
||||
|
||||
void ncclTestGinPutLatencyPingPongLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
void ncclTestGinPutLatencyPingPongLaunch(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, size_t numElems, int iters);
|
||||
void ncclTestGinPutSignalLatencyPingPongLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
void ncclTestGinPutSignalLatencyPingPongLaunch(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, size_t numElems, int iters);
|
||||
|
||||
void ncclTestGinSignalLatencyPingPongLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
void ncclTestGinSignalLatencyPingPongLaunch(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, size_t numElems, int iters);
|
||||
|
||||
@@ -6,12 +6,12 @@
|
||||
#include "gin_context.h"
|
||||
|
||||
template <bool skipCreditCheck, ncclGinResourceSharingMode rsm>
|
||||
__global__ void ginPutCounterPingKernel(ncclDevComm comm, ncclDevResourceHandle hBuf, int iters, size_t numElems,
|
||||
__global__ void ginPutCounterPingKernel(ncclDevComm comm, ncclDevResourceHandle devBufHandle, int iters, size_t numElems,
|
||||
int queueDepth) {
|
||||
#if __CUDA_ARCH__ >= 700
|
||||
ncclTeam team = ncclTeamWorld(comm);
|
||||
ncclGin gin(comm, 0, rsm);
|
||||
ncclSymPtr<int> sbuf = (ncclSymPtr<int>)ncclGetResourceBuffer(comm, hBuf);
|
||||
ncclSymPtr<int> sbuf = (ncclSymPtr<int>)ncclGetResourceBuffer(comm, devBufHandle);
|
||||
ncclSymPtr<int> dbuf = sbuf;
|
||||
|
||||
constexpr uint32_t optFlags =
|
||||
@@ -35,11 +35,11 @@ __global__ void ginPutCounterPingKernel(ncclDevComm comm, ncclDevResourceHandle
|
||||
}
|
||||
|
||||
template <ncclGinResourceSharingMode rsm>
|
||||
static void ginPutCounterLatencyPingLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
static void ginPutCounterLatencyPingLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, size_t numElems, int iters) {
|
||||
const int queueDepth = args->queueDepth;
|
||||
#define LAUNCH_PING_PUT_COUNTER(SKIP) \
|
||||
ginPutCounterPingKernel<SKIP, rsm><<<1, 1, 0, stream>>>(dcomm, hBuf, iters, numElems, queueDepth)
|
||||
ginPutCounterPingKernel<SKIP, rsm><<<1, 1, 0, stream>>>(dcomm, devBufHandle, iters, numElems, queueDepth)
|
||||
|
||||
if (args->ginSkipCreditCheck) LAUNCH_PING_PUT_COUNTER(true);
|
||||
else LAUNCH_PING_PUT_COUNTER(false);
|
||||
@@ -48,13 +48,13 @@ static void ginPutCounterLatencyPingLaunchRsm(ncclDevComm dcomm, ncclDevResource
|
||||
CUDACHECK_FATAL(cudaGetLastError());
|
||||
}
|
||||
|
||||
void ncclTestGinPutCounterLatencyPingLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
void ncclTestGinPutCounterLatencyPingLaunch(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, size_t numElems, int iters) {
|
||||
if (args->ginRsm == ncclGinRsmThread) {
|
||||
ginPutCounterLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_THREAD>(dcomm, hBuf, stream, args, numElems, iters);
|
||||
ginPutCounterLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_THREAD>(dcomm, devBufHandle, stream, args, numElems, iters);
|
||||
} else if (args->ginRsm == ncclGinRsmCta) {
|
||||
ginPutCounterLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, hBuf, stream, args, numElems, iters);
|
||||
ginPutCounterLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, devBufHandle, stream, args, numElems, iters);
|
||||
} else {
|
||||
ginPutCounterLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, hBuf, stream, args, numElems, iters);
|
||||
ginPutCounterLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, devBufHandle, stream, args, numElems, iters);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -6,11 +6,11 @@
|
||||
#include "gin_context.h"
|
||||
|
||||
template <bool skipCreditCheck, ncclGinResourceSharingMode rsm>
|
||||
__global__ void ginPutPingKernel(ncclDevComm comm, ncclDevResourceHandle hBuf, int iters, size_t numElems) {
|
||||
__global__ void ginPutPingKernel(ncclDevComm comm, ncclDevResourceHandle devBufHandle, int iters, size_t numElems) {
|
||||
#if __CUDA_ARCH__ >= 700
|
||||
ncclTeam team = ncclTeamWorld(comm);
|
||||
ncclGin gin(comm, 0, rsm);
|
||||
ncclSymPtr<int> sbuf = (ncclSymPtr<int>)ncclGetResourceBuffer(comm, hBuf);
|
||||
ncclSymPtr<int> sbuf = (ncclSymPtr<int>)ncclGetResourceBuffer(comm, devBufHandle);
|
||||
ncclSymPtr<int> dbuf = sbuf;
|
||||
|
||||
constexpr uint32_t optFlags =
|
||||
@@ -25,10 +25,10 @@ __global__ void ginPutPingKernel(ncclDevComm comm, ncclDevResourceHandle hBuf, i
|
||||
}
|
||||
|
||||
template <ncclGinResourceSharingMode rsm>
|
||||
static void ginPutLatencyPingLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
static void ginPutLatencyPingLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, size_t numElems, int iters) {
|
||||
#define LAUNCH_PING_PUT(SKIP) \
|
||||
ginPutPingKernel<SKIP, rsm><<<1, 1, 0, stream>>>(dcomm, hBuf, iters, numElems)
|
||||
ginPutPingKernel<SKIP, rsm><<<1, 1, 0, stream>>>(dcomm, devBufHandle, iters, numElems)
|
||||
|
||||
if (args->ginSkipCreditCheck) LAUNCH_PING_PUT(true);
|
||||
else LAUNCH_PING_PUT(false);
|
||||
@@ -37,13 +37,13 @@ static void ginPutLatencyPingLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle
|
||||
CUDACHECK_FATAL(cudaGetLastError());
|
||||
}
|
||||
|
||||
void ncclTestGinPutLatencyPingLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
void ncclTestGinPutLatencyPingLaunch(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, size_t numElems, int iters) {
|
||||
if (args->ginRsm == ncclGinRsmThread) {
|
||||
ginPutLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_THREAD>(dcomm, hBuf, stream, args, numElems, iters);
|
||||
ginPutLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_THREAD>(dcomm, devBufHandle, stream, args, numElems, iters);
|
||||
} else if (args->ginRsm == ncclGinRsmCta) {
|
||||
ginPutLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, hBuf, stream, args, numElems, iters);
|
||||
ginPutLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, devBufHandle, stream, args, numElems, iters);
|
||||
} else {
|
||||
ginPutLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, hBuf, stream, args, numElems, iters);
|
||||
ginPutLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, devBufHandle, stream, args, numElems, iters);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -6,12 +6,12 @@
|
||||
#include "gin_context.h"
|
||||
|
||||
template <bool skipCreditCheck, ncclGinResourceSharingMode rsm>
|
||||
__global__ void ginPutPingPongKernel(ncclDevComm comm, ncclDevResourceHandle hBuf, int iters, size_t numElems,
|
||||
__global__ void ginPutPingPongKernel(ncclDevComm comm, ncclDevResourceHandle devBufHandle, int iters, size_t numElems,
|
||||
int queueDepth) {
|
||||
#if __CUDA_ARCH__ >= 700
|
||||
ncclTeam team = ncclTeamWorld(comm);
|
||||
ncclGin gin(comm, 0, rsm);
|
||||
ncclSymPtr<int> sbuf = (ncclSymPtr<int>)ncclGetResourceBuffer(comm, hBuf);
|
||||
ncclSymPtr<int> sbuf = (ncclSymPtr<int>)ncclGetResourceBuffer(comm, devBufHandle);
|
||||
ncclSymPtr<int> dbuf = sbuf;
|
||||
|
||||
constexpr uint32_t optFlags =
|
||||
@@ -19,7 +19,7 @@ __global__ void ginPutPingPongKernel(ncclDevComm comm, ncclDevResourceHandle hBu
|
||||
|
||||
const int flushEvery = queueDepth / 2 < 1 ? 1 : queueDepth / 2;
|
||||
|
||||
int* lastElem = (int*)ncclGetResourceBufferLocalPointer(comm, hBuf) + (numElems - 1);
|
||||
int* lastElem = (int*)ncclGetResourceBufferLocalPointer(comm, devBufHandle) + (numElems - 1);
|
||||
|
||||
int serverSign = 1;
|
||||
int clientSign = -1;
|
||||
@@ -47,11 +47,11 @@ __global__ void ginPutPingPongKernel(ncclDevComm comm, ncclDevResourceHandle hBu
|
||||
}
|
||||
|
||||
template <ncclGinResourceSharingMode rsm>
|
||||
static void ginPutLatencyPingPongLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
static void ginPutLatencyPingPongLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, size_t numElems, int iters) {
|
||||
const int queueDepth = args->queueDepth;
|
||||
#define LAUNCH_PING_PONG_PUT(SKIP) \
|
||||
ginPutPingPongKernel<SKIP, rsm><<<1, 1, 0, stream>>>(dcomm, hBuf, iters, numElems, queueDepth)
|
||||
ginPutPingPongKernel<SKIP, rsm><<<1, 1, 0, stream>>>(dcomm, devBufHandle, iters, numElems, queueDepth)
|
||||
|
||||
if (args->ginSkipCreditCheck) LAUNCH_PING_PONG_PUT(true);
|
||||
else LAUNCH_PING_PONG_PUT(false);
|
||||
@@ -60,13 +60,13 @@ static void ginPutLatencyPingPongLaunchRsm(ncclDevComm dcomm, ncclDevResourceHan
|
||||
CUDACHECK_FATAL(cudaGetLastError());
|
||||
}
|
||||
|
||||
void ncclTestGinPutLatencyPingPongLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
void ncclTestGinPutLatencyPingPongLaunch(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, size_t numElems, int iters) {
|
||||
if (args->ginRsm == ncclGinRsmThread) {
|
||||
ginPutLatencyPingPongLaunchRsm<NCCL_GIN_RESOURCE_SHARING_THREAD>(dcomm, hBuf, stream, args, numElems, iters);
|
||||
ginPutLatencyPingPongLaunchRsm<NCCL_GIN_RESOURCE_SHARING_THREAD>(dcomm, devBufHandle, stream, args, numElems, iters);
|
||||
} else if (args->ginRsm == ncclGinRsmCta) {
|
||||
ginPutLatencyPingPongLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, hBuf, stream, args, numElems, iters);
|
||||
ginPutLatencyPingPongLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, devBufHandle, stream, args, numElems, iters);
|
||||
} else {
|
||||
ginPutLatencyPingPongLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, hBuf, stream, args, numElems, iters);
|
||||
ginPutLatencyPingPongLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, devBufHandle, stream, args, numElems, iters);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -19,6 +19,7 @@ int main(int argc, char** argv) {
|
||||
caps.allowAggregateRequests = false;
|
||||
caps.allowBidirectional = false;
|
||||
caps.allowMultiCtaThreads = false;
|
||||
caps.canUseHostMemory = false;
|
||||
ncclTestGinConfigureArgs(&args, &caps);
|
||||
|
||||
ginBenchmark_t bench = {};
|
||||
|
||||
@@ -16,6 +16,7 @@ int main(int argc, char** argv) {
|
||||
caps.allowAggregateRequests = false;
|
||||
caps.allowBidirectional = false;
|
||||
caps.allowMultiCtaThreads = false;
|
||||
caps.canUseHostMemory = false;
|
||||
ncclTestGinConfigureArgs(&args, &caps);
|
||||
|
||||
ginBenchmark_t bench = {};
|
||||
|
||||
@@ -6,11 +6,11 @@
|
||||
#include "gin_context.h"
|
||||
|
||||
template <bool skipCreditCheck, bool strongSignal, ncclGinResourceSharingMode rsm>
|
||||
__global__ void ginPutSignalPingKernel(ncclDevComm comm, ncclDevResourceHandle hBuf, int iters, size_t numElems) {
|
||||
__global__ void ginPutSignalPingKernel(ncclDevComm comm, ncclDevResourceHandle devBufHandle, int iters, size_t numElems) {
|
||||
#if __CUDA_ARCH__ >= 700
|
||||
ncclTeam team = ncclTeamWorld(comm);
|
||||
ncclGin gin(comm, 0, rsm);
|
||||
ncclSymPtr<int> sbuf = (ncclSymPtr<int>)ncclGetResourceBuffer(comm, hBuf);
|
||||
ncclSymPtr<int> sbuf = (ncclSymPtr<int>)ncclGetResourceBuffer(comm, devBufHandle);
|
||||
ncclSymPtr<int> dbuf = sbuf;
|
||||
|
||||
constexpr uint32_t optFlags =
|
||||
@@ -32,10 +32,10 @@ __global__ void ginPutSignalPingKernel(ncclDevComm comm, ncclDevResourceHandle h
|
||||
}
|
||||
|
||||
template <ncclGinResourceSharingMode rsm>
|
||||
static void ginPutSignalLatencyPingLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
static void ginPutSignalLatencyPingLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, size_t numElems, int iters) {
|
||||
#define LAUNCH_PING_PUT_SIGNAL(SKIP, STRONG) \
|
||||
ginPutSignalPingKernel<SKIP, STRONG, rsm><<<1, 1, 0, stream>>>(dcomm, hBuf, iters, numElems)
|
||||
ginPutSignalPingKernel<SKIP, STRONG, rsm><<<1, 1, 0, stream>>>(dcomm, devBufHandle, iters, numElems)
|
||||
|
||||
if (args->ginSkipCreditCheck) {
|
||||
if (args->ginStrongSignal) LAUNCH_PING_PUT_SIGNAL(true, true);
|
||||
@@ -49,13 +49,13 @@ static void ginPutSignalLatencyPingLaunchRsm(ncclDevComm dcomm, ncclDevResourceH
|
||||
CUDACHECK_FATAL(cudaGetLastError());
|
||||
}
|
||||
|
||||
void ncclTestGinPutSignalLatencyPingLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
void ncclTestGinPutSignalLatencyPingLaunch(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, size_t numElems, int iters) {
|
||||
if (args->ginRsm == ncclGinRsmThread) {
|
||||
ginPutSignalLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_THREAD>(dcomm, hBuf, stream, args, numElems, iters);
|
||||
ginPutSignalLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_THREAD>(dcomm, devBufHandle, stream, args, numElems, iters);
|
||||
} else if (args->ginRsm == ncclGinRsmCta) {
|
||||
ginPutSignalLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, hBuf, stream, args, numElems, iters);
|
||||
ginPutSignalLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, devBufHandle, stream, args, numElems, iters);
|
||||
} else {
|
||||
ginPutSignalLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, hBuf, stream, args, numElems, iters);
|
||||
ginPutSignalLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, devBufHandle, stream, args, numElems, iters);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -6,12 +6,12 @@
|
||||
#include "gin_context.h"
|
||||
|
||||
template <bool skipCreditCheck, bool strongSignal, ncclGinResourceSharingMode rsm>
|
||||
__global__ void ginPutSignalPingPongKernel(ncclDevComm comm, ncclDevResourceHandle hBuf, int iters, size_t numElems,
|
||||
__global__ void ginPutSignalPingPongKernel(ncclDevComm comm, ncclDevResourceHandle devBufHandle, int iters, size_t numElems,
|
||||
int queueDepth) {
|
||||
#if __CUDA_ARCH__ >= 700
|
||||
ncclTeam team = ncclTeamWorld(comm);
|
||||
ncclGin gin(comm, 0, rsm);
|
||||
ncclSymPtr<int> sbuf = (ncclSymPtr<int>)ncclGetResourceBuffer(comm, hBuf);
|
||||
ncclSymPtr<int> sbuf = (ncclSymPtr<int>)ncclGetResourceBuffer(comm, devBufHandle);
|
||||
ncclSymPtr<int> dbuf = sbuf;
|
||||
|
||||
constexpr uint32_t optFlags =
|
||||
@@ -52,11 +52,11 @@ __global__ void ginPutSignalPingPongKernel(ncclDevComm comm, ncclDevResourceHand
|
||||
}
|
||||
|
||||
template <ncclGinResourceSharingMode rsm>
|
||||
static void ginPutSignalLatencyPingPongLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
static void ginPutSignalLatencyPingPongLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, size_t numElems, int iters) {
|
||||
const int queueDepth = args->queueDepth;
|
||||
#define LAUNCH_PING_PONG_PUT_SIGNAL(SKIP, STRONG) \
|
||||
ginPutSignalPingPongKernel<SKIP, STRONG, rsm><<<1, 1, 0, stream>>>(dcomm, hBuf, iters, numElems, queueDepth)
|
||||
ginPutSignalPingPongKernel<SKIP, STRONG, rsm><<<1, 1, 0, stream>>>(dcomm, devBufHandle, iters, numElems, queueDepth)
|
||||
|
||||
if (args->ginSkipCreditCheck) {
|
||||
if (args->ginStrongSignal) LAUNCH_PING_PONG_PUT_SIGNAL(true, true);
|
||||
@@ -70,13 +70,13 @@ static void ginPutSignalLatencyPingPongLaunchRsm(ncclDevComm dcomm, ncclDevResou
|
||||
CUDACHECK_FATAL(cudaGetLastError());
|
||||
}
|
||||
|
||||
void ncclTestGinPutSignalLatencyPingPongLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
void ncclTestGinPutSignalLatencyPingPongLaunch(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, size_t numElems, int iters) {
|
||||
if (args->ginRsm == ncclGinRsmThread) {
|
||||
ginPutSignalLatencyPingPongLaunchRsm<NCCL_GIN_RESOURCE_SHARING_THREAD>(dcomm, hBuf, stream, args, numElems, iters);
|
||||
ginPutSignalLatencyPingPongLaunchRsm<NCCL_GIN_RESOURCE_SHARING_THREAD>(dcomm, devBufHandle, stream, args, numElems, iters);
|
||||
} else if (args->ginRsm == ncclGinRsmCta) {
|
||||
ginPutSignalLatencyPingPongLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, hBuf, stream, args, numElems, iters);
|
||||
ginPutSignalLatencyPingPongLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, devBufHandle, stream, args, numElems, iters);
|
||||
} else {
|
||||
ginPutSignalLatencyPingPongLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, hBuf, stream, args, numElems, iters);
|
||||
ginPutSignalLatencyPingPongLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, devBufHandle, stream, args, numElems, iters);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -6,7 +6,7 @@
|
||||
#include "gin_context.h"
|
||||
|
||||
template <bool skipCreditCheck, bool strongSignal, ncclGinResourceSharingMode rsm>
|
||||
__global__ void ginSignalPingKernel(ncclDevComm comm, ncclDevResourceHandle hBuf, int iters) {
|
||||
__global__ void ginSignalPingKernel(ncclDevComm comm, ncclDevResourceHandle devBufHandle, int iters) {
|
||||
#if __CUDA_ARCH__ >= 700
|
||||
ncclTeam team = ncclTeamWorld(comm);
|
||||
ncclGin gin(comm, 0, rsm);
|
||||
@@ -30,10 +30,10 @@ __global__ void ginSignalPingKernel(ncclDevComm comm, ncclDevResourceHandle hBuf
|
||||
}
|
||||
|
||||
template <ncclGinResourceSharingMode rsm>
|
||||
static void ginSignalLatencyPingLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
static void ginSignalLatencyPingLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, int iters) {
|
||||
#define LAUNCH_PING_SIGNAL(SKIP, STRONG) \
|
||||
ginSignalPingKernel<SKIP, STRONG, rsm><<<1, 1, 0, stream>>>(dcomm, hBuf, iters)
|
||||
ginSignalPingKernel<SKIP, STRONG, rsm><<<1, 1, 0, stream>>>(dcomm, devBufHandle, iters)
|
||||
|
||||
if (args->ginSkipCreditCheck) {
|
||||
if (args->ginStrongSignal) LAUNCH_PING_SIGNAL(true, true);
|
||||
@@ -47,13 +47,13 @@ static void ginSignalLatencyPingLaunchRsm(ncclDevComm dcomm, ncclDevResourceHand
|
||||
CUDACHECK_FATAL(cudaGetLastError());
|
||||
}
|
||||
|
||||
void ncclTestGinSignalLatencyPingLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
void ncclTestGinSignalLatencyPingLaunch(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, size_t numElems, int iters) {
|
||||
if (args->ginRsm == ncclGinRsmThread) {
|
||||
ginSignalLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_THREAD>(dcomm, hBuf, stream, args, iters);
|
||||
ginSignalLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_THREAD>(dcomm, devBufHandle, stream, args, iters);
|
||||
} else if (args->ginRsm == ncclGinRsmCta) {
|
||||
ginSignalLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, hBuf, stream, args, iters);
|
||||
ginSignalLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, devBufHandle, stream, args, iters);
|
||||
} else {
|
||||
ginSignalLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, hBuf, stream, args, iters);
|
||||
ginSignalLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, devBufHandle, stream, args, iters);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -6,7 +6,7 @@
|
||||
#include "gin_context.h"
|
||||
|
||||
template <bool skipCreditCheck, bool strongSignal, ncclGinResourceSharingMode rsm>
|
||||
__global__ void ginSignalPingPongKernel(ncclDevComm comm, ncclDevResourceHandle hBuf, int iters,
|
||||
__global__ void ginSignalPingPongKernel(ncclDevComm comm, ncclDevResourceHandle devBufHandle, int iters,
|
||||
int queueDepth) {
|
||||
#if __CUDA_ARCH__ >= 700
|
||||
ncclTeam team = ncclTeamWorld(comm);
|
||||
@@ -50,11 +50,11 @@ __global__ void ginSignalPingPongKernel(ncclDevComm comm, ncclDevResourceHandle
|
||||
}
|
||||
|
||||
template <ncclGinResourceSharingMode rsm>
|
||||
static void ginSignalLatencyPingPongLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
static void ginSignalLatencyPingPongLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, int iters) {
|
||||
const int queueDepth = args->queueDepth;
|
||||
#define LAUNCH_PING_PONG_SIGNAL(SKIP, STRONG) \
|
||||
ginSignalPingPongKernel<SKIP, STRONG, rsm><<<1, 1, 0, stream>>>(dcomm, hBuf, iters, queueDepth)
|
||||
ginSignalPingPongKernel<SKIP, STRONG, rsm><<<1, 1, 0, stream>>>(dcomm, devBufHandle, iters, queueDepth)
|
||||
|
||||
if (args->ginSkipCreditCheck) {
|
||||
if (args->ginStrongSignal) LAUNCH_PING_PONG_SIGNAL(true, true);
|
||||
@@ -68,13 +68,13 @@ static void ginSignalLatencyPingPongLaunchRsm(ncclDevComm dcomm, ncclDevResource
|
||||
CUDACHECK_FATAL(cudaGetLastError());
|
||||
}
|
||||
|
||||
void ncclTestGinSignalLatencyPingPongLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
void ncclTestGinSignalLatencyPingPongLaunch(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, size_t numElems, int iters) {
|
||||
if (args->ginRsm == ncclGinRsmThread) {
|
||||
ginSignalLatencyPingPongLaunchRsm<NCCL_GIN_RESOURCE_SHARING_THREAD>(dcomm, hBuf, stream, args, iters);
|
||||
ginSignalLatencyPingPongLaunchRsm<NCCL_GIN_RESOURCE_SHARING_THREAD>(dcomm, devBufHandle, stream, args, iters);
|
||||
} else if (args->ginRsm == ncclGinRsmCta) {
|
||||
ginSignalLatencyPingPongLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, hBuf, stream, args, iters);
|
||||
ginSignalLatencyPingPongLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, devBufHandle, stream, args, iters);
|
||||
} else {
|
||||
ginSignalLatencyPingPongLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, hBuf, stream, args, iters);
|
||||
ginSignalLatencyPingPongLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, devBufHandle, stream, args, iters);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -16,6 +16,7 @@ int main(int argc, char** argv) {
|
||||
caps.allowAggregateRequests = false;
|
||||
caps.allowBidirectional = false;
|
||||
caps.allowMultiCtaThreads = false;
|
||||
caps.canUseHostMemory = false;
|
||||
ncclTestGinConfigureArgs(&args, &caps);
|
||||
|
||||
ginBenchmark_t bench = {};
|
||||
|
||||
@@ -16,6 +16,7 @@ int main(int argc, char** argv) {
|
||||
caps.allowAggregateRequests = false;
|
||||
caps.allowBidirectional = false;
|
||||
caps.allowMultiCtaThreads = false;
|
||||
caps.canUseHostMemory = false;
|
||||
ncclTestGinConfigureArgs(&args, &caps);
|
||||
|
||||
ginBenchmark_t bench = {};
|
||||
|
||||
@@ -6,7 +6,7 @@
|
||||
#include "gin_context.h"
|
||||
|
||||
template <bool skipCreditCheck, bool aggregateRequests, ncclGinResourceSharingMode rsm>
|
||||
__global__ void ginGetBwKernel(ncclDevComm comm, ncclDevResourceHandle hBuf, int iters, size_t numElems,
|
||||
__global__ void ginGetBwKernel(ncclDevComm comm, ncclDevResourceHandle devBufHandle, int iters, size_t numElems,
|
||||
int queueDepth, size_t maxElems) {
|
||||
#if __CUDA_ARCH__ >= 700
|
||||
const int tag = blockIdx.x;
|
||||
@@ -15,7 +15,7 @@ __global__ void ginGetBwKernel(ncclDevComm comm, ncclDevResourceHandle hBuf, int
|
||||
ncclGin gin(comm, tag, rsm);
|
||||
|
||||
const size_t slots = maxElems / numElems;
|
||||
ncclSymPtr<int> buf = (ncclSymPtr<int>)ncclGetResourceBuffer(comm, hBuf);
|
||||
ncclSymPtr<int> buf = (ncclSymPtr<int>)ncclGetResourceBuffer(comm, devBufHandle);
|
||||
buf += (size_t)(threadIdx.x % slots) * numElems;
|
||||
ncclSymPtr<int> lbuf = buf;
|
||||
ncclSymPtr<int> rbuf = buf;
|
||||
@@ -56,12 +56,12 @@ __global__ void ginGetBwKernel(ncclDevComm comm, ncclDevResourceHandle hBuf, int
|
||||
}
|
||||
|
||||
template <ncclGinResourceSharingMode rsm>
|
||||
static void ginGetBWLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
static void ginGetBWLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, size_t numElems, int iters) {
|
||||
const int queueDepth = args->queueDepth;
|
||||
const size_t maxElems = args->maxBytes / sizeof(int);
|
||||
#define LAUNCH_BW_GET(SKIP, AG) \
|
||||
ginGetBwKernel<SKIP, AG, rsm><<<args->numCtas, args->numThreads, 0, stream>>>(dcomm, hBuf, iters, numElems, queueDepth, maxElems)
|
||||
ginGetBwKernel<SKIP, AG, rsm><<<args->numCtas, args->numThreads, 0, stream>>>(dcomm, devBufHandle, iters, numElems, queueDepth, maxElems)
|
||||
|
||||
if (args->ginSkipCreditCheck) {
|
||||
if (args->ginAggregateRequests) LAUNCH_BW_GET(true, true);
|
||||
@@ -75,11 +75,11 @@ static void ginGetBWLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cud
|
||||
CUDACHECK_FATAL(cudaGetLastError());
|
||||
}
|
||||
|
||||
void ncclTestGinGetBWLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
void ncclTestGinGetBWLaunch(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, size_t numElems, int iters) {
|
||||
if (args->ginRsm == ncclGinRsmCta) {
|
||||
ginGetBWLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, hBuf, stream, args, numElems, iters);
|
||||
ginGetBWLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, devBufHandle, stream, args, numElems, iters);
|
||||
} else {
|
||||
ginGetBWLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, hBuf, stream, args, numElems, iters);
|
||||
ginGetBWLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, devBufHandle, stream, args, numElems, iters);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -16,6 +16,7 @@ int main(int argc, char** argv) {
|
||||
caps.allowAggregateRequests = true;
|
||||
caps.allowBidirectional = true;
|
||||
caps.allowMultiCtaThreads = true;
|
||||
caps.canUseHostMemory = false;
|
||||
ncclTestGinConfigureArgs(&args, &caps);
|
||||
|
||||
ginBenchmark_t bench = {};
|
||||
|
||||
@@ -6,7 +6,7 @@
|
||||
#include "gin_context.h"
|
||||
|
||||
template <bool skipCreditCheck, bool aggregateRequests, ncclGinResourceSharingMode rsm>
|
||||
__global__ void ginPutBwKernel(ncclDevComm comm, ncclDevResourceHandle hBuf, int iters, size_t numElems,
|
||||
__global__ void ginPutBwKernel(ncclDevComm comm, ncclDevResourceHandle devBufHandle, int iters, size_t numElems,
|
||||
int queueDepth, size_t maxElems) {
|
||||
#if __CUDA_ARCH__ >= 700
|
||||
const int tag = blockIdx.x;
|
||||
@@ -15,7 +15,7 @@ __global__ void ginPutBwKernel(ncclDevComm comm, ncclDevResourceHandle hBuf, int
|
||||
ncclGin gin(comm, tag, rsm);
|
||||
|
||||
const size_t slots = maxElems / numElems;
|
||||
ncclSymPtr<int> buf = (ncclSymPtr<int>)ncclGetResourceBuffer(comm, hBuf);
|
||||
ncclSymPtr<int> buf = (ncclSymPtr<int>)ncclGetResourceBuffer(comm, devBufHandle);
|
||||
buf += (size_t)(threadIdx.x % slots) * numElems;
|
||||
ncclSymPtr<int> sbuf = buf;
|
||||
ncclSymPtr<int> dbuf = buf;
|
||||
@@ -55,12 +55,12 @@ __global__ void ginPutBwKernel(ncclDevComm comm, ncclDevResourceHandle hBuf, int
|
||||
}
|
||||
|
||||
template <ncclGinResourceSharingMode rsm>
|
||||
static void ginPutBWLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
static void ginPutBWLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, size_t numElems, int iters) {
|
||||
const int queueDepth = args->queueDepth;
|
||||
const size_t maxElems = args->maxBytes / sizeof(int);
|
||||
#define LAUNCH_BW_PUT(SKIP, AG) \
|
||||
ginPutBwKernel<SKIP, AG, rsm><<<args->numCtas, args->numThreads, 0, stream>>>(dcomm, hBuf, iters, numElems, queueDepth, maxElems)
|
||||
ginPutBwKernel<SKIP, AG, rsm><<<args->numCtas, args->numThreads, 0, stream>>>(dcomm, devBufHandle, iters, numElems, queueDepth, maxElems)
|
||||
|
||||
if (args->ginSkipCreditCheck) {
|
||||
if (args->ginAggregateRequests) LAUNCH_BW_PUT(true, true);
|
||||
@@ -74,11 +74,11 @@ static void ginPutBWLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cud
|
||||
CUDACHECK_FATAL(cudaGetLastError());
|
||||
}
|
||||
|
||||
void ncclTestGinPutBWLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
void ncclTestGinPutBWLaunch(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, size_t numElems, int iters) {
|
||||
if (args->ginRsm == ncclGinRsmCta) {
|
||||
ginPutBWLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, hBuf, stream, args, numElems, iters);
|
||||
ginPutBWLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, devBufHandle, stream, args, numElems, iters);
|
||||
} else {
|
||||
ginPutBWLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, hBuf, stream, args, numElems, iters);
|
||||
ginPutBWLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, devBufHandle, stream, args, numElems, iters);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -16,6 +16,7 @@ int main(int argc, char** argv) {
|
||||
caps.allowAggregateRequests = true;
|
||||
caps.allowBidirectional = true;
|
||||
caps.allowMultiCtaThreads = true;
|
||||
caps.canUseHostMemory = false;
|
||||
ncclTestGinConfigureArgs(&args, &caps);
|
||||
|
||||
ginBenchmark_t bench = {};
|
||||
|
||||
@@ -6,7 +6,7 @@
|
||||
#include "gin_context.h"
|
||||
|
||||
template <bool skipCreditCheck, bool aggregateRequests, ncclGinResourceSharingMode rsm>
|
||||
__global__ void ginPutCounterBwKernel(ncclDevComm comm, ncclDevResourceHandle hBuf, int iters, size_t numElems,
|
||||
__global__ void ginPutCounterBwKernel(ncclDevComm comm, ncclDevResourceHandle devBufHandle, int iters, size_t numElems,
|
||||
int queueDepth, size_t maxElems) {
|
||||
#if __CUDA_ARCH__ >= 700
|
||||
const int tag = blockIdx.x;
|
||||
@@ -15,7 +15,7 @@ __global__ void ginPutCounterBwKernel(ncclDevComm comm, ncclDevResourceHandle hB
|
||||
ncclGin gin(comm, tag, rsm);
|
||||
|
||||
const size_t slots = maxElems / numElems;
|
||||
ncclSymPtr<int> buf = (ncclSymPtr<int>)ncclGetResourceBuffer(comm, hBuf);
|
||||
ncclSymPtr<int> buf = (ncclSymPtr<int>)ncclGetResourceBuffer(comm, devBufHandle);
|
||||
buf += (size_t)(threadIdx.x % slots) * numElems;
|
||||
ncclSymPtr<int> sbuf = buf;
|
||||
ncclSymPtr<int> dbuf = buf;
|
||||
@@ -60,12 +60,12 @@ __global__ void ginPutCounterBwKernel(ncclDevComm comm, ncclDevResourceHandle hB
|
||||
}
|
||||
|
||||
template <ncclGinResourceSharingMode rsm>
|
||||
static void ginPutCounterBWLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
static void ginPutCounterBWLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, size_t numElems, int iters) {
|
||||
const int queueDepth = args->queueDepth;
|
||||
const size_t maxElems = args->maxBytes / sizeof(int);
|
||||
#define LAUNCH_BW_PUT_COUNTER(SKIP, AG) \
|
||||
ginPutCounterBwKernel<SKIP, AG, rsm><<<args->numCtas, args->numThreads, 0, stream>>>(dcomm, hBuf, iters, numElems, queueDepth, maxElems)
|
||||
ginPutCounterBwKernel<SKIP, AG, rsm><<<args->numCtas, args->numThreads, 0, stream>>>(dcomm, devBufHandle, iters, numElems, queueDepth, maxElems)
|
||||
|
||||
if (args->ginSkipCreditCheck) {
|
||||
if (args->ginAggregateRequests) LAUNCH_BW_PUT_COUNTER(true, true);
|
||||
@@ -79,11 +79,11 @@ static void ginPutCounterBWLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle hB
|
||||
CUDACHECK_FATAL(cudaGetLastError());
|
||||
}
|
||||
|
||||
void ncclTestGinPutCounterBWLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
void ncclTestGinPutCounterBWLaunch(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, size_t numElems, int iters) {
|
||||
if (args->ginRsm == ncclGinRsmCta) {
|
||||
ginPutCounterBWLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, hBuf, stream, args, numElems, iters);
|
||||
ginPutCounterBWLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, devBufHandle, stream, args, numElems, iters);
|
||||
} else {
|
||||
ginPutCounterBWLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, hBuf, stream, args, numElems, iters);
|
||||
ginPutCounterBWLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, devBufHandle, stream, args, numElems, iters);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -7,7 +7,7 @@
|
||||
|
||||
template <bool skipCreditCheck, bool aggregateRequests, bool strongSignal,
|
||||
ncclGinResourceSharingMode rsm>
|
||||
__global__ void ginPutSignalBwKernel(ncclDevComm comm, ncclDevResourceHandle hBuf, int iters, size_t numElems,
|
||||
__global__ void ginPutSignalBwKernel(ncclDevComm comm, ncclDevResourceHandle devBufHandle, int iters, size_t numElems,
|
||||
int queueDepth, size_t maxElems) {
|
||||
#if __CUDA_ARCH__ >= 700
|
||||
const int tag = blockIdx.x;
|
||||
@@ -16,7 +16,7 @@ __global__ void ginPutSignalBwKernel(ncclDevComm comm, ncclDevResourceHandle hBu
|
||||
ncclGin gin(comm, tag, rsm);
|
||||
|
||||
const size_t slots = maxElems / numElems;
|
||||
ncclSymPtr<int> buf = (ncclSymPtr<int>)ncclGetResourceBuffer(comm, hBuf);
|
||||
ncclSymPtr<int> buf = (ncclSymPtr<int>)ncclGetResourceBuffer(comm, devBufHandle);
|
||||
buf += (size_t)(threadIdx.x % slots) * numElems;
|
||||
ncclSymPtr<int> sbuf = buf;
|
||||
ncclSymPtr<int> dbuf = buf;
|
||||
@@ -75,12 +75,12 @@ __global__ void ginPutSignalBwKernel(ncclDevComm comm, ncclDevResourceHandle hBu
|
||||
}
|
||||
|
||||
template <ncclGinResourceSharingMode rsm>
|
||||
static void ginPutSignalBWLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
static void ginPutSignalBWLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, size_t numElems, int iters) {
|
||||
const int queueDepth = args->queueDepth;
|
||||
const size_t maxElems = args->maxBytes / sizeof(int);
|
||||
#define LAUNCH_BW_PUT_SIGNAL(SKIP, AG, STRONG) \
|
||||
ginPutSignalBwKernel<SKIP, AG, STRONG, rsm><<<args->numCtas, args->numThreads, 0, stream>>>(dcomm, hBuf, iters, numElems, queueDepth, maxElems)
|
||||
ginPutSignalBwKernel<SKIP, AG, STRONG, rsm><<<args->numCtas, args->numThreads, 0, stream>>>(dcomm, devBufHandle, iters, numElems, queueDepth, maxElems)
|
||||
|
||||
if (args->ginSkipCreditCheck) {
|
||||
if (args->ginAggregateRequests) {
|
||||
@@ -104,11 +104,11 @@ static void ginPutSignalBWLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle hBu
|
||||
CUDACHECK_FATAL(cudaGetLastError());
|
||||
}
|
||||
|
||||
void ncclTestGinPutSignalBWLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
void ncclTestGinPutSignalBWLaunch(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, size_t numElems, int iters) {
|
||||
if (args->ginRsm == ncclGinRsmCta) {
|
||||
ginPutSignalBWLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, hBuf, stream, args, numElems, iters);
|
||||
ginPutSignalBWLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, devBufHandle, stream, args, numElems, iters);
|
||||
} else {
|
||||
ginPutSignalBWLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, hBuf, stream, args, numElems, iters);
|
||||
ginPutSignalBWLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, devBufHandle, stream, args, numElems, iters);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -7,7 +7,7 @@
|
||||
|
||||
template <bool skipCreditCheck, bool aggregateRequests, bool strongSignal,
|
||||
ncclGinResourceSharingMode rsm>
|
||||
__global__ void ginSignalBwKernel(ncclDevComm comm, ncclDevResourceHandle hBuf, int iters, int queueDepth) {
|
||||
__global__ void ginSignalBwKernel(ncclDevComm comm, ncclDevResourceHandle devBufHandle, int iters, int queueDepth) {
|
||||
#if __CUDA_ARCH__ >= 700
|
||||
const int tag = blockIdx.x;
|
||||
ncclTeam team = ncclTeamWorld(comm);
|
||||
@@ -66,11 +66,11 @@ __global__ void ginSignalBwKernel(ncclDevComm comm, ncclDevResourceHandle hBuf,
|
||||
}
|
||||
|
||||
template <ncclGinResourceSharingMode rsm>
|
||||
static void ginSignalBWLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
static void ginSignalBWLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, int iters) {
|
||||
const int queueDepth = args->queueDepth;
|
||||
#define LAUNCH_BW_SIGNAL(SKIP, AG, STRONG) \
|
||||
ginSignalBwKernel<SKIP, AG, STRONG, rsm><<<args->numCtas, args->numThreads, 0, stream>>>(dcomm, hBuf, iters, queueDepth)
|
||||
ginSignalBwKernel<SKIP, AG, STRONG, rsm><<<args->numCtas, args->numThreads, 0, stream>>>(dcomm, devBufHandle, iters, queueDepth)
|
||||
|
||||
if (args->ginSkipCreditCheck) {
|
||||
if (args->ginAggregateRequests) {
|
||||
@@ -94,11 +94,11 @@ static void ginSignalBWLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle hBuf,
|
||||
CUDACHECK_FATAL(cudaGetLastError());
|
||||
}
|
||||
|
||||
void ncclTestGinSignalBWLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
void ncclTestGinSignalBWLaunch(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, size_t numElems, int iters) {
|
||||
if (args->ginRsm == ncclGinRsmCta) {
|
||||
ginSignalBWLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, hBuf, stream, args, iters);
|
||||
ginSignalBWLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, devBufHandle, stream, args, iters);
|
||||
} else {
|
||||
ginSignalBWLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, hBuf, stream, args, iters);
|
||||
ginSignalBWLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, devBufHandle, stream, args, iters);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -19,6 +19,7 @@ int main(int argc, char** argv) {
|
||||
caps.allowAggregateRequests = true;
|
||||
caps.allowBidirectional = true;
|
||||
caps.allowMultiCtaThreads = true;
|
||||
caps.canUseHostMemory = false;
|
||||
ncclTestGinConfigureArgs(&args, &caps);
|
||||
|
||||
ginBenchmark_t bench = {};
|
||||
|
||||
@@ -13,15 +13,15 @@
|
||||
* ginBenchmark_t::run directly.
|
||||
*/
|
||||
|
||||
void ncclTestGinPutBWLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
void ncclTestGinPutBWLaunch(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, size_t numElems, int iters);
|
||||
void ncclTestGinPutSignalBWLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
void ncclTestGinPutSignalBWLaunch(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, size_t numElems, int iters);
|
||||
void ncclTestGinPutCounterBWLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
void ncclTestGinPutCounterBWLaunch(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, size_t numElems, int iters);
|
||||
|
||||
void ncclTestGinGetBWLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
void ncclTestGinGetBWLaunch(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, size_t numElems, int iters);
|
||||
|
||||
void ncclTestGinSignalBWLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
|
||||
void ncclTestGinSignalBWLaunch(ncclDevComm dcomm, ncclDevResourceHandle devBufHandle, cudaStream_t stream,
|
||||
const ginArgs_t* args, size_t numElems, int iters);
|
||||
|
||||
@@ -25,6 +25,8 @@ static void printUsage(const char* argv0) {
|
||||
" --gin_strong_signal Use a Strong signal instead of Weak, for tests that support both (default: off, i.e. weak)\n"
|
||||
" --gin_ag Enable ncclGinOptFlagsAggregateRequests on throughput tests (default: off)\n"
|
||||
" --gin_bd Bidirectional bandwidth: both ranks send to each other; reported metric is the sum of each rank's average (throughput tests only, default: off)\n"
|
||||
" --local_mem_type <device|host> Memory type of local rank (default: device)\n"
|
||||
" --remote_mem_type <device|host> Memory type of remote rank (default: device)\n"
|
||||
" --gin_rsm <thread|cta|gpu> GIN resource sharing mode (default: gpu)\n"
|
||||
" --gin_tx_depth <depth> GIN Send (queue) depth (default: 1024)\n",
|
||||
argv0, sizeof(int));
|
||||
@@ -87,6 +89,13 @@ static ncclGinOp_t parseOp(const char* str) {
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
|
||||
static ncclGinMemoryType_t parseMemoryType(const char* str, bool isLocal) {
|
||||
if (strcmp(str, "device") == 0) return ncclGinMemoryDevice;
|
||||
if (strcmp(str, "host") == 0) return ncclGinMemoryHost;
|
||||
fprintf(stderr, "Error: invalid value for --%s: '%s' (expected device or host)\n", isLocal ? "local_mem_type" : "remote_mem_type", str);
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
|
||||
const char* ncclTestGinOpName(ncclGinOp_t op) {
|
||||
switch (op) {
|
||||
case ncclGinOpPut: return "put";
|
||||
@@ -105,6 +114,10 @@ const char* ncclTestGinRsmName(ncclGinRsm_t rsm) {
|
||||
}
|
||||
}
|
||||
|
||||
const char* ncclTestGinMemoryTypeName(ncclGinMemoryType_t memory) {
|
||||
return memory == ncclGinMemoryHost ? "host" : "device";
|
||||
}
|
||||
|
||||
void ncclTestGinParseArgs(int argc, char** argv, ginArgs_t* args) {
|
||||
args->minBytes = sizeof(int);
|
||||
args->maxBytes = 4 * 1024 * 1024;
|
||||
@@ -118,6 +131,8 @@ void ncclTestGinParseArgs(int argc, char** argv, ginArgs_t* args) {
|
||||
args->ginStrongSignal = false;
|
||||
args->ginAggregateRequests = false;
|
||||
args->ginBidirectional = false;
|
||||
args->localMemoryType = ncclGinMemoryDevice;
|
||||
args->remoteMemoryType = ncclGinMemoryDevice;
|
||||
args->queueDepth = 1024;
|
||||
args->ginRsm = ncclGinRsmUnset;
|
||||
args->ginOp = ncclGinOpUnset;
|
||||
@@ -131,6 +146,8 @@ void ncclTestGinParseArgs(int argc, char** argv, ginArgs_t* args) {
|
||||
{"gin_strong_signal", no_argument, NULL, 0},
|
||||
{"gin_ag", no_argument, NULL, 0},
|
||||
{"gin_bd", no_argument, NULL, 0},
|
||||
{"local_mem_type", required_argument, NULL, 0},
|
||||
{"remote_mem_type", required_argument, NULL, 0},
|
||||
{"gin_rsm", required_argument, NULL, 0},
|
||||
{"gin_tx_depth", required_argument, NULL, 0},
|
||||
{"num_ctas", required_argument, NULL, 'c'},
|
||||
@@ -160,6 +177,10 @@ void ncclTestGinParseArgs(int argc, char** argv, ginArgs_t* args) {
|
||||
args->ginAggregateRequests = true;
|
||||
} else if (strcmp(longOpts[idx].name, "gin_bd") == 0) {
|
||||
args->ginBidirectional = true;
|
||||
} else if (strcmp(longOpts[idx].name, "local_mem_type") == 0) {
|
||||
args->localMemoryType = parseMemoryType(optarg, true);
|
||||
} else if (strcmp(longOpts[idx].name, "remote_mem_type") == 0) {
|
||||
args->remoteMemoryType = parseMemoryType(optarg, false);
|
||||
} else if (strcmp(longOpts[idx].name, "gin_rsm") == 0) {
|
||||
args->ginRsm = parseRsm(optarg);
|
||||
} else if (strcmp(longOpts[idx].name, "gin_tx_depth") == 0) {
|
||||
@@ -291,6 +312,10 @@ void ncclTestGinConfigureArgs(ginArgs_t* args, const ginTestCaps_t* caps) {
|
||||
fprintf(stderr, "Error: --gin_bd is only supported by throughput (bandwidth) tests\n");
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
if (!caps->canUseHostMemory && (args->localMemoryType == ncclGinMemoryHost || args->remoteMemoryType == ncclGinMemoryHost)) {
|
||||
fprintf(stderr, "Error: --remote_mem_type host/--local_mem_type host is only supported by ginGetBW_perf\n");
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
if (!caps->allowMultiCtaThreads && (args->numCtas != 1 || args->numThreads != 1)) {
|
||||
fprintf(stderr, "Error: Setting > 1 cta or thread is only supported by throughput tests\n");
|
||||
exit(EXIT_FAILURE);
|
||||
|
||||
@@ -21,6 +21,11 @@ typedef enum {
|
||||
ncclGinOpPutCount,
|
||||
} ncclGinOp_t;
|
||||
|
||||
typedef enum {
|
||||
ncclGinMemoryDevice = 0,
|
||||
ncclGinMemoryHost,
|
||||
} ncclGinMemoryType_t;
|
||||
|
||||
#define GIN_OP_BIT(op) (1u << (unsigned)(op))
|
||||
|
||||
typedef struct {
|
||||
@@ -36,6 +41,8 @@ typedef struct {
|
||||
bool ginStrongSignal;
|
||||
bool ginAggregateRequests;
|
||||
bool ginBidirectional;
|
||||
ncclGinMemoryType_t localMemoryType;
|
||||
ncclGinMemoryType_t remoteMemoryType;
|
||||
int queueDepth;
|
||||
ncclGinRsm_t ginRsm;
|
||||
ncclGinOp_t ginOp;
|
||||
@@ -59,6 +66,7 @@ typedef struct {
|
||||
* allowBidirectional – --gin_bd (throughput only).
|
||||
* allowMultiCtaThreads – -c/-t other than 1 (throughput only;
|
||||
* latency requires a single CTA and thread).
|
||||
* canUseHostMemory – --local_mem_type/--remote_mem_type options for CPU-backed memory support.
|
||||
*/
|
||||
typedef struct {
|
||||
ncclGinRsm_t defaultRsm;
|
||||
@@ -69,9 +77,11 @@ typedef struct {
|
||||
bool allowAggregateRequests;
|
||||
bool allowBidirectional;
|
||||
bool allowMultiCtaThreads;
|
||||
bool canUseHostMemory;
|
||||
} ginTestCaps_t;
|
||||
|
||||
void ncclTestGinParseArgs(int argc, char** argv, ginArgs_t* args);
|
||||
void ncclTestGinConfigureArgs(ginArgs_t* args, const ginTestCaps_t* caps);
|
||||
const char* ncclTestGinOpName(ncclGinOp_t op);
|
||||
const char* ncclTestGinRsmName(ncclGinRsm_t rsm);
|
||||
const char* ncclTestGinMemoryTypeName(ncclGinMemoryType_t memory);
|
||||
|
||||
@@ -13,6 +13,7 @@ static void printConfiguration(const ginArgs_t* args, const ginBenchmark_t* benc
|
||||
if (args->ginSkipCreditCheck) printf(", skip_credit_check");
|
||||
if (args->ginAggregateRequests) printf(", aggregate_requests");
|
||||
if (args->ginBidirectional) printf(", bidirectional");
|
||||
printf(", local mem type=%s, remote mem type=%s", ncclTestGinMemoryTypeName(args->localMemoryType), ncclTestGinMemoryTypeName(args->remoteMemoryType));
|
||||
if (bench->type == GIN_BENCHMARK_TYPE_THROUGHPUT) printf(", ctas=%d, threads=%d", args->numCtas, args->numThreads);
|
||||
printf("\n");
|
||||
}
|
||||
@@ -112,7 +113,7 @@ void ncclTestGinPerfRun(int argc, char** argv, const ginArgs_t* args, const ginB
|
||||
ncclComm_t comm;
|
||||
NCCLCHECK_FATAL(ncclCommInitRankConfig(&comm, nRanks, id, rank, &config));
|
||||
|
||||
ginContext_t ctx;
|
||||
ginContext_t ctx = {};
|
||||
ncclTestGinDevCommCreate(comm, args, &ctx);
|
||||
|
||||
const bool bidir = args->ginBidirectional;
|
||||
@@ -135,7 +136,7 @@ void ncclTestGinPerfRun(int argc, char** argv, const ginArgs_t* args, const ginB
|
||||
|
||||
//-----Warmup-----
|
||||
if (args->warmupIters > 0 && participates) {
|
||||
bench->run(ctx.dcomm, ctx.hBuf, stream, args, numElems, args->warmupIters);
|
||||
bench->run(ctx.dcomm, ctx.devBufHandle, ctx.hostBufWindow, stream, args, numElems, args->warmupIters);
|
||||
CUDACHECK_FATAL(cudaStreamSynchronize(stream));
|
||||
}
|
||||
|
||||
@@ -143,7 +144,7 @@ void ncclTestGinPerfRun(int argc, char** argv, const ginArgs_t* args, const ginB
|
||||
MPICHECK_FATAL(MPI_Barrier(MPI_COMM_WORLD));
|
||||
|
||||
if (measures) CUDACHECK_FATAL(cudaEventRecord(start, stream));
|
||||
if (participates) bench->run(ctx.dcomm, ctx.hBuf, stream, args, numElems, args->iters);
|
||||
if (participates) bench->run(ctx.dcomm, ctx.devBufHandle, ctx.hostBufWindow, stream, args, numElems, args->iters);
|
||||
if (measures) CUDACHECK_FATAL(cudaEventRecord(stop, stream));
|
||||
if (participates) CUDACHECK_FATAL(cudaStreamSynchronize(stream));
|
||||
|
||||
|
||||
@@ -32,6 +32,16 @@
|
||||
} \
|
||||
} while (0)
|
||||
|
||||
#define CUCHECK(cmd) do { \
|
||||
CUresult _e = (cmd); \
|
||||
if (_e != CUDA_SUCCESS) { \
|
||||
const char* _s = nullptr; \
|
||||
cuGetErrorString(_e, &_s); \
|
||||
fprintf(stderr, "Failed: CUDA driver error %s:%d '%s'\n", __FILE__, __LINE__, _s); \
|
||||
return ncclUnhandledCudaError; \
|
||||
} \
|
||||
} while (0)
|
||||
|
||||
#define NCCLCHECK(cmd) do { \
|
||||
ncclResult_t _r = (cmd); \
|
||||
if (_r != ncclSuccess) { \
|
||||
@@ -73,7 +83,7 @@
|
||||
* Benchmark callback invoked by ncclTestGinPerfRun for every (size, iters) point.
|
||||
*
|
||||
* dcomm – device comm handle
|
||||
* hBuf – registered buffer handle
|
||||
* devBufHandle – registered buffer handle
|
||||
* stream – CUDA stream to launch into
|
||||
* args – parsed CLI args
|
||||
* numElems – elements (ints) to transfer this size point
|
||||
@@ -81,7 +91,8 @@
|
||||
*/
|
||||
typedef void (*ginRunFn_t)(
|
||||
ncclDevComm dcomm,
|
||||
ncclDevResourceHandle hBuf,
|
||||
ncclDevResourceHandle devBufHandle,
|
||||
ncclWindow_t hostBufWindow,
|
||||
cudaStream_t stream,
|
||||
const ginArgs_t* args,
|
||||
size_t numElems,
|
||||
|
||||
@@ -1,6 +1,70 @@
|
||||
#include "gin_context.h"
|
||||
#include "common.h"
|
||||
|
||||
#include <cuda.h>
|
||||
|
||||
static ncclResult_t ncclTestGinHostAlloc(void** ptr, size_t size) {
|
||||
#if CUDART_VERSION >= 12020
|
||||
size_t granularity = 0;
|
||||
CUdevice currentDev;
|
||||
CUmemAllocationProp prop = {};
|
||||
CUmemAccessDesc accessDesc = {};
|
||||
CUmemGenericAllocationHandle handle;
|
||||
int cudaDev, cpuNumaNodeId = -1;
|
||||
|
||||
CUDACHECK(cudaGetDevice(&cudaDev));
|
||||
CUCHECK(cuDeviceGet(¤tDev, cudaDev));
|
||||
CUCHECK(cuDeviceGetAttribute(&cpuNumaNodeId, CU_DEVICE_ATTRIBUTE_HOST_NUMA_ID, currentDev));
|
||||
|
||||
if (cpuNumaNodeId < 0) cpuNumaNodeId = 0;
|
||||
prop.location.type = CU_MEM_LOCATION_TYPE_HOST_NUMA;
|
||||
prop.location.id = cpuNumaNodeId;
|
||||
prop.type = CU_MEM_ALLOCATION_TYPE_PINNED;
|
||||
prop.requestedHandleTypes = CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR;
|
||||
CUCHECK(cuMemGetAllocationGranularity(&granularity, &prop, CU_MEM_ALLOC_GRANULARITY_MINIMUM));
|
||||
|
||||
size = ((size + granularity - 1) / granularity) * granularity;
|
||||
CUCHECK(cuMemCreate(&handle, size, &prop, 0));
|
||||
CUCHECK(cuMemAddressReserve((CUdeviceptr*)ptr, size, granularity, 0, 0));
|
||||
CUCHECK(cuMemMap((CUdeviceptr)*ptr, size, 0, handle, 0));
|
||||
|
||||
accessDesc.location.type = CU_MEM_LOCATION_TYPE_DEVICE;
|
||||
accessDesc.location.id = cudaDev;
|
||||
accessDesc.flags = CU_MEM_ACCESS_FLAGS_PROT_READWRITE;
|
||||
CUCHECK(cuMemSetAccess((CUdeviceptr)*ptr, size, &accessDesc, 1));
|
||||
|
||||
accessDesc.location.type = CU_MEM_LOCATION_TYPE_HOST_NUMA;
|
||||
accessDesc.location.id = cpuNumaNodeId;
|
||||
CUCHECK(cuMemSetAccess((CUdeviceptr)*ptr, size, &accessDesc, 1));
|
||||
|
||||
return ncclSuccess;
|
||||
#else /* CUDART_VERSION >= 12020 */
|
||||
(void)ptr;
|
||||
(void)size;
|
||||
return ncclSystemError;
|
||||
#endif /* CUDART_VERSION >= 12020 */
|
||||
}
|
||||
|
||||
static ncclResult_t ncclTestGinHostFree(void* ptr) {
|
||||
#if CUDART_VERSION >= 12020
|
||||
if (ptr == nullptr) return ncclSuccess;
|
||||
|
||||
CUmemGenericAllocationHandle handle;
|
||||
size_t size = 0;
|
||||
|
||||
CUCHECK(cuMemRetainAllocationHandle(&handle, ptr));
|
||||
CUCHECK(cuMemRelease(handle));
|
||||
CUCHECK(cuMemGetAddressRange(nullptr, &size, (CUdeviceptr)ptr));
|
||||
CUCHECK(cuMemUnmap((CUdeviceptr)ptr, size));
|
||||
CUCHECK(cuMemRelease(handle));
|
||||
CUCHECK(cuMemAddressFree((CUdeviceptr)ptr, size));
|
||||
return ncclSuccess;
|
||||
#else /* CUDART_VERSION >= 12020 */
|
||||
(void)ptr;
|
||||
return ncclSystemError;
|
||||
#endif /* CUDART_VERSION >= 12020 */
|
||||
}
|
||||
|
||||
/*
|
||||
* The build only produces these benchmarks against headers new enough for the
|
||||
* GIN requirements below, so this guards the remaining case the build system
|
||||
@@ -26,6 +90,11 @@ size_t ncclTestGinMaxBufferBytes(void) {
|
||||
void ncclTestGinDevCommCreate(ncclComm_t comm, const ginArgs_t* args, ginContext_t* ctx) {
|
||||
checkNcclVersion();
|
||||
|
||||
if (args->localMemoryType == ncclGinMemoryHost || args->remoteMemoryType == ncclGinMemoryHost) {
|
||||
NCCLCHECK_FATAL(ncclTestGinHostAlloc(&ctx->hostBuf, args->maxBytes));
|
||||
NCCLCHECK_FATAL(ncclCommWindowRegister(comm, ctx->hostBuf, args->maxBytes, &ctx->hostBufWindow, NCCL_WIN_GIN_ONLY));
|
||||
}
|
||||
|
||||
ncclDevCommRequirements_t reqs = NCCL_DEV_COMM_REQUIREMENTS_INITIALIZER;
|
||||
reqs.ginContextCount = args->numCtas;
|
||||
reqs.ginQueueDepth = args->queueDepth;
|
||||
@@ -38,7 +107,7 @@ void ncclTestGinDevCommCreate(ncclComm_t comm, const ginArgs_t* args, ginContext
|
||||
|
||||
ncclDevResourceRequirements bufReq = {};
|
||||
bufReq.bufferSize = args->maxBytes;
|
||||
bufReq.outBufferHandle = &ctx->hBuf;
|
||||
bufReq.outBufferHandle = &ctx->devBufHandle;
|
||||
bufReq.next = reqs.resourceRequirementsList;
|
||||
reqs.resourceRequirementsList = &bufReq;
|
||||
|
||||
@@ -47,4 +116,6 @@ void ncclTestGinDevCommCreate(ncclComm_t comm, const ginArgs_t* args, ginContext
|
||||
|
||||
void ncclTestGinDevCommDestroy(ncclComm_t comm, ginContext_t* ctx) {
|
||||
NCCLCHECK_FATAL(ncclDevCommDestroy(comm, &ctx->dcomm));
|
||||
if (ctx->hostBufWindow != nullptr) NCCLCHECK_FATAL(ncclCommWindowDeregister(comm, ctx->hostBufWindow));
|
||||
if (ctx->hostBuf != nullptr) NCCLCHECK_FATAL(ncclTestGinHostFree(ctx->hostBuf));
|
||||
}
|
||||
|
||||
@@ -6,7 +6,9 @@
|
||||
|
||||
typedef struct {
|
||||
ncclDevComm dcomm;
|
||||
ncclDevResourceHandle hBuf;
|
||||
ncclDevResourceHandle devBufHandle;
|
||||
void* hostBuf;
|
||||
ncclWindow_t hostBufWindow;
|
||||
} ginContext_t;
|
||||
|
||||
size_t ncclTestGinMaxBufferBytes(void);
|
||||
|
||||
Reference in New Issue
Block a user