Getting Started#

Introduction#

cuBLASMp aims to provide GPU-accelerated PBLAS-like basic linear algebra functionality.
cuBLASMp leverages the 2D block-cyclic data layout for load balancing and to maximize compatibility with PBLAS routines.
The library assumes the data is available in device memory. It is the responsibility of the developer to allocate memory and to copy data between GPU memory and CPU memory using standard CUDA runtime APIs, such as cudaMalloc(), cudaMallocAsync(), and cudaMemcpyAsync(). For cublasMpMatmul, allocating the device workspace with NCCL symmetric memory, or registering a compatible allocation with cublasMpBufferRegister, enables the high-performance AllGather + Matmul, Matmul + ReduceScatter, and Matmul + AllReduce algorithms; without it, the library uses no_overlap (see Memory Management).
Python users can access cuBLASMp through nvmath-python, either via the higher-level Distributed APIs or the low-level nvmath.bindings.cublasMp module.

Hardware and Software Requirements#

GPU Architectures

Compute Capability 7.5 (Turing) and above (CUDA 13)

Compute Capability 7.0 (Volta) and above (CUDA 12)

CUDA

13.x

12.9.0 or above recommended, 12.x compatible

CPU Architectures

x86_64, arm64-sbsa

Operating Systems

Linux

  • Recommended NVIDIA InfiniBand solutions for accelerated inter-node communication

Note

Single-Process Multi-GPU (SPMG) execution is supported with CUDA 13+.

Required Packages

Note

cuBLASMp versions older than 0.8.0 also require NVSHMEM v3.3.24 and above. Starting from cuBLASMp 0.8.0, NVSHMEM is no longer required as the library uses NCCL Symmetric Memory instead.

Recommended Packages

Data Layout of Local Matrices#

cuBLASMp assumes that local matrices are stored in column-major format.

Workflow#

cuBLASMp’s workflow can be broken down as follows:
1. Create a NCCL communicator: NCCL Initialization.
2. Initialize the library handle: cublasMpCreate().
3. Initialize grid descriptors: cublasMpGridCreate(). The NCCL communicator passed to this call must contain exactly nprow * npcol ranks.
4. Initialize matrix descriptors: cublasMpMatrixDescriptorCreate().
5. Query the host and device buffer sizes for a given routine.
6. Allocate workspace buffers. For cublasMpMatmul, use cublasMpMalloc() or allocate with ncclMemAlloc / CUDA VMM and register the allocation with cublasMpBufferRegister() to enable the communication-computation overlap in the AllGather + Matmul, Matmul + ReduceScatter, and Matmul + AllReduce algorithms. Standard cudaMalloc is also accepted, but then the library either uses no_overlap or returns CUBLASMP_STATUS_NOT_SUPPORTED for explicit pipelined algorithm requests. For all other routines, use standard cudaMalloc or any other CUDA memory allocator. The host workspace and input/output matrix buffers (A, B, C, D) always use standard allocators.
7. Execute the routine to perform the desired computation.
8. Synchronize local stream to make sure the result is available, if required: cudaStreamSynchronize().
9. Deallocate host and device workspace.
10. Destroy matrix descriptors: cublasMpMatrixDescriptorDestroy().
11. Destroy grid descriptors: cublasMpGridDestroy().
12. Destroy cuBLASMp library handle: cublasMpDestroy().
13. Destroy the NCCL communicator: NCCL Initialization.

CUDA Graphs#

cuBLASMp supports CUDA Graphs for all routines except matrix redistribution operations (GEMR2D and TRMR2D).

Note

AllGather + Matmul and Matmul + ReduceScatter require NCCL 2.30.4 or above for CUDA Graphs support. NCCL 2.30.7 or above is preferred; see the relevant Troubleshooting known issue.

Single-Process Multi-GPU (SPMG)#

Single-Process Multi-GPU (SPMG) execution lets one process drive multiple cuBLASMp ranks by using one host thread per local GPU. The process grid still counts NCCL/cuBLASMp ranks, not operating-system processes. For example, two MPI processes with four SPMG threads each create up to eight cuBLASMp ranks. An operation can use all eight ranks, or it can use a subset of those ranks, such as four ranks with a 1 x 4, 2 x 2, or 4 x 1 process grid.

In SPMG, every participating host thread is expected to participate in the NCCL communicator and in each collective cuBLASMp API call for its rank. A typical SPMG worker thread should:

  • Select its CUDA device with cudaSetDevice() before creating CUDA, NCCL, or cuBLASMp objects.

  • Initialize its own ncclComm_t with that thread’s rank in the participating NCCL communicator.

  • Create its own CUDA stream and cublasMpHandle_t.

  • Create its own cublasMpGrid_t objects from that thread’s NCCL communicator handle.

  • Create its own matrix descriptors, matmul descriptors, device workspace, and local matrix buffers for that rank.

  • Enter cuBLASMp calls in the same collective order as the other participating threads.

Do not share one cuBLASMp handle, grid, matrix descriptor, matmul descriptor, stream, or workspace across SPMG threads. Treat these objects as rank-local: create, use, synchronize, and destroy them from the thread that owns the corresponding CUDA device and NCCL rank. If one participating thread skips a collective cuBLASMp call while the others enter it, the operation follows the same failure mode as a missing NCCL rank and can hang.

Code Samples#

Code samples can be found in the CUDALibrarySamples repository.