diff --git a/lib/kernels/src/kernels/optimizer_kernels.cc b/lib/kernels/src/kernels/optimizer_kernels.cc index 3d92d9ddc9..00b12a24bf 100644 --- a/lib/kernels/src/kernels/optimizer_kernels.cc +++ b/lib/kernels/src/kernels/optimizer_kernels.cc @@ -23,18 +23,41 @@ void sgd_update_task(device_stream_t const &stream, sgd_v_ptr = sgd_v.value().get_float_ptr(); } - gpu_sgd_nccl_update_task( - /*stream=*/stream.require_gpu(), - /*lr=*/lr, - /*momentum=*/momentum, - /*nesterov=*/nesterov, - /*weight_decay=*/weight_decay, - /*handle=*/handle.require_for_gpu(), - /*weight_grad_ptr=*/weight_grad.get_float_ptr(), - /*size=*/ - get_num_elements(weight_grad.shape.dims).int_from_positive_int(), - /*weight_ptr=*/weight.get_float_ptr(), - /*sgd_v_ptr=*/sgd_v_ptr); + PerDeviceFFHandle const &gpu_handle = handle.require_for_gpu(); +#ifdef FF_USE_NCCL + bool have_nccl_comm = (gpu_handle.ncclComm != nullptr); +#else + bool have_nccl_comm = false; +#endif + + if (have_nccl_comm) { + gpu_sgd_nccl_update_task( + /*stream=*/stream.require_gpu(), + /*lr=*/lr, + /*momentum=*/momentum, + /*nesterov=*/nesterov, + /*weight_decay=*/weight_decay, + /*handle=*/gpu_handle, + /*weight_grad_ptr=*/weight_grad.get_float_ptr(), + /*size=*/ + get_num_elements(weight_grad.shape.dims).int_from_positive_int(), + /*weight_ptr=*/weight.get_float_ptr(), + /*sgd_v_ptr=*/sgd_v_ptr); + } else { + // No communicator (single-rank), so the gradient allreduce is a no-op. + gpu_sgd_ps_update_task( + /*stream=*/stream.require_gpu(), + /*lr=*/lr, + /*momentum=*/momentum, + /*nesterov=*/nesterov, + /*weight_decay=*/weight_decay, + /*weight_grad_ptr=*/weight_grad.get_float_ptr(), + /*size=*/ + get_num_elements(weight_grad.shape.dims).int_from_positive_int(), + /*num_replicas=*/num_replicas, + /*weight_ptr=*/weight.get_float_ptr(), + /*sgd_v_ptr=*/sgd_v_ptr); + } } else { ASSERT(stream.is_cpu()); ASSERT(handle.is_for_cpu()); diff --git a/lib/kernels/src/managed_per_device_ff_handle.cc b/lib/kernels/src/managed_per_device_ff_handle.cc index cccc46d6bf..a6fc855403 100644 --- a/lib/kernels/src/managed_per_device_ff_handle.cc +++ b/lib/kernels/src/managed_per_device_ff_handle.cc @@ -18,10 +18,16 @@ ManagedPerDeviceFFHandle::ManagedPerDeviceFFHandle( checkCUDA(cudaMalloc(&this->handle->workSpace, this->handle->workSpaceSize)); #ifdef FF_USE_NCCL - ncclUniqueId ncclId; - checkNCCL(ncclGetUniqueId(&ncclId)); - checkNCCL(ncclCommInitRank( - &handle->ncclComm, num_ranks, ncclId, my_rank)); // todo generalize + // A single-rank communicator can only ever perform no-op collectives, but + // still costs ~500MiB of device memory, so skip creating it entirely. + if (num_ranks == 1) { + this->handle->ncclComm = nullptr; + } else { + ncclUniqueId ncclId; + checkNCCL(ncclGetUniqueId(&ncclId)); + checkNCCL(ncclCommInitRank( + &handle->ncclComm, num_ranks, ncclId, my_rank)); // todo generalize + } #endif } @@ -41,7 +47,9 @@ ManagedPerDeviceFFHandle::~ManagedPerDeviceFFHandle() { checkCUBLAS(cublasDestroy(this->handle->blas)); checkCUDA(cudaFree(this->handle->workSpace)); #ifdef FF_USE_NCCL - checkNCCL(ncclCommDestroy(this->handle->ncclComm)); + if (this->handle->ncclComm != nullptr) { + checkNCCL(ncclCommDestroy(this->handle->ncclComm)); + } #endif delete this->handle; }