upstream-archive/llama.cpp
1
0
Fork 0
mirror of https://github.com/ggml-org/llama.cpp.git synced 2026-08-06 01:50:50 -05:00
Code Issues Packages Projects Releases Wiki Activity

support cuda virtual devices

Browse Source
This commit is contained in:
Anav Prasad
2026-06-25 10:26:56 +00:00
parent c264f65ff9
commit a8eca04c3f
2 changed files with 134 additions and 35 deletions
Download Patch File Download Diff File Expand all files Collapse all files

8
ggml/src/ggml-cuda/common.cuh
View File

@@ -1115,7 +1115,8 @@ struct ggml_cuda_type_traits<GGML_TYPE_IQ3_S> {
//////////////////////
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;

161
ggml/src/ggml-cuda/ggml-cuda.cu
View File

@@ -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(&current_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<std::pair<int, std::string>> 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<CUmemAccessDesc> 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);
}