SALT Systems Lab

Projects / GPU Data Management

warpSpeed

Concurrent GPU hash tables

CUDA

Overview

warpSpeed is a library of concurrent, thread-safe GPU hash tables, together with the building blocks for writing new ones. Every table has a similar API, and the repository includes a benchmark suite that compares the tables against existing state-of-the-art designs.

Tables are initialized with a fixed number of keys, support key-value pairs up to 16 bytes, and are tiled for performance.

Key features

Eight table designs

Chaining, cuckoo, double hashing, iceberg hashing, and power-of-two-choice hashing, several with metadata variants.

Uniform API

upsert_replace, upsert_function, find_with_reference, find_pair, and remove.

Lock-free compound ops

Every operation has a _no_lock variant for building compound operations inside a critical section.

Thorough benchmarks

Load, scaling, aging, caching, sparse-tensor, and adversarial-race benchmarks.

Core device API

__device__ bool upsert_replace(const tile_type& tile, const Key& key, const Val& val);
__device__ bool find_with_reference(tile_type tile, Key key, Val& val);
__device__ bool remove(tile_type tile, Key key);

// Lock-free variants for compound operations:
uint64_t b = get_lock_bucket(tile, key);
stall_lock(tile, b);
upsert_replace_no_lock(tile, key, val);
unlock(tile, b);