From a8eca04c3f8daf0e44c79ade592ceb683d86cbf9 Mon Sep 17 00:00:00 2001 From: Anav Prasad Date: Thu, 25 Jun 2026 10:26:56 +0000 Subject: [PATCH] support cuda virtual devices --- ggml/src/ggml-cuda/common.cuh | 8 +- ggml/src/ggml-cuda/ggml-cuda.cu | 161 +++++++++++++++++++++++++------- 2 files changed, 134 insertions(+), 35 deletions(-) diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index 290dc4aff2..3f9a4f496f 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -1115,7 +1115,8 @@ struct ggml_cuda_type_traits { ////////////////////// struct ggml_cuda_device_info { - int device_count; + int device_count; // number of (possibly virtual) devices exposed to the rest of ggml + int physical_device_count; // number of physical CUDA devices actually present struct cuda_device_info { int cc; // compute capability @@ -1128,6 +1129,8 @@ struct ggml_cuda_device_info { size_t total_vram; int warp_size; // Number of threads in a dispatch bool supports_cooperative_launch; // whether cooperative launch is supported + int physical_device; // backing physical CUDA device for this (virtual) device + int physical_share_count; // number of (virtual) devices sharing this device's physical GPU }; cuda_device_info devices[GGML_CUDA_MAX_DEVICES] = {}; @@ -1140,6 +1143,9 @@ const ggml_cuda_device_info & ggml_cuda_info(); void ggml_cuda_set_device(int device); int ggml_cuda_get_device(); +// map a (possibly virtual) device id to the physical CUDA device that backs it +int ggml_cuda_get_physical_device(int device); + struct ggml_cuda_pool { virtual ~ggml_cuda_pool() = default; diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index f6632498c5..9245313c74 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -107,14 +107,17 @@ void ggml_cuda_error(const char * stmt, const char * func, const char * file, in // this is faster on Windows // probably because the Windows CUDA libraries forget to make this check before invoking the drivers void ggml_cuda_set_device(int device) { + // translate the (possibly virtual) device id to the physical CUDA device that backs it + const int physical_device = ggml_cuda_get_physical_device(device); + int current_device; CUDA_CHECK(cudaGetDevice(¤t_device)); - if (device == current_device) { + if (physical_device == current_device) { return; } - CUDA_CHECK(cudaSetDevice(device)); + CUDA_CHECK(cudaSetDevice(physical_device)); } int ggml_cuda_get_device() { @@ -123,6 +126,12 @@ int ggml_cuda_get_device() { return id; } +int ggml_cuda_get_physical_device(int device) { + const ggml_cuda_device_info & info = ggml_cuda_info(); + GGML_ASSERT(device >= 0 && device < info.device_count); + return info.devices[device].physical_device; +} + static cudaError_t ggml_cuda_device_malloc(void ** ptr, size_t size, int device) { ggml_cuda_set_device(device); cudaError_t err; @@ -205,48 +214,90 @@ static int ggml_cuda_parse_id(char devName[]) { static ggml_cuda_device_info ggml_cuda_init() { ggml_cuda_device_info info = {}; - cudaError_t err = cudaGetDeviceCount(&info.device_count); + cudaError_t err = cudaGetDeviceCount(&info.physical_device_count); if (err != cudaSuccess) { GGML_LOG_ERROR("%s: failed to initialize " GGML_CUDA_NAME ": %s\n", __func__, cudaGetErrorString(err)); return info; } - GGML_ASSERT(info.device_count <= GGML_CUDA_MAX_DEVICES); + GGML_ASSERT(info.physical_device_count <= GGML_CUDA_MAX_DEVICES); + + // by default expose exactly the physical devices; GGML_CUDA_DEVICES can request a different + // number of (virtual) devices to emulate multi-GPU systems on a machine with fewer GPUs + info.device_count = info.physical_device_count; + + const char * devices_env = getenv("GGML_CUDA_DEVICES"); + if (devices_env != nullptr && info.physical_device_count > 0) { + const int requested = atoi(devices_env); + if (requested > 0) { + info.device_count = requested; + } else { + GGML_LOG_WARN("%s: ignoring invalid GGML_CUDA_DEVICES=\"%s\"\n", __func__, devices_env); + } + } + + if (info.device_count > GGML_CUDA_MAX_DEVICES) { + GGML_LOG_WARN("%s: requested %d devices, clamping to GGML_CUDA_MAX_DEVICES=%d\n", + __func__, info.device_count, GGML_CUDA_MAX_DEVICES); + info.device_count = GGML_CUDA_MAX_DEVICES; + } + + // map each (virtual) device to a backing physical device (round-robin) and + // store virtual device count per physical device + int physical_share_count[GGML_CUDA_MAX_DEVICES] = {}; + GGML_ASSERT(info.device_count == 0 || info.physical_device_count > 0); + for (int id = 0; id < info.device_count; ++id) { + info.devices[id].physical_device = id % info.physical_device_count; + physical_share_count[info.devices[id].physical_device]++; + } int64_t total_vram = 0; - for (int id = 0; id < info.device_count; ++id) { + for (int id = 0; id < info.physical_device_count; ++id) { cudaDeviceProp prop; CUDA_CHECK(cudaGetDeviceProperties(&prop, id)); total_vram += prop.totalGlobalMem; } GGML_LOG_INFO("%s: found %d " GGML_CUDA_NAME " devices (Total VRAM: %zu MiB):\n", - __func__, info.device_count, (size_t)(total_vram / (1024 * 1024))); + __func__, info.physical_device_count, (size_t)(total_vram / (1024 * 1024))); + if (info.device_count != info.physical_device_count) { + GGML_LOG_INFO("%s: emulating %d virtual device(s) on %d physical device(s) (GGML_CUDA_DEVICES)\n", + __func__, info.device_count, info.physical_device_count); + } total_vram = 0; std::vector> turing_devices_without_mma; for (int id = 0; id < info.device_count; ++id) { + const int physical_id = info.devices[id].physical_device; + int device_vmm = 0; #if defined(GGML_USE_VMM) CUdevice device; - CU_CHECK(cuDeviceGet(&device, id)); + CU_CHECK(cuDeviceGet(&device, physical_id)); CU_CHECK(cuDeviceGetAttribute(&device_vmm, CU_DEVICE_ATTRIBUTE_VIRTUAL_MEMORY_MANAGEMENT_SUPPORTED, device)); if (device_vmm) { CUmemAllocationProp alloc_prop = {}; alloc_prop.type = CU_MEM_ALLOCATION_TYPE_PINNED; alloc_prop.location.type = CU_MEM_LOCATION_TYPE_DEVICE; - alloc_prop.location.id = id; + alloc_prop.location.id = physical_id; CU_CHECK(cuMemGetAllocationGranularity(&info.devices[id].vmm_granularity, &alloc_prop, CU_MEM_ALLOC_GRANULARITY_RECOMMENDED)); } #endif // defined(GGML_USE_VMM) info.devices[id].vmm = !!device_vmm; cudaDeviceProp prop; - CUDA_CHECK(cudaGetDeviceProperties(&prop, id)); + CUDA_CHECK(cudaGetDeviceProperties(&prop, physical_id)); + + // a virtual device owns only a share of its physical GPU's memory; report that share so the + // logged per-device VRAM sums to the physical total above. + GGML_ASSERT(physical_share_count[physical_id] > 0); + info.devices[id].physical_share_count = physical_share_count[physical_id]; + const size_t device_vram = prop.totalGlobalMem / info.devices[id].physical_share_count; + const size_t device_vram_mib = device_vram / (1024 * 1024); info.default_tensor_split[id] = total_vram; - total_vram += prop.totalGlobalMem; + total_vram += device_vram; info.devices[id].integrated = false; // Temporarily disabled due to issues with corrupted output (e.g. #15034) info.devices[id].nsm = prop.multiProcessorCount; info.devices[id].smpb = prop.sharedMemPerBlock; @@ -254,7 +305,7 @@ static ggml_cuda_device_info ggml_cuda_init() { #ifndef GGML_USE_MUSA int supports_coop_launch = 0; - CUDA_CHECK(cudaDeviceGetAttribute(&supports_coop_launch, cudaDevAttrCooperativeLaunch, id)); + CUDA_CHECK(cudaDeviceGetAttribute(&supports_coop_launch, cudaDevAttrCooperativeLaunch, physical_id)); info.devices[id].supports_cooperative_launch = !!supports_coop_launch; #else info.devices[id].supports_cooperative_launch = false; @@ -277,7 +328,7 @@ static ggml_cuda_device_info ggml_cuda_init() { GGML_LOG_INFO(" Device %d: %s, %s (0x%x), VMM: %s, Wave Size: %d, VRAM: %zu MiB\n", id, prop.name, prop.gcnArchName, info.devices[id].cc & 0xffff, device_vmm ? "yes" : "no", prop.warpSize, - (size_t)(prop.totalGlobalMem / (1024 * 1024))); + device_vram_mib); #elif defined(GGML_USE_MUSA) // FIXME: Ensure compatibility with varying warp sizes across different MUSA archs. info.devices[id].warp_size = 32; @@ -286,13 +337,13 @@ static ggml_cuda_device_info ggml_cuda_init() { info.devices[id].cc += prop.minor * 0x10; GGML_LOG_INFO(" Device %d: %s, compute capability %d.%d, VMM: %s, VRAM: %zu MiB\n", id, prop.name, prop.major, prop.minor, device_vmm ? "yes" : "no", - (size_t)(prop.totalGlobalMem / (1024 * 1024))); + device_vram_mib); #else info.devices[id].smpbo = prop.sharedMemPerBlockOptin; info.devices[id].cc = 100*prop.major + 10*prop.minor; GGML_LOG_INFO(" Device %d: %s, compute capability %d.%d, VMM: %s, VRAM: %zu MiB\n", id, prop.name, prop.major, prop.minor, device_vmm ? "yes" : "no", - (size_t)(prop.totalGlobalMem / (1024 * 1024))); + device_vram_mib); std::string device_name(prop.name); if (device_name == "NVIDIA GeForce MX450") { turing_devices_without_mma.push_back({ id, device_name }); @@ -307,7 +358,7 @@ static ggml_cuda_device_info ggml_cuda_init() { // TODO: Check for future drivers the default scheduling strategy and // remove this call again when cudaDeviceScheduleSpin is default. if (prop.major == 12 && prop.minor == 1) { - CUDA_CHECK(cudaSetDevice(id)); + CUDA_CHECK(cudaSetDevice(physical_id)); CUDA_CHECK(cudaSetDeviceFlags(cudaDeviceScheduleSpin)); } @@ -332,9 +383,9 @@ static ggml_cuda_device_info ggml_cuda_init() { // CUBLAS_CHECK(cublasLoggerConfigure(1, 1, 0, nullptr)); if (getenv("GGML_CUDA_P2P") != nullptr) { - for (int id = 0; id < info.device_count; ++id) { - ggml_cuda_set_device(id); - for (int id_other = 0; id_other < info.device_count; ++id_other) { + for (int id = 0; id < info.physical_device_count; ++id) { + CUDA_CHECK(cudaSetDevice(id)); + for (int id_other = 0; id_other < info.physical_device_count; ++id_other) { if (id == id_other) { continue; } @@ -479,6 +530,7 @@ struct ggml_cuda_pool_vmm : public ggml_cuda_pool { static const size_t CUDA_POOL_VMM_MAX_SIZE = 1ull << 35; // 32 GB int device; + int physical_device; CUdeviceptr pool_addr = 0; size_t pool_used = 0; size_t pool_size = 0; @@ -489,6 +541,7 @@ struct ggml_cuda_pool_vmm : public ggml_cuda_pool { explicit ggml_cuda_pool_vmm(int device) : device(device), + physical_device(ggml_cuda_get_physical_device(device)), granularity(ggml_cuda_info().devices[device].vmm_granularity) { } @@ -524,7 +577,7 @@ struct ggml_cuda_pool_vmm : public ggml_cuda_pool { CUmemAllocationProp prop = {}; prop.type = CU_MEM_ALLOCATION_TYPE_PINNED; prop.location.type = CU_MEM_LOCATION_TYPE_DEVICE; - prop.location.id = device; + prop.location.id = physical_device; CUmemGenericAllocationHandle handle; CU_CHECK(cuMemCreate(&handle, reserve_size, &prop, 0)); @@ -553,20 +606,28 @@ struct ggml_cuda_pool_vmm : public ggml_cuda_pool { // NCCL implicitly enables peer access (cudaDeviceEnablePeerAccess), and // GGML_CUDA_P2P enables it explicitly. Unlike cudaMalloc buffers, VMM // allocations do not become peer-accessible from that alone, so access - // must be granted explicitly here. + // must be granted explicitly here. With virtual devices, grant access + // on the backing *physical* devices (deduplicated, since several + // virtual devices can map to the same physical GPU). std::vector access_descs; + bool physical_seen[GGML_CUDA_MAX_DEVICES] = {}; const int device_count = ggml_cuda_info().device_count; for (int id = 0; id < device_count; ++id) { - if (id != device) { + const int id_physical = ggml_cuda_get_physical_device(id); + if (id_physical != physical_device) { int can_access_peer = 0; - CUDA_CHECK(cudaDeviceCanAccessPeer(&can_access_peer, id, device)); + CUDA_CHECK(cudaDeviceCanAccessPeer(&can_access_peer, id_physical, physical_device)); if (!can_access_peer) { continue; } } + if (physical_seen[id_physical]) { + continue; + } + physical_seen[id_physical] = true; CUmemAccessDesc access = {}; access.location.type = CU_MEM_LOCATION_TYPE_DEVICE; - access.location.id = id; + access.location.id = id_physical; access.flags = CU_MEM_ACCESS_FLAGS_PROT_READWRITE; access_descs.push_back(access); } @@ -575,7 +636,7 @@ struct ggml_cuda_pool_vmm : public ggml_cuda_pool { // set access for non P2P CUmemAccessDesc access = {}; access.location.type = CU_MEM_LOCATION_TYPE_DEVICE; - access.location.id = device; + access.location.id = physical_device; access.flags = CU_MEM_ACCESS_FLAGS_PROT_READWRITE; CU_CHECK(cuMemSetAccess(start_ptr, reserve_size, &access, 1)); } @@ -751,13 +812,17 @@ static bool ggml_backend_cuda_buffer_cpy_tensor(ggml_backend_buffer_t buffer, co if (ggml_backend_buffer_is_cuda(src->buffer)) { ggml_backend_cuda_buffer_context * src_ctx = (ggml_backend_cuda_buffer_context *)src->buffer->context; ggml_backend_cuda_buffer_context * dst_ctx = (ggml_backend_cuda_buffer_context *)dst->buffer->context; - if (src_ctx->device == dst_ctx->device) { + // compare the backing physical devices: distinct virtual devices may share one physical GPU, + // in which case a same-device copy (not a peer copy) is required + const int src_physical = ggml_cuda_get_physical_device(src_ctx->device); + const int dst_physical = ggml_cuda_get_physical_device(dst_ctx->device); + if (src_physical == dst_physical) { CUDA_CHECK(cudaMemcpyAsync(dst->data, src->data, ggml_nbytes(src), cudaMemcpyDeviceToDevice, cudaStreamPerThread)); } else { #ifdef GGML_CUDA_NO_PEER_COPY return false; #else - CUDA_CHECK(cudaMemcpyPeerAsync(dst->data, dst_ctx->device, src->data, src_ctx->device, ggml_nbytes(src), cudaStreamPerThread)); + CUDA_CHECK(cudaMemcpyPeerAsync(dst->data, dst_physical, src->data, src_physical, ggml_nbytes(src), cudaStreamPerThread)); #endif } CUDA_CHECK(cudaStreamSynchronize(cudaStreamPerThread)); @@ -2355,13 +2420,17 @@ static bool ggml_backend_cuda_cpy_tensor_async(ggml_backend_t backend_src, ggml_ if (backend_src != backend_dst) { // copy on src stream - if (cuda_ctx_src->device == cuda_ctx_dst->device) { + // compare the backing physical devices: distinct virtual devices may share one physical GPU, + // in which case a same-device copy (not a peer copy) is required + const int src_physical = ggml_cuda_get_physical_device(cuda_ctx_src->device); + const int dst_physical = ggml_cuda_get_physical_device(cuda_ctx_dst->device); + if (src_physical == dst_physical) { CUDA_CHECK(cudaMemcpyAsync(dst->data, src->data, ggml_nbytes(dst), cudaMemcpyDeviceToDevice, cuda_ctx_src->stream())); } else { #ifdef GGML_CUDA_NO_PEER_COPY return false; #else - CUDA_CHECK(cudaMemcpyPeerAsync(dst->data, cuda_ctx_dst->device, src->data, cuda_ctx_src->device, ggml_nbytes(dst), cuda_ctx_src->stream())); + CUDA_CHECK(cudaMemcpyPeerAsync(dst->data, dst_physical, src->data, src_physical, ggml_nbytes(dst), cuda_ctx_src->stream())); #endif // GGML_CUDA_NO_PEER_COPY } @@ -4345,14 +4414,25 @@ int ggml_backend_cuda_get_device_count() { void ggml_backend_cuda_get_device_description(int device, char * description, size_t description_size) { cudaDeviceProp prop; - CUDA_CHECK(cudaGetDeviceProperties(&prop, device)); + CUDA_CHECK(cudaGetDeviceProperties(&prop, ggml_cuda_get_physical_device(device))); snprintf(description, description_size, "%s", prop.name); } +static int ggml_cuda_physical_device_share_count(int device) { + const ggml_cuda_device_info & info = ggml_cuda_info(); + GGML_ASSERT(device >= 0 && device < info.device_count); + return info.devices[device].physical_share_count; +} + void ggml_backend_cuda_get_device_memory(int device, size_t * free, size_t * total) { ggml_cuda_set_device(device); CUDA_CHECK(cudaMemGetInfo(free, total)); + + // virtual devices sharing one physical GPU share its memory pool; split it between them + const int share_count = ggml_cuda_physical_device_share_count(device); + *free /= share_count; + *total /= share_count; } bool ggml_backend_cuda_register_host_buffer(void * buffer, size_t size) { @@ -4496,7 +4576,7 @@ static void ggml_backend_cuda_device_get_memory(ggml_backend_dev_t dev, size_t * #if defined(__linux__) // Check if this is a UMA (Unified Memory Architecture) system cudaDeviceProp prop; - CUDA_CHECK(cudaGetDeviceProperties(&prop, ctx->device)); + CUDA_CHECK(cudaGetDeviceProperties(&prop, ggml_cuda_get_physical_device(ctx->device))); // Check if UMA is explicitly enabled via environment variable bool uma_env = getenv("GGML_CUDA_ENABLE_UNIFIED_MEMORY") != nullptr; @@ -4515,13 +4595,17 @@ static void ggml_backend_cuda_device_get_memory(ggml_backend_dev_t dev, size_t * } #endif // defined(__linux__) + // virtual devices sharing one physical GPU share its memory pool; split it between them + const int share_count = ggml_cuda_physical_device_share_count(ctx->device); + *free /= share_count; + *total /= share_count; } static enum ggml_backend_dev_type ggml_backend_cuda_device_get_type(ggml_backend_dev_t dev) { ggml_backend_cuda_device_context * ctx = (ggml_backend_cuda_device_context *) dev->context; cudaDeviceProp prop; - CUDA_CHECK(cudaGetDeviceProperties(&prop, ctx->device)); + CUDA_CHECK(cudaGetDeviceProperties(&prop, ggml_cuda_get_physical_device(ctx->device))); return prop.integrated ? GGML_BACKEND_DEVICE_TYPE_IGPU @@ -5167,18 +5251,27 @@ ggml_backend_reg_t ggml_backend_cuda_reg() { ggml_backend_cuda_reg_context * ctx = new ggml_backend_cuda_reg_context; const int min_batch_size = getenv("GGML_OP_OFFLOAD_MIN_BATCH") ? atoi(getenv("GGML_OP_OFFLOAD_MIN_BATCH")) : 32; - for (int i = 0; i < ggml_cuda_info().device_count; i++) { + const ggml_cuda_device_info & info = ggml_cuda_info(); + const bool virtual_devices = info.device_count > info.physical_device_count; + + for (int i = 0; i < info.device_count; i++) { + const int physical_id = info.devices[i].physical_device; + ggml_backend_cuda_device_context * dev_ctx = new ggml_backend_cuda_device_context; dev_ctx->device = i; dev_ctx->name = GGML_CUDA_NAME + std::to_string(i); cudaDeviceProp prop; - CUDA_CHECK(cudaGetDeviceProperties(&prop, i)); + CUDA_CHECK(cudaGetDeviceProperties(&prop, physical_id)); dev_ctx->description = prop.name; char pci_bus_id[32] = {}; - CUDA_CHECK(cudaDeviceGetPCIBusId(pci_bus_id, sizeof(pci_bus_id), i)); + CUDA_CHECK(cudaDeviceGetPCIBusId(pci_bus_id, sizeof(pci_bus_id), physical_id)); dev_ctx->pci_bus_id = pci_bus_id; + if (virtual_devices) { + // make the pci bus id unique for virtual devices + dev_ctx->pci_bus_id += "-v" + std::to_string(i); + } for (char & c : dev_ctx->pci_bus_id) { c = std::tolower(c); }