--numa mirror: mirror model weights to every Numa node in the system#16000
--numa mirror: mirror model weights to every Numa node in the system#16000dbsanfte wants to merge 25 commits into
--numa mirror: mirror model weights to every Numa node in the system#16000Conversation
- Achieved 5% inference speed improvement (14.6 -> 15.3 t/s) - Clean explicit NUMA setup during model loading - Ultra-minimal hot path with thread-local NUMA node access - Working NUMA mirrors for all model weights - Performance: text generation improved, prompt processing needs optimization Performance Results (Qwen3-30B-A3B): - Text Generation: 14.6 -> 15.3 t/s (+5% improvement) - Prompt Processing: 176 -> 152 t/s (14% regression - needs investigation) Technical Implementation: - tensor_data(): O(1) NUMA-aware access via thread-local ggml_current_numa_node - tensor_set_data_with_numa_mirrors(): Explicit NUMA setup for model weights - NUMA coordinator: Thread binding and memory locality - Clean separation: model loading (explicit setup) vs inference (fast access)
dbsanfte
requested review from
0cc4m,
JohannesGaessler and
taronaeo
as code owners
github-actions
Bot
added
testing
|
Physical core detection was very broken in |
|
@Ph0rk0z Actually there is something extra you have to do in but this was only done for the original hack as I needed to save the tensor names in the cache, and we should later be able to patch a different function (which I forget the name of now) that |
|
For that I think you can -no-ooae |
You have to run |
|
Everything I have is on one node and those are the only "real" cores on it. |
|
Seems like there is probably a fair potential for tuning: |
|
Not sure it's doing much yet... proc is a QQ89 so like 82xx stepfun-ai_Step-3.5-Flash-Q4_K_L layer with the patch, 0-23 main: n_kv_max = 65536, n_batch = 2048, n_ubatch = 2048, flash_attn = 1, n_gpu_layers = 46, n_threads = 48, n_threads_batch = 48
Without Patch
Graph no patch
Graph patch 0-23
Graph patch 0-11
What should be tuned? |
First make sure it's actually doing something - during (offloaded) PP: are you seeing a single core in |
|
it might also be worth adding some |
|
I am seeing the system load at half every time. I have a CPU use and temperature on the task bar. Also watching pcm-memory throughput which seems around 54-58 gb/s. It definitely compiles because I see it: |
The problem this solves might not be a problem for you then? It's worth calculating your PCI-E bandwidth to get an idea of the theoretical maximum, then watch In my case: I get around |
|
I have 4x gpu, on this model the problem is the system feeding them because I only see 33-35% load with stepfun, even with graph. It's difficult to get a peak on nvtop due to it only really climbing to the gigs during prompt processing. I should do a trace to see where the bottleneck is. On fully offloaded models it was definitely PCIE. |
|
@jukofyork I setup 8 threads on the cores connected to GPU. During PP in htop I observe those 8 threads on the assigned cpunums at e.g. 50%, plus one 100%, but the 100% core is on the other (i.e. wrong) CPU. 😖 |
Yeah, it seems the best way to avoid this is to leave the first 2 cores free on the CPU connected to the GPU: // ==================================================================
// CUDA pinned-buffer patch for ik_llama.cpp
// Replaces the CPU→GPU memcpy for "_exps" tensors with:
// 1. Parallel CPU gather (pinned NUMA-local threads)
// 2. Double-buffered async DMA (overlap CPU work with GPU transfer)
// ==================================================================
#include <pthread.h>
#include <thread>
#include <algorithm>
#include <mutex>
#include <vector>
#include <sstream>
#include <string>
#include <cstdlib>
#include <cassert>
#include <queue>
#include <unordered_map>
#include <condition_variable>
#include <functional>
#include <memory>
// ------------------------------------------------------------------
// Tunables
// ------------------------------------------------------------------
static constexpr size_t BUFFER_BYTES = 512ULL * 1024 * 1024; // 0.5GB staging buffer
static constexpr size_t NUM_BUFFERS = 2; // ping-pong double buffer
// ---- CPU thread settings (edit these) -----------------------------
static constexpr int CPU_THREAD_START = 2; // skip cores 0-1 (OS / main)
static constexpr int CPU_THREAD_COUNT = 8; // exactly this many worker threads
// ------------------------------------------------------------------
// ------------------------------------------------------------------
// 1. CPU core list (hardcoded — no env vars, no auto-detect)
// ------------------------------------------------------------------
static const std::vector<int> & get_gpu_local_cores() {
static const std::vector<int> cores = []{
std::vector<int> v;
v.reserve(CPU_THREAD_COUNT);
for (int i = CPU_THREAD_START; i < CPU_THREAD_START + CPU_THREAD_COUNT; ++i) {
v.push_back(i);
}
return v;
}();
return cores;
}
// ------------------------------------------------------------------
// 2. Persistent thread pool (per process, not per call)
// ------------------------------------------------------------------
class LocalThreadPool {
public:
explicit LocalThreadPool(const std::vector<int> & cores) {
const int n = static_cast<int>(cores.size());
for (int i = 0; i < n; ++i) {
workers_.emplace_back([this, core = cores[i]] {
cpu_set_t mask;
CPU_ZERO(&mask);
if (core >= 0 && core < CPU_SETSIZE) CPU_SET(core, &mask);
pthread_setaffinity_np(pthread_self(), sizeof(mask), &mask);
while (true) {
std::function<void()> task;
{
std::unique_lock<std::mutex> lock(mtx_);
cv_.wait(lock, [this]{ return stop_ || !tasks_.empty(); });
if (stop_ && tasks_.empty()) return;
task = std::move(tasks_.front());
tasks_.pop();
}
task();
}
});
}
}
void enqueue(std::function<void()> task) {
{
std::unique_lock<std::mutex> lock(mtx_);
tasks_.push(std::move(task));
}
cv_.notify_one();
}
void wait() {
{
std::unique_lock<std::mutex> lock(mtx_);
cv_.wait(lock, [this]{ return tasks_.empty(); });
}
}
~LocalThreadPool() {
{
std::unique_lock<std::mutex> lock(mtx_);
stop_ = true;
}
cv_.notify_all();
for (auto & t : workers_) if (t.joinable()) t.join();
}
private:
std::vector<std::thread> workers_;
std::queue<std::function<void()>> tasks_;
std::mutex mtx_;
std::condition_variable cv_;
bool stop_ = false;
};
// singleton: created on first chunked memcpy, reused forever
static LocalThreadPool & get_pool() {
static LocalThreadPool pool(get_gpu_local_cores());
return pool;
}
// ------------------------------------------------------------------
// 3. Parallel CPU memcpy
// ------------------------------------------------------------------
static void parallel_memcpy_node_local(void * dst, const void * src, size_t n) {
const auto & cores = get_gpu_local_cores();
const int num_threads = static_cast<int>(cores.size());
if (n < 4ULL * 1024 * 1024 || num_threads <= 1) {
memcpy(dst, src, n);
return;
}
auto & pool = get_pool();
std::mutex done_mtx;
int done = 0;
std::condition_variable done_cv;
for (int t = 0; t < num_threads; ++t) {
size_t off = (n * t) / num_threads;
size_t len = (n * (t + 1)) / num_threads - off;
pool.enqueue([=, &done, &done_mtx, &done_cv]{
memcpy((char *)dst + off, (const char *)src + off, len);
{
std::unique_lock<std::mutex> lock(done_mtx);
++done;
}
done_cv.notify_one();
});
}
{
std::unique_lock<std::mutex> lock(done_mtx);
done_cv.wait(lock, [&]{ return done == num_threads; });
}
}
// ------------------------------------------------------------------
// 4. GPU staging scratch area
// ------------------------------------------------------------------
struct PinnedScratch {
void * scratch = nullptr;
cudaEvent_t evt[NUM_BUFFERS] = {};
std::mutex xfer_mtx;
PinnedScratch() {
CUDA_CHECK(cudaHostAlloc(&scratch, NUM_BUFFERS * BUFFER_BYTES, cudaHostAllocDefault));
for (size_t i = 0; i < NUM_BUFFERS; ++i) {
CUDA_CHECK(cudaEventCreate(&evt[i]));
}
}
};
static PinnedScratch & get_scratch(int cuda_device) {
static std::unordered_map<int, std::unique_ptr<PinnedScratch>> table;
static std::mutex mtx;
std::lock_guard<std::mutex> lock(mtx);
auto & ps = table[cuda_device];
if (!ps) {
cpu_set_t cpuset;
CPU_ZERO(&cpuset);
const auto & cores = get_gpu_local_cores();
CPU_SET(cores.empty() ? 0 : cores[0], &cpuset);
pthread_setaffinity_np(pthread_self(), sizeof(cpuset), &cpuset);
ps = std::make_unique<PinnedScratch>();
}
return *ps;
}
// ------------------------------------------------------------------
// 5. Hook: ggml CUDA tensor upload
// ------------------------------------------------------------------
GGML_CALL static void ggml_backend_cuda_buffer_set_tensor(
ggml_backend_buffer_t buffer,
ggml_tensor * tensor,
const void * data,
size_t offset,
size_t size)
{
ggml_backend_cuda_buffer_context * ctx =
(ggml_backend_cuda_buffer_context *)buffer->context;
ggml_cuda_set_device(ctx->device);
if (strstr(tensor->name, "_exps") == nullptr) {
CUDA_CHECK(cudaMemcpyAsync((char *)tensor->data + offset,
data,
size,
cudaMemcpyHostToDevice,
cudaStreamPerThread));
CUDA_CHECK(cudaStreamSynchronize(cudaStreamPerThread));
return;
}
auto & ps = get_scratch(ctx->device);
const size_t CHUNK = BUFFER_BYTES;
std::lock_guard<std::mutex> lock(ps.xfer_mtx);
const size_t num_chunks = (size + CHUNK - 1) / CHUNK;
for (size_t i = 0; i < num_chunks; ++i) {
size_t pos = i * CHUNK;
size_t chunk = std::min(CHUNK, size - pos);
size_t cur = i % NUM_BUFFERS;
if (i >= NUM_BUFFERS) {
CUDA_CHECK(cudaEventSynchronize(ps.evt[cur]));
}
char * cur_buf = (char *)ps.scratch + cur * CHUNK;
const char * src = (const char *)data + pos;
parallel_memcpy_node_local(cur_buf, src, chunk);
if (i > 0) {
size_t prev_i = i - 1;
size_t prev_pos = prev_i * CHUNK;
size_t prev_chunk = std::min(CHUNK, size - prev_pos);
size_t prev = prev_i % NUM_BUFFERS;
char * prev_buf = (char *)ps.scratch + prev * CHUNK;
char * prev_dst = (char *)tensor->data + offset + prev_pos;
CUDA_CHECK(cudaMemcpyAsync(prev_dst, prev_buf, prev_chunk,
cudaMemcpyHostToDevice,
cudaStreamPerThread));
CUDA_CHECK(cudaEventRecord(ps.evt[prev], cudaStreamPerThread));
}
}
if (num_chunks > 0) {
size_t last_i = num_chunks - 1;
size_t last_pos = last_i * CHUNK;
size_t last_chunk = std::min(CHUNK, size - last_pos);
size_t last_cur = last_i % NUM_BUFFERS;
char * last_buf = (char *)ps.scratch + last_cur * CHUNK;
char * last_dst = (char *)tensor->data + offset + last_pos;
CUDA_CHECK(cudaMemcpyAsync(last_dst, last_buf, last_chunk,
cudaMemcpyHostToDevice,
cudaStreamPerThread));
}
CUDA_CHECK(cudaStreamSynchronize(cudaStreamPerThread));
}As before, find and replace fit properly with your specific NUMA architecture eg: (use the Using a fixed set of worker threads seems to have helped my There are some other CUDA functions that might help (eg: cudaMemPrefetchAsync), but I doubt if I'll get time to try anything in the near future - at least now we don't need 2x the RAM! 😁 |
|
Also I wouldn't bother trying to increase the buffer (or number of buffers) - it will just add extra latency and from what I've read there might be some finite limit to the amount of memory you are allowed to allocate with |
|
There's some slight uplift on this one: Stock:
New patch 10c
I chose 10 cores vs 8 because I have 24 per CPU. Just for fun I ran it without --no-ooae
I figured it would crash? But it did not. |
IIRC, it uses a different function and never calls this one with the
I wouldn't expect much change though, as at least for |
Also, I would try and play with the buffer size and thread range, as even between my Broadwell E5-2699v4 machines (2x QP link) and my Cascade Lake Xeon Gold 6248 machine (3x UPI link), there seems to quite a difference in what settings work or not. |
|
I tried with qwen and saw a slightly higher 4.0gb/s transfer. So sadly this is still all minor changes. What worked better on the 6248? That's closer to my procs. |
|
Just seeing these new posts lol. I've been in
Doesn't interleave just make the memory pool of all the NUMA nodes behave as a single unified pool? So if data is properly distributed for NUMA, it would be better to disable it? |
Yeah, I don't use it but others have. I find I can get the best TG by flushing the buffers and using |
Sorry missed this: I'm just using the same settings as I posted above (8 threads and 2x 512MB buffers). |
|
mmap leaves me with lower speeds since the models load off HDD. I suppose after a bunch of use it should even out. |
|
I use --numa distribute w/drop_caches protocol and mmap. AFAIK I have tried every combination and this gives the highest decode throughput w/ik_llama and llama.cpp. (Kimi Q4, 2x EPYC 9b14 NPS4, 1x R6KPMQ) I tried Juk's follow-up above re: avoiding the low position cores, and it does indeed seem to resolve the main-thread-on-wrong-cpu problem. However, I tried various different arrangements of the thread pool, incl.
I expected to observe some arrangement on CPU0 (GPU attached to CPU0) to give a higher throughput, and any arrangement on CPU1 to give a significantly lower throughput. Instead I observe very little difference between any. I'm not sure what to make of this. Slightly off-topic (but maybe interesting for NUMA users):The fastest arrangement I found wrt. thread count, placement and buffer size is still not as fast as sglang/kt-kernel. sg/kt's layerwise-offload transfers the weights in 0.330s / layer, or ~20s. For Kimi-Q4 that's 531GiB of _exp tensors, 531 / 19.8s = 26.81GiB/s. Compute for a fully-loaded 32768 batch is ~0.120s/layer, so 32768 / ((0.330+0.120)*60(layers)) = ~1200t/s prefill. w/o layerwise prefill it's 50-100t/s depending on context. Setting the threshold accordingly gives pretty good latency for e.g. tool calls. At low context ik_llama's decode throughput is better by 10% or so, but that is lost once 100K is reached and sglang pulls ahead significantly. The build/install quality of sg/kt leaves much to be desired, but once a working commit/library-build out/CUDA combo etc. is found the program itself is stable and runs w/o error for days or weeks. ...and even more off-topic (but still maybe interesting for NUMA users):Meanwhile lvllm and lsglang appear to have matured a lot wrt. build/install quality. lvllm w/DSv4-Flash gives ~45-50t/s decode w/MTP on my setup, prefill ~1000t/s (by wall-clock measurement, the logs report values that are nonsensically high). There's high latency as the system is first running since tilelang continues to compile additional kernel variations for quite some time. Eventually this settles and the system behaves normally wrt. TTFT. lsglang w/Kimi gives me incoherent/gibberish output, reason unknown. I integrated kt-kernel into a DwarfStar fork here using the MXFP4 weights. This gives ~350t/s prefill and ~34t/s decode, and frees up GPU space for e.g. 500K context in ~30GB VRAM, or 1M in ~50GB. lvllm is significantly faster than this implementation, but DS4 is still interesting for its integration wrt. cache/chat/tool handling. (It also runs -Pro w/~100t/s prefill and ~10t/s decode) |
The last time I tried this, it improved decode speed but murdered my prefill speed (compared to --numa distribute + interleave + no mmap, or just keeping everything on a single socket if it fits). I haven't yet tested juk's new workaround for that.
On my setup, llama.cpp actually has faster offload-to-GPU prefill speed than sglang/kt-kernel, at least for normal-ish batch sizes that I encounter in practice. My sglang/kt-kernel layerwise-offload speed is a bit slower than yours, like 0.380s / layer, but in llama.cpp, I can routinely near-saturate my PCIe 5.0 x16 bandwidth, hitting like 50GiB/s. I'm not sure if it breaks even at some point when I hit batch sizes of like 32768 in sglang/kt-kernel compared to the 4096 I have set in llama.cpp for reasonable compute buffer sizes. sglang/kt-kernel has excellent decode speed with true multi-NUMA tensor parallelism though. Also AMX support for really speedy prefill on models quanted to their bespoke format, but I kinda don't trust the quality of per-channel RTN int4 quant, and 8bit is rather fat and slows decode a lot.
Yeah a bunch of stuff broke with the updates for DSv4 lol, I think I finally have ol' Kimi K2.6 working again on latest commit sglang/kt-kernel.
Interesting, might be another backend for me to look into. |
|
fastllm had the best decode for me and I saw all my bandwidth. Not being an epyc user, I am in the shit position of putting everything on one CPU not helping. QQ89 with 24c is simply not enough. I also get some bottleneck from PCIE3 and one PLX per 2 GPUs. Need to try lvllm and lsglang with some mid size model like qwen-235b. QOL on these backends is terrible compared to IK so I take the path of least resistance. |
|
Sadly I don't really know enough about CUDA or the workings of GGML, but I've had another idea for this that might work but at the cost of quite a lot of extra VRAM:
I doubt I will be able to try this as can't afford an extra 9GB of VRAM, but also because it is probably very brittle and the optimisation of skipping the final MoE tensors for all but the final token in a PP batch will add complexity, etc. It may turn out that the extra cost in VRAM can be somewhat offset by using much smaller |
|
I've released my inferencing engine called Llaminar. It solves the NUMA problem by using OpenMPI and should scale to large clusters over Infiniband although I lack the hardware to test it. Also supports CUDA and ROCm, with the fastest Mi50 kernels that I've seen. https://github.com/Llaminar/llaminar Feel free to have a play. |
|
I have avx512 but no vnni so stuck in /wait |
After being reminded of the loss I revisited this. Here is an implementation of the double-buffer with prefetch from NUMA shards that gives nearly the same perf as NPS0 pinned tensors. I also tried an impl moving the buffers to VRAM, but the perf wasn't as great at first glance and quite expensive when that space can now be used as KV for 1M ctx models that weren't available earlier in this thread. Example invocation: MEMORY_CONCURRENCY_IN_NUMA_NODES=8
COMPUTE_POWER_IN_THREADS_PER_NUMA=6
M=/model/GLM-5.2/unsloth/UD-Q4_K_XL/GLM-5.2-UD-Q4_K_XL-00001-of-00011.gguf
N_THREADS=$[ $MEMORY_CONCURRENCY_IN_NUMA_NODES * $COMPUTE_POWER_IN_THREADS_PER_NUMA ]
N_BATCH=8192
GGML_CUDA_NO_PINNED=1 \
GGML_CUDA_EXPS_READAHEAD_THREADS=10 \
./ik_llama.cpp/build/bin/llama-server \
--host 0.0.0.0 --port 4972 --webui llamacpp \
--numa distribute \
-t $N_THREADS \
--cache-ram 200000 \
-b $N_BATCH -ub $N_BATCH -amb 512 \
-mla 1 --dsa -fidx -ctk f16 \
-wgt 1 \
-ngl 999 -ot exps=CPU \
--jinja --parallel-tool-calls \
--chat-template-kwargs '{"enable_thinking": true, "reasoning_effort": "max"}' \
-c $[ 2**18 ] \
--metrics \
--alias GLM-5.2 \
-m "$M"(not shown: drop_caches, anyone reading this should know this protocol) |
|
Nice! I'm away from home for the next week or so but will give this a try when I get back. |
7 participants
This PR adds a new
--numa mirroroption which mirrors model weights to each Numa node on the system, and uses a thread-local var in the OMP threadpool to select the correct mirror copy local to the thread at runtime, to eliminate cross-socket traffic.Build instructions:
To test:
Test system is a two-socket Xeon 6238R Cascade Lake, with 768GB of DDR4-2933 (6 channels per socket).
Without
--numa mirror:With
--numa mirror:Intel PCM tool during mirror inference showing both sockets using local mem:
There's still a bit of cross-socket traffic (5%) because only model weights are mirrored, not tensors created at inference time. I'll play with that, maybe mirroring those aggressively will help too, or maybe not. Right now anything created at inference time just gets set to live on Node 0.