GGML_ABORT(GGML_CUDA_NAME " error");
}
+// map a (possibly virtual) device id to the physical CUDA device that backs it
+static 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;
+}
+
// 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() {
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);
- int64_t total_vram = 0;
+ // 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), assign each its index
+ // among the (virtual) devices sharing that physical GPU, and store the per-physical share count
+ 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;
+ info.devices[id].virtual_index = physical_share_count[info.devices[id].physical_device]++;
+ }
+
+ int64_t total_vram = 0;
+ 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;
#if defined(GGML_USE_HIP)
info.devices[id].integrated = prop.integrated;
#else
#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;
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;
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 });
// 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));
}
// 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;
}
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;
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) {
}
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));
// 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);
}
// 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));
}
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));
static void ggml_backend_cuda_comm_init_nccl(ggml_backend_cuda_comm_context * ret) {
#ifdef GGML_USE_NCCL
+ // Disabling NCCL path when CUDA virtual devices are in use since NCCL requires one distinct physical GPU per rank.
+ const ggml_cuda_device_info & info = ggml_cuda_info();
+ if (info.device_count > info.physical_device_count) {
+ GGML_LOG_WARN("NCCL disabled: virtual devices in use; "
+ "falling back to internal AllReduce\n");
+ ggml_backend_cuda_comm_init_internal(ret);
+ return;
+ }
+
const size_t n = ret->dev_ids.size();
ret->comms.resize(n);
ncclResult_t rc = ncclCommInitAll(ret->comms.data(), (int) n, ret->dev_ids.data());
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
}
return ggml_cuda_info().device_count;
}
-void ggml_backend_cuda_get_device_description(int device, char * description, size_t description_size) {
+static std::string ggml_cuda_device_description(int device) {
cudaDeviceProp prop;
- CUDA_CHECK(cudaGetDeviceProperties(&prop, device));
- snprintf(description, description_size, "%s", prop.name);
+ CUDA_CHECK(cudaGetDeviceProperties(&prop, ggml_cuda_get_physical_device(device)));
+
+ const ggml_cuda_device_info & info = ggml_cuda_info();
+ std::string description = prop.name;
+ if (info.device_count > info.physical_device_count) {
+ description += " (physical device " + std::to_string(info.devices[device].physical_device) +
+ ", virtual device " + std::to_string(info.devices[device].virtual_index) + ")";
+ }
+ return description;
+}
+
+void ggml_backend_cuda_get_device_description(int device, char * description, size_t description_size) {
+ snprintf(description, description_size, "%s", ggml_cuda_device_description(device).c_str());
+}
+
+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) {
#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;
}
#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
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));
- dev_ctx->description = prop.name;
+ dev_ctx->description = ggml_cuda_device_description(i);
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);
}
overview / pkg / ggml / sources / llama.cpp / commitdiff