diff --git a/src/Makefile b/src/Makefile
index 375905a..85e85c1 100644
--- a/src/Makefile
+++ b/src/Makefile
@@ -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
diff --git a/src/device_api/gin/CMakeLists.txt b/src/device_api/gin/CMakeLists.txt
index 37f741d..6497752 100644
--- a/src/device_api/gin/CMakeLists.txt
+++ b/src/device_api/gin/CMakeLists.txt
@@ -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++
         )
diff --git a/src/device_api/gin/README.md b/src/device_api/gin/README.md
index c8a693c..b099688 100644
--- a/src/device_api/gin/README.md
+++ b/src/device_api/gin/README.md
@@ -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
 
diff --git a/src/device_api/gin/latency/ginGetLatency_ping.cu b/src/device_api/gin/latency/ginGetLatency_ping.cu
index 0ddca22..cd8437a 100644
--- a/src/device_api/gin/latency/ginGetLatency_ping.cu
+++ b/src/device_api/gin/latency/ginGetLatency_ping.cu
@@ -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);
   }
 }
diff --git a/src/device_api/gin/latency/ginGetLatency_ping_main.cu b/src/device_api/gin/latency/ginGetLatency_ping_main.cu
index a9a89a8..925897f 100644
--- a/src/device_api/gin/latency/ginGetLatency_ping_main.cu
+++ b/src/device_api/gin/latency/ginGetLatency_ping_main.cu
@@ -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 = {};
diff --git a/src/device_api/gin/latency/ginLatencyTests.h b/src/device_api/gin/latency/ginLatencyTests.h
index 4536aa7..7eb952f 100644
--- a/src/device_api/gin/latency/ginLatencyTests.h
+++ b/src/device_api/gin/latency/ginLatencyTests.h
@@ -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);
diff --git a/src/device_api/gin/latency/ginPutCounterLatency_ping.cu b/src/device_api/gin/latency/ginPutCounterLatency_ping.cu
index 2355d34..e79cf4a 100644
--- a/src/device_api/gin/latency/ginPutCounterLatency_ping.cu
+++ b/src/device_api/gin/latency/ginPutCounterLatency_ping.cu
@@ -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);
   }
 }
diff --git a/src/device_api/gin/latency/ginPutLatency_ping.cu b/src/device_api/gin/latency/ginPutLatency_ping.cu
index 1466499..36f1aa9 100644
--- a/src/device_api/gin/latency/ginPutLatency_ping.cu
+++ b/src/device_api/gin/latency/ginPutLatency_ping.cu
@@ -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);
   }
 }
diff --git a/src/device_api/gin/latency/ginPutLatency_pingPong.cu b/src/device_api/gin/latency/ginPutLatency_pingPong.cu
index 4a9cf22..075dab5 100644
--- a/src/device_api/gin/latency/ginPutLatency_pingPong.cu
+++ b/src/device_api/gin/latency/ginPutLatency_pingPong.cu
@@ -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);
   }
 }
diff --git a/src/device_api/gin/latency/ginPutLatency_pingPong_main.cu b/src/device_api/gin/latency/ginPutLatency_pingPong_main.cu
index 2483558..e7f60f3 100644
--- a/src/device_api/gin/latency/ginPutLatency_pingPong_main.cu
+++ b/src/device_api/gin/latency/ginPutLatency_pingPong_main.cu
@@ -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 = {};
diff --git a/src/device_api/gin/latency/ginPutLatency_ping_main.cu b/src/device_api/gin/latency/ginPutLatency_ping_main.cu
index 1907c2e..434af13 100644
--- a/src/device_api/gin/latency/ginPutLatency_ping_main.cu
+++ b/src/device_api/gin/latency/ginPutLatency_ping_main.cu
@@ -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 = {};
diff --git a/src/device_api/gin/latency/ginPutSignalLatency_ping.cu b/src/device_api/gin/latency/ginPutSignalLatency_ping.cu
index 8480ea9..12c1c1d 100644
--- a/src/device_api/gin/latency/ginPutSignalLatency_ping.cu
+++ b/src/device_api/gin/latency/ginPutSignalLatency_ping.cu
@@ -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);
   }
 }
diff --git a/src/device_api/gin/latency/ginPutSignalLatency_pingPong.cu b/src/device_api/gin/latency/ginPutSignalLatency_pingPong.cu
index 0db6a4d..e6e922c 100644
--- a/src/device_api/gin/latency/ginPutSignalLatency_pingPong.cu
+++ b/src/device_api/gin/latency/ginPutSignalLatency_pingPong.cu
@@ -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);
   }
 }
diff --git a/src/device_api/gin/latency/ginSignalLatency_ping.cu b/src/device_api/gin/latency/ginSignalLatency_ping.cu
index f32ef35..d8fc5a1 100644
--- a/src/device_api/gin/latency/ginSignalLatency_ping.cu
+++ b/src/device_api/gin/latency/ginSignalLatency_ping.cu
@@ -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);
   }
 }
diff --git a/src/device_api/gin/latency/ginSignalLatency_pingPong.cu b/src/device_api/gin/latency/ginSignalLatency_pingPong.cu
index 89416fb..cc79dcd 100644
--- a/src/device_api/gin/latency/ginSignalLatency_pingPong.cu
+++ b/src/device_api/gin/latency/ginSignalLatency_pingPong.cu
@@ -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);
   }
 }
diff --git a/src/device_api/gin/latency/ginSignalLatency_pingPong_main.cu b/src/device_api/gin/latency/ginSignalLatency_pingPong_main.cu
index c2f2ddd..2e2d7b7 100644
--- a/src/device_api/gin/latency/ginSignalLatency_pingPong_main.cu
+++ b/src/device_api/gin/latency/ginSignalLatency_pingPong_main.cu
@@ -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 = {};
diff --git a/src/device_api/gin/latency/ginSignalLatency_ping_main.cu b/src/device_api/gin/latency/ginSignalLatency_ping_main.cu
index e18f2c8..27c1830 100644
--- a/src/device_api/gin/latency/ginSignalLatency_ping_main.cu
+++ b/src/device_api/gin/latency/ginSignalLatency_ping_main.cu
@@ -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 = {};
diff --git a/src/device_api/gin/throughput/ginGetBW.cu b/src/device_api/gin/throughput/ginGetBW.cu
index 04d5b5d..831e7e8 100644
--- a/src/device_api/gin/throughput/ginGetBW.cu
+++ b/src/device_api/gin/throughput/ginGetBW.cu
@@ -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);
   }
 }
diff --git a/src/device_api/gin/throughput/ginGetBW_main.cu b/src/device_api/gin/throughput/ginGetBW_main.cu
index 8edf622..55ebe0d 100644
--- a/src/device_api/gin/throughput/ginGetBW_main.cu
+++ b/src/device_api/gin/throughput/ginGetBW_main.cu
@@ -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 = {};
diff --git a/src/device_api/gin/throughput/ginPutBW.cu b/src/device_api/gin/throughput/ginPutBW.cu
index 623cdc0..c717734 100644
--- a/src/device_api/gin/throughput/ginPutBW.cu
+++ b/src/device_api/gin/throughput/ginPutBW.cu
@@ -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);
   }
 }
diff --git a/src/device_api/gin/throughput/ginPutBW_main.cu b/src/device_api/gin/throughput/ginPutBW_main.cu
index d871d45..507c8cf 100644
--- a/src/device_api/gin/throughput/ginPutBW_main.cu
+++ b/src/device_api/gin/throughput/ginPutBW_main.cu
@@ -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 = {};
diff --git a/src/device_api/gin/throughput/ginPutCounterBW.cu b/src/device_api/gin/throughput/ginPutCounterBW.cu
index 3a0353e..67255c9 100644
--- a/src/device_api/gin/throughput/ginPutCounterBW.cu
+++ b/src/device_api/gin/throughput/ginPutCounterBW.cu
@@ -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);
   }
 }
diff --git a/src/device_api/gin/throughput/ginPutSignalBW.cu b/src/device_api/gin/throughput/ginPutSignalBW.cu
index 9e2ce0f..e5a50b5 100644
--- a/src/device_api/gin/throughput/ginPutSignalBW.cu
+++ b/src/device_api/gin/throughput/ginPutSignalBW.cu
@@ -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);
   }
 }
diff --git a/src/device_api/gin/throughput/ginSignalBW.cu b/src/device_api/gin/throughput/ginSignalBW.cu
index b897645..381f889 100644
--- a/src/device_api/gin/throughput/ginSignalBW.cu
+++ b/src/device_api/gin/throughput/ginSignalBW.cu
@@ -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);
   }
 }
diff --git a/src/device_api/gin/throughput/ginSignalBW_main.cu b/src/device_api/gin/throughput/ginSignalBW_main.cu
index 18fda77..3814a8c 100644
--- a/src/device_api/gin/throughput/ginSignalBW_main.cu
+++ b/src/device_api/gin/throughput/ginSignalBW_main.cu
@@ -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 = {};
diff --git a/src/device_api/gin/throughput/ginThroughputTests.h b/src/device_api/gin/throughput/ginThroughputTests.h
index ade39bb..cae9bcc 100644
--- a/src/device_api/gin/throughput/ginThroughputTests.h
+++ b/src/device_api/gin/throughput/ginThroughputTests.h
@@ -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);
diff --git a/src/device_api/gin/utils/args.cc b/src/device_api/gin/utils/args.cc
index c842d60..fe8217f 100644
--- a/src/device_api/gin/utils/args.cc
+++ b/src/device_api/gin/utils/args.cc
@@ -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);
diff --git a/src/device_api/gin/utils/args.h b/src/device_api/gin/utils/args.h
index 275be12..6add6c7 100644
--- a/src/device_api/gin/utils/args.h
+++ b/src/device_api/gin/utils/args.h
@@ -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);
diff --git a/src/device_api/gin/utils/common.cc b/src/device_api/gin/utils/common.cc
index a6557c9..8e0c03a 100644
--- a/src/device_api/gin/utils/common.cc
+++ b/src/device_api/gin/utils/common.cc
@@ -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));
 
diff --git a/src/device_api/gin/utils/common.h b/src/device_api/gin/utils/common.h
index 5273586..c952029 100644
--- a/src/device_api/gin/utils/common.h
+++ b/src/device_api/gin/utils/common.h
@@ -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,
diff --git a/src/device_api/gin/utils/gin_context.cc b/src/device_api/gin/utils/gin_context.cc
index 53d7369..9bd1355 100644
--- a/src/device_api/gin/utils/gin_context.cc
+++ b/src/device_api/gin/utils/gin_context.cc
@@ -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(&currentDev, 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));
 }
diff --git a/src/device_api/gin/utils/gin_context.h b/src/device_api/gin/utils/gin_context.h
index ad8cb94..af8607f 100644
--- a/src/device_api/gin/utils/gin_context.h
+++ b/src/device_api/gin/utils/gin_context.h
@@ -6,7 +6,9 @@
 
 typedef struct {
   ncclDevComm dcomm;
-  ncclDevResourceHandle hBuf;
+  ncclDevResourceHandle devBufHandle;
+  void* hostBuf;
+  ncclWindow_t hostBufWindow;
 } ginContext_t;
 
 size_t ncclTestGinMaxBufferBytes(void);