I have a program where I use unified memory and then launch multiple kernels on multiple GPUs, each will process subset of an array let’s say, it is working well on v100, h100 but not on RTX6000
#include <cuda_runtime.h>
#include <cstdlib>
#include <iostream>
#define CHECK_CUDA(call) \
do { \
cudaError_t err__ = (call); \
if (err__ != cudaSuccess) { \
std::cerr << "CUDA error at " << __FILE__ << ":" << __LINE__ \
<< " : " << cudaGetErrorString(err__) << std::endl; \
std::exit(EXIT_FAILURE); \
} \
} while (0)
static void prefetch_to_device(const void* ptr,
size_t bytes,
int device,
cudaStream_t stream = 0)
{
#if CUDART_VERSION >= 13000
cudaMemLocation loc{};
loc.type = cudaMemLocationTypeDevice;
loc.id = device;
CHECK_CUDA(cudaMemPrefetchAsync(ptr, bytes, loc, 0, stream));
#else
CHECK_CUDA(cudaMemPrefetchAsync(ptr, bytes, device, stream));
#endif
}
static void prefetch_to_cpu(const void* ptr,
size_t bytes,
cudaStream_t stream = 0)
{
#if CUDART_VERSION >= 13000
cudaMemLocation loc{};
loc.type = cudaMemLocationTypeHost;
loc.id = 0;
CHECK_CUDA(cudaMemPrefetchAsync(ptr, bytes, loc, 0, stream));
#else
CHECK_CUDA(cudaMemPrefetchAsync(ptr, bytes, cudaCpuDeviceId, stream));
#endif
}
static void advise_preferred_device(const void* ptr,
size_t bytes,
int device)
{
#if CUDART_VERSION >= 13000
cudaMemLocation loc{};
loc.type = cudaMemLocationTypeDevice;
loc.id = device;
// CHECK_CUDA(cudaMemAdvise(ptr,
// bytes,
// cudaMemAdviseSetPreferredLocation,
// loc));
#else
CHECK_CUDA(cudaMemAdvise(ptr,
bytes,
cudaMemAdviseSetPreferredLocation,
device));
#endif
}
__global__ void init_on_gpu0(int *a, int n)
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n)
a[i] = i;
}
__global__ void use_on_gpu1(int *a, int n)
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n)
a[i] = a[i] + 1000;
}
__global__ void check_on_gpu1(const int *a, int n, int *errors)
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n)
{
int expected = i + 1000;
if (a[i] != expected)
atomicAdd(errors, 1);
}
}
static void print_device_props(int dev)
{
cudaDeviceProp p{};
CHECK_CUDA(cudaGetDeviceProperties(&p, dev));
std::cout << "GPU" << dev << ": " << p.name << "\n"
<< " unifiedAddressing=" << p.unifiedAddressing << "\n"
<< " managedMemory=" << p.managedMemory << "\n"
<< " concurrentManagedAccess=" << p.concurrentManagedAccess << "\n"
<< " pageableMemoryAccess=" << p.pageableMemoryAccess << "\n"
<< " pageableMemoryAccessUsesHostPageTables="
<< p.pageableMemoryAccessUsesHostPageTables << "\n";
}
static void print_ptr_attr(const char *name, const void *ptr)
{
cudaPointerAttributes attr{};
cudaError_t err = cudaPointerGetAttributes(&attr, ptr);
if (err != cudaSuccess)
{
std::cout << name << " attr failed: "
<< cudaGetErrorString(err) << std::endl;
cudaGetLastError();
return;
}
std::cout << name
<< " ptr=" << ptr
<< " type=" << static_cast<int>(attr.type)
<< " device=" << attr.device
<< " devicePointer=" << attr.devicePointer
<< " hostPointer=" << attr.hostPointer
<< std::endl;
}
int main()
{
int device_count = 0;
CHECK_CUDA(cudaGetDeviceCount(&device_count));
if (device_count < 2)
{
std::cerr << "Need at least 2 GPUs\n";
return EXIT_FAILURE;
}
const int gpu0 = 0;
const int gpu1 = 1;
print_device_props(gpu0);
print_device_props(gpu1);
int can01 = 0;
int can10 = 0;
CHECK_CUDA(cudaDeviceCanAccessPeer(&can01, gpu0, gpu1));
CHECK_CUDA(cudaDeviceCanAccessPeer(&can10, gpu1, gpu0));
std::cout << "P2P GPU0 -> GPU1 can=" << can01 << "\n";
std::cout << "P2P GPU1 -> GPU0 can=" << can10 << "\n";
if (can01)
{
CHECK_CUDA(cudaSetDevice(gpu0));
cudaError_t e = cudaDeviceEnablePeerAccess(gpu1, 0);
if (e == cudaErrorPeerAccessAlreadyEnabled)
cudaGetLastError();
else
CHECK_CUDA(e);
}
if (can10)
{
CHECK_CUDA(cudaSetDevice(gpu1));
cudaError_t e = cudaDeviceEnablePeerAccess(gpu0, 0);
if (e == cudaErrorPeerAccessAlreadyEnabled)
cudaGetLastError();
else
CHECK_CUDA(e);
}
constexpr int N = 1 << 20;
constexpr int BLOCK = 256;
const int GRID = (N + BLOCK - 1) / BLOCK;
int *a = nullptr;
int *errors = nullptr;
CHECK_CUDA(cudaSetDevice(gpu0));
CHECK_CUDA(cudaMallocManaged(&a, N * sizeof(int)));
CHECK_CUDA(cudaMallocManaged(&errors, sizeof(int)));
print_ptr_attr("a", a);
print_ptr_attr("errors", errors);
advise_preferred_device(a, N * sizeof(int), gpu0);
advise_preferred_device(errors, sizeof(int), gpu0);
std::cout << "\nPrefetching a/errors to GPU0...\n";
prefetch_to_device(a, N * sizeof(int), gpu0);
prefetch_to_device(errors, sizeof(int), gpu0);
CHECK_CUDA(cudaSetDevice(gpu0));
CHECK_CUDA(cudaDeviceSynchronize());
std::cout << "Initializing on GPU0...\n";
init_on_gpu0<<<GRID, BLOCK>>>(a, N);
CHECK_CUDA(cudaGetLastError());
CHECK_CUDA(cudaDeviceSynchronize());
CHECK_CUDA(cudaSetDevice(gpu1));
prefetch_to_device(a, N * sizeof(int), gpu1);
prefetch_to_device(errors, sizeof(int), gpu1);
CHECK_CUDA(cudaDeviceSynchronize());
std::cout << "Using on GPU1...\n";
use_on_gpu1<<<GRID, BLOCK>>>(a, N);
CHECK_CUDA(cudaGetLastError());
CHECK_CUDA(cudaDeviceSynchronize());
CHECK_CUDA(cudaMemset(errors, 0, sizeof(int)));
check_on_gpu1<<<GRID, BLOCK>>>(a, N, errors);
CHECK_CUDA(cudaGetLastError());
CHECK_CUDA(cudaDeviceSynchronize());
std::cout << "Prefetching results to CPU...\n";
prefetch_to_cpu(errors, sizeof(int));
prefetch_to_cpu(a, N * sizeof(int));
CHECK_CUDA(cudaSetDevice(gpu1));
CHECK_CUDA(cudaDeviceSynchronize());
if (*errors == 0)
{
std::cout << "PASS: GPU1 correctly read and modified unified memory\n";
}
else
{
std::cout << "FAIL: errors=" << *errors << "\n";
}
std::cout << "a[0]=" << a[0] << "\n";
std::cout << "a[1]=" << a[1] << "\n";
std::cout << "a[123]=" << a[123] << "\n";
std::cout << "a[N-1]=" << a[N - 1] << "\n";
CHECK_CUDA(cudaFree(a));
CHECK_CUDA(cudaFree(errors));
return 0;
}
On v100
GPU0: Tesla V100S-PCIE-32GB
unifiedAddressing=1
managedMemory=1
concurrentManagedAccess=1
pageableMemoryAccess=0
pageableMemoryAccessUsesHostPageTables=0
GPU1: Tesla V100S-PCIE-32GB
unifiedAddressing=1
managedMemory=1
concurrentManagedAccess=1
pageableMemoryAccess=0
pageableMemoryAccessUsesHostPageTables=0
P2P GPU0 -> GPU1 can=1
P2P GPU1 -> GPU0 can=1
a ptr=0x7f1e54000000 type=3 device=0 devicePointer=0x7f1e54000000 hostPointer=0x7f1e54000000
errors ptr=0x7f1e54400000 type=3 device=0 devicePointer=0x7f1e54400000 hostPointer=0x7f1e54400000
Prefetching a/errors to GPU0...
Initializing on GPU0...
Prefetching a/errors to GPU1...
Using on GPU1...
Prefetching results to CPU...
PASS: GPU1 correctly read and modified unified memory
a[0]=1000
a[1]=1001
a[123]=1123
a[N-1]=1049575
on RTX6000Pro
GPU0: NVIDIA RTX PRO 6000 Blackwell Server Edition
unifiedAddressing=1
managedMemory=1
concurrentManagedAccess=1
pageableMemoryAccess=0
pageableMemoryAccessUsesHostPageTables=0
GPU1: NVIDIA RTX PRO 6000 Blackwell Server Edition
unifiedAddressing=1
managedMemory=1
concurrentManagedAccess=1
pageableMemoryAccess=0
pageableMemoryAccessUsesHostPageTables=0
P2P GPU0 -> GPU1 can=1
P2P GPU1 -> GPU0 can=1
a ptr=0x7f79a4000000 type=3 device=0 devicePointer=0x7f79a4000000 hostPointer=0x7f79a4000000
errors ptr=0x7f79a4400000 type=3 device=0 devicePointer=0x7f79a4400000 hostPointer=0x7f79a4400000
Prefetching a/errors to GPU0...
Initializing on GPU0...
Using on GPU1...
Prefetching results to CPU...
FAIL: errors=1048575
a[0]=1000
a[1]=1000
a[123]=1000
a[N-1]=1000
What I can see is that GPU1 cannot see the updates done by GPU0.