Gallatin · SALT Systems Lab

Overview

Gallatin is a general-purpose GPU allocator. GPU threads can malloc and free device-side allocations inside kernels with no host calls and no synchronization with the host. It scales to GPUs of any size and supports allocations of any size, up to all of GPU DRAM.

GPU memory is split into segments. Segments can be formatted into smaller blocks, which are further divided into slices, or several contiguous segments can be combined into one large allocation. Van Emde Boas trees give fast successor search over segments, so Gallatin always picks the lowest-ID segment, which keeps fragmentation low and maximizes memory reuse.

Key features

In-kernel malloc/free

Device-side allocation from any thread, with no host round-trips.

Any size

Allocations range from a few bytes up to all of device memory.

Tunable

Template parameters for segment, minimum, and maximum sizes. Tuning can more than double peak throughput for fixed-size workloads.

Header-only

Depends only on CUDA and Cooperative Groups. Add it with CPM in one snippet.

Add to CMake and allocate in a kernel

# CMakeLists.txt
CPMAddPackage(NAME gallatin GITHUB_REPOSITORY saltsystemslab/gallatin GIT_TAG origin/main)
target_link_libraries(${EXE_NAME} PRIVATE gallatin)

// main.cu
#include <gallatin/allocators/global_allocator.cuh>
using namespace gallatin::allocators;

__global__ void kernel() {
  void* p = global_malloc(256);
  /* ... */
  global_free(p);
}

init_global_allocator(16ULL << 30, /*seed=*/42);   // host
kernel<<<blocks, threads>>>();
free_global_allocator();