ggml: backend-agnostic tensor parallelism (experimental)

[JohannesGaessler]

This PR adds initial support for tensor parallelism, enabled via specifying --split-mode tensor. This should be considered as an experimental feature that is not yet production ready. In principle the implementation is backend-agnostic, in practice as of right now only the CUDA backend has received the necessary extensions and performance optimizations to make the performance better than --split-mode layer (in some cases).

The preexisting --split-mode row could already parallelize some matrix multiplications in the CUDA backend but this required a synchronization after every single operation. As a consequence the overhead is so large that it is only really worthwhile for old and slow GPUs like P40s where adding a bit of latency between operations makes relatively little difference to the overall runtime. The new implementation works by adding a new "meta" backend that internally wraps multiple conventional ggml backends. When given a compute graph the meta backend then automatically infers how the data is split based on the ggml compute graph and only schedules a synchronization at the necessary points. And the external interface for a meta backend is the same as for any other ggml backend. So in practice the meta backend allows ggml to use multiple GPUs in the same way as a single GPU. Importantly all of this is done at the ggml backend level and there are no hard dependencies for any extensions beyond what already exists on master (but without extensions the performance may be so bad that there is no point).

Current Support Status

What currently works:

• Multiple CUDA GPUs work.
• For good performance, make sure that NCCL is installed.
• Performance should become comparatively better for deeper contexts where there is more work for each individual GPU.
• The "ROCm" backend works since it is just the CUDA code translated via HIP. On the hardware combinations that I have (RX 6800 + MI50 or RX 9060 XT + MI100) the performance is bad vs. the -sm layer baseline though.
• Any number of GPUs is supported (limited to an arbitrary maximum of 16 GPUs for practical reasons, ping me if you have a use case that needs a higher limit).
• The preexisting argument --tensor-split can be used to distribute arbitrary fractions of the data across GPUs.
• In my testing --n-gpu-layers was working correctly in combination with -sm tensor though I suspect that there are some scenarios where this combination is broken.
• Most models should work correctly with --split-mode tensor, for more details see below.

What currently doesn't work:

• There is currently an issue with a VRAM leak due to CUDA graphs being continually created but never destroyed (negligible for -sm layer but not for -sm tensor).
• Vulkan technically works at short contexts but the performance is bad, at long contexts there are also stability issues.
• All other backends may work but should be assumed to be broken or unusable by default.
• Going forward the parallelization of NUMA nodes for better CPU performance is planned. As of right now there is no support.
• FlashAttention must be enabled and KV cache quantiztation must be disabled, or otherwise the meta backend will be unable to correctly infer the data splits between GPUs.
• Support for -fit is not implemented, you may need to set the --ctx-size manually.
• Support for backend sampling is not implemented, I'm not yet sure what the best way to go about it is.
• Memory use in general is inefficient due to overallocations.

Model support

For quality control of this PR I added end-to-end tests in test-llama-archs to assert that --split-mode row produces consistent results with the CPU backend. Based on these tests most of the model architectures in llama.cpp are working correctly. However, the following model architectures are definitely broken with -sm tensor:

Details

• bitnet
• deepseek2
• falcon_h1
• gemma3n
• glm_dsa
• granite_hybrid
• grok
• jamba
• kimi_linear lfm2
• lfm2moe
• mamba
• mamba2
• minicpm
• minimax_m2
• mistral4
• mpt
• nemotron_h
• nemotron_h_moe
• olmo2
• olmoe
• plamo2
• t5

Furthermore, for the following model architectures the tests themselves are broken and they probably also will not work:

Details

• arwkv7
• bert
• deepseek2ocr
• eurobert
• embed
• gemma_embedding
• jina_bert_v2
• jina_bert_v3
• modern_bert
• neo_bert
• nomic_bert_moe
• plm
• rwkv6
• rwkv6qwen2
• rwkv7
• t5encoder

Performance

Performance on 2x RTX 4090

model	test	t/s -sm layer	t/s -sm tensor	Speedup
gemma4 26B A4B Q4_0	pp512	8682.29	3428.43	0.39
gemma4 26B A4B Q4_0	pp2048	12917.31	3245.69	0.25
gemma4 26B A4B Q4_0	tg128	197.11	138.76	0.70
gemma4 26B A4B Q4_0	pp512 @ d65536	3892.61	2507.52	0.64
gemma4 26B A4B Q4_0	pp2048 @ d65536	6714.55	2575.19	0.38
gemma4 26B A4B Q4_0	tg128 @ d65536	146.11	118.64	0.81
gemma4 31B Q4_K_M	pp512	2998.02	2846.09	0.95
gemma4 31B Q4_K_M	pp2048	4475.39	2725.37	0.61
gemma4 31B Q4_K_M	tg128	43.81	59.45	1.36
gemma4 31B Q4_K_M	pp512 @ d65536	1308.09	1706.04	1.30
gemma4 31B Q4_K_M	pp2048 @ d65536	2100.01	1760.90	0.84
gemma4 31B Q4_K_M	tg128 @ d65536	34.03	48.80	1.43
gemma4 31B Q8_0	pp512	2736.66	2877.47	1.05
gemma4 31B Q8_0	pp2048	4117.76	2708.27	0.66
gemma4 31B Q8_0	tg128	26.74	40.98	1.53
gemma4 31B Q8_0	pp512 @ d65536	1259.75	1696.77	1.35
gemma4 31B Q8_0	pp2048 @ d65536	2013.56	1751.84	0.87
gemma4 31B Q8_0	tg128 @ d65536	22.75	35.72	1.57
gpt-oss 20B MXFP4 MoE	pp512	9552.27	9298.15	0.97
gpt-oss 20B MXFP4 MoE	pp2048	13987.72	8734.39	0.62
gpt-oss 20B MXFP4 MoE	tg128	260.27	269.68	1.04
gpt-oss 20B MXFP4 MoE	pp512 @ d65536	4331.62	5209.42	1.20
gpt-oss 20B MXFP4 MoE	pp2048 @ d65536	7116.87	5563.52	0.78
gpt-oss 20B MXFP4 MoE	tg128 @ d65536	172.17	204.50	1.19
llama 1B F16	pp512	47450.84	30410.70	0.64
llama 1B F16	pp2048	73060.68	28377.32	0.39
llama 1B F16	tg128	324.01	374.62	1.16
llama 1B F16	pp512 @ d65536	8604.25	9659.83	1.12
llama 1B F16	pp2048 @ d65536	15157.90	11101.89	0.73
llama 1B F16	tg128 @ d65536	181.99	248.22	1.36
llama 8B F16	pp512	10624.32	8119.05	0.76
llama 8B F16	pp2048	15854.74	7654.63	0.48
llama 8B F16	tg128	60.28	96.63	1.60
llama 8B F16	pp512 @ d65536	3043.24	3779.64	1.24
llama 8B F16	pp2048 @ d65536	4791.03	4032.86	0.84
llama 8B F16	tg128 @ d65536	38.81	66.63	1.72
qwen35 27B Q4_0	pp512	3221.77	2927.23	0.91
qwen35 27B Q4_0	pp2048	4286.78	2712.77	0.63
qwen35 27B Q4_0	tg128	51.06	61.82	1.21
qwen35 27B Q4_0	pp512 @ d65536	1843.62	2083.90	1.13
qwen35 27B Q4_0	pp2048 @ d65536	2486.88	2080.78	0.84
qwen35 27B Q4_0	tg128 @ d65536	41.25	53.48	1.30
qwen35 27B Q8_0	pp512	2767.40	2839.12	1.03
qwen35 27B Q8_0	pp2048	3727.66	2639.67	0.71
qwen35 27B Q8_0	tg128	30.61	43.03	1.41
qwen35 27B Q8_0	pp512 @ d65536	1692.06	2044.16	1.21
qwen35 27B Q8_0	pp2048 @ d65536	2295.14	2038.38	0.89
qwen35 27B Q8_0	tg128 @ d65536	26.78	38.54	1.44
qwen35moe 35B.A3B Q4_0	pp512	7347.60	7161.24	0.97
qwen35moe 35B.A3B Q4_0	pp2048	9872.26	6217.26	0.63
qwen35moe 35B.A3B Q4_0	tg128	205.74	141.41	0.69
qwen35moe 35B.A3B Q4_0	pp512 @ d65536	4139.60	4654.36	1.12
qwen35moe 35B.A3B Q4_0	pp2048 @ d65536	5500.91	4644.75	0.84
qwen35moe 35B.A3B Q4_0	tg128 @ d65536	154.43	125.49	0.81
qwen35moe 35B.A3B Q8_0	pp512	6947.57	6989.91	1.01
qwen35moe 35B.A3B Q8_0	pp2048	9451.57	6068.24	0.64
qwen35moe 35B.A3B Q8_0	tg128	158.23	126.58	0.80
qwen35moe 35B.A3B Q8_0	pp512 @ d65536	4019.37	4429.08	1.10
qwen35moe 35B.A3B Q8_0	pp2048 @ d65536	5318.73	4534.64	0.85
qwen35moe 35B.A3B Q8_0	tg128 @ d65536	126.66	107.34	0.85

Performance on 4x RTX 4090

model	test	t/s -sm layer	t/s -sm tensor	Speedup
gpt-oss 120B MXFP4 MoE	pp512	3822.24	5446.21	1.42
gpt-oss 120B MXFP4 MoE	pp8192	4675.65	5211.46	1.11
gpt-oss 120B MXFP4 MoE	tg128	180.36	179.79	1.00
gpt-oss 120B MXFP4 MoE	pp512 @ d65536	2308.57	3538.01	1.53
gpt-oss 120B MXFP4 MoE	pp8192 @ d65536	3034.32	3762.05	1.24
gpt-oss 120B MXFP4 MoE	tg128 @ d65536	118.65	147.89	1.25
llama 1B F16	pp512	44766.52	18753.39	0.42
llama 1B F16	pp8192	50398.93	17823.13	0.35
llama 1B F16	tg128	314.75	409.87	1.30
llama 1B F16	pp512 @ d65536	7528.70	7263.59	0.96
llama 1B F16	pp8192 @ d65536	11519.92	7805.30	0.68
llama 1B F16	tg128 @ d65536	178.13	307.56	1.73
llama 8B F16	pp512	10409.04	4262.81	0.41
llama 8B F16	pp8192	11816.43	4112.37	0.35
llama 8B F16	tg128	59.88	143.65	2.40
llama 8B F16	pp512 @ d65536	2884.76	2707.52	0.94
llama 8B F16	pp8192 @ d65536	3844.32	2915.65	0.76
llama 8B F16	tg128 @ d65536	38.65	104.49	2.70
llama 70B Q4_K_M	pp512	1439.14	1313.79	0.91
llama 70B Q4_K_M	pp8192	1701.36	1264.31	0.74
llama 70B Q4_K_M	tg128	21.25	48.31	2.27
llama 70B Q4_K_M	pp512 @ d65536	574.89	929.93	1.62
llama 70B Q4_K_M	pp8192 @ d65536	726.52	922.87	1.27
llama 70B Q4_K_M	tg128 @ d65536	14.26	36.68	2.57
glm4moe 106B.A12B IQ4_XS	pp512	2664.37	2983.24	1.12
glm4moe 106B.A12B IQ4_XS	pp8192	2931.28	2762.00	0.94
glm4moe 106B.A12B IQ4_XS	tg128	98.64	115.14	1.17
glm4moe 106B.A12B IQ4_XS	pp512 @ d65536	653.92	1400.18	2.14
glm4moe 106B.A12B IQ4_XS	pp8192 @ d65536	799.05	1435.96	1.80
glm4moe 106B.A12B IQ4_XS	tg128 @ d65536	42.64	79.20	1.86

All GPUs are connected via 16x PCIe 4.0. pp512 is representative of single-GPU performance since pipelining is impossible, pp2048 and pp8192 are representative of multi-GPU performance with pipelining.

By default models are run with --split-mode layer which runs the GPUs sequentially. For the prompt this is efficient because the tokens can be pipelined with minimal synchronization overhead. However, when generating tokens for a single concurrent request there is no speedup because the GPUs have to run sequentially. --split-mode tensor works for all cases but the synchronization overhead may be prohibitive. The key factor is how long the base runtime is vs. the synchronization overhead; if the model runs slow by default adding a bit of overhead hurts less, if it is already fast adding even a small amount of overhead is really noticeable. So --split-mode tensor works comparatively better for slow GPUs with fast interconnect running large, dense models at large quantiztations and deep contexts, and comparatively worse for fast GPUs with slow interconnect running small, sparse models at small quantizations and shallow contexts.
Similarly, because the runtime per token for the prompt is much shorter than when generating them, the speedup for generation is also much better than for the prompt.

Original PR description

Details

This PR adds support for backend-agnostic tensor parallelism, enabled via specifying `--split-mode tensor`. This is done by adding a new "meta" backend that internally wraps multiple "simple" backends but can be used in the same way as a regular ggml backend.

ggml Backend Interface Changes

This PR extends the ggml backend interface with some new functions for tensor copying:

• set_tensor_2d_async/get_tensor_2d_async which are equivalent to memcpy2DAsync in CUDA. This is not needed for the computation of the meta backend itself but rather for setting/getting weights or the output. Currently not implemented, as a workaround the one-dimensional version is used in a loop.
• shfl_tensor_async to allow two ggml backends to exchange two tensors and to synchronize on the completion of the exchange. As a fallback cpy_tensor_async can be used but this has a higher latency because the copy in one direction can only start once the one in the other direction has finished. Needed for a generic AllReduce between ggml backends. Implemented.
• allreduce_tensor_async to allow ggml backends to specify a backend-specific way to do an AllReduce operation. Intended to be used for NCCL support in cooperation with @gaugarg-nv . Not yet implemented.

@slaren please provide feedback regarding whether you agree that these operations should be in the ggml backend interface. For context, all of them are optional and can use existing operations as a fallback.

Meta Backend Implementation Details

The meta backend implements an entire ggml backend stack starting from a meta device. The meta device is created from multiple simple backend devices as well as a function to determine how the data should be split across devices ("split states"). Backend buffer types, buffers, and backends are created as per usual. When calling ggml_backend_graph_compute the code infers the split states of the nodes in the compute graph based on the split states assigned for the weights/kv cache. The basic pattern is to make all tensors mirrored by default. For the weight matrices, do a split in dimension 1, then a split in dimension 0, then an AllReduce. For a transformer this means two AllReduce operations, one after the attention and one after the FFN. The attention is effectively split by dimension 0, which equates to a split by attention head.

An generic AllReduce operation is performed in the meta backend by splitting the graph into subgraphs. After a subgraph is executed, call shfl_tensor_async to make backends exchange partial results, and then have them execute auxiliary graphs that contain only a GGML_ADD operation to combine the results.

The memory allocation for the compute graph is rather tricky - the way I solved it is to allocate the memory for the meta backend as per usual and to then transplant the calculated addresses relative to the backend buffer base pointer to the underlying simple backends. Because the simple tensors only require a fraction of the full memory this yields correct results, though it does result in overallocation for the compute graphs. For the weights/kv cache the memory allocation for the meta backend is done via a new function ggml_backend_meta_alloc_ctx_tensors_from_buft to prevent duplicated weights (which are much larger in size). I'm not yet sure what the best approach will be long-term, I think the graph allocation code in ggml-alloc.c will need to be adjusted.

Current Issues/Limitations

• Only 1 or 2 GPUs are supported. Note: 1 GPU is not actually any faster, this is only useful for testing whether the code works correctly.
• All GPUs must have an equal share of the data, --tensor-split has no effect.
• Only dense models are supported. The LLaMA 3 models seem to be working correctly, I have not yet tested others.
• Support for llama_params_fit is not implemented so the context size has to be set manually.
• Without FlashAttention the code will probably crash because some transition between split states is not yet implemented.
• In principle all backends should work. CUDA does in my testing, Vulkan however does not. I think there may be some issues with deadlock between the GPUs. @jeffbolznv @0cc4m if you could take a look it would be appreciated.
• Memory for the ggml contexts is being overallocated.
• Performance is (presumably) still suboptimal vs. NCCL.
• I'm currently using tensor names to determine how to split individual tensors. I think it would be preferable to use some sort of enum instead (which we already seem to have for loading tensors). This should also be used for llama_params_fit.
• I'm currently setting ggml_tensor::data to dummy values since that is what is checked in ggml-alloc.c to determine whether or not a tensor is considered allocated. This dummy value should never actually be used for any computations but I don't consider this a good solution.
• I'm not putting meta devices into the ggml backend device registry (which I think is correct).

Performance

LLaMA 3 on 2x RTX 4090

model	test	t/s -sm layer	t/s -sm row	t/s -sm tensor
llama 8B Q4_0	pp512	12550.97	2997.41	6305.68
llama 8B Q4_0	pp2048	18788.11	2970.83	6300.12
llama 8B Q4_0	tg128	175.43	67.00	101.98
llama 8B Q4_0	pp512 @ d32768	5099.65	2200.50	4137.73
llama 8B Q4_0	pp2048 @ d32768	7925.21	2242.84	4337.54
llama 8B Q4_0	tg128 @ d32768	96.78	49.76	102.81
llama 8B Q4_0	pp512 @ d65536	3154.69	1748.05	3139.40
llama 8B Q4_0	pp2048 @ d65536	4996.19	1806.82	3404.70
llama 8B Q4_0	tg128 @ d65536	67.11	40.27	83.62
llama 8B Q4_0	pp512 @ d131072	1800.72	1243.57	2152.01
llama 8B Q4_0	pp2048 @ d131072	2867.83	1294.61	2238.01
llama 8B Q4_0	tg128 @ d131072	41.66	29.19	59.53
llama 8B F16	pp512	10578.54	1591.55	5950.44
llama 8B F16	pp2048	15890.45	1581.04	6005.96
llama 8B F16	tg128	60.07	46.32	70.64
llama 8B F16	pp512 @ d32768	4745.19	1330.02	4032.19
llama 8B F16	pp2048 @ d32768	**7329.28 **	1347.02	4279.02
llama 8B F16	tg128 @ d32768	47.03	38.02	64.63
llama 8B F16	pp512 @ d65536	3033.10	1154.54	3100.10
llama 8B F16	pp2048 @ d65536	4782.02	1176.88	3341.03
llama 8B F16	tg128 @ d65536	38.72	32.19	56.48
llama 8B F16	pp512 @ d131072	1735.63	905.39	2090.03
llama 8B F16	pp2048 @ d131072	2782.42	936.55	2288.65
llama 8B F16	tg128 @ d131072	28.60	24.79	43.32
llama 70B Q3_K - Small	pp512	1287.50	582.81	1072.80
llama 70B Q3_K - Small	pp2048	1954.83	590.45	1069.28
llama 70B Q3_K - Small	tg128	27.97	20.36	29.65
llama 70B Q3_K - Small	pp512 @ d32768	776.80	458.17	812.77
llama 70B Q3_K - Small	pp2048 @ d32768	1185.82	459.43	824.99
llama 70B Q3_K - Small	tg128 @ d32768	20.85	16.19	29.39

Generally speaking it can be observed that parallelizing larger models has better performance than parallelizing smaller models. Similarly, parallelizing the model becomes more worthwhile as the context depth increases. This makes sense as both of these result in a larger workload per GPU vs. the overhead from parallelization. Token generation benefits more from parallelization than prompt processing because the amount of data that needs to be transferred between GPUs is proportional to batch size - long-term it may make sense to implement support for FP16/BF16 compute types which would count the I/O in half vs. FP32. For pp512 pipeline parallelism is effectively disabled while for pp2048 it's enabled. With pipeline parallelism -sm layer is still faster than -sm tensor even at high context depths.

[jeffbolznv]

In principle all backends should work. CUDA does in my testing, Vulkan however does not. I think there may be some issues with deadlock between the GPUs. @jeffbolznv @0cc4m if you could take a look it would be appreciated.

I'm not seeing a deadlock, just a crash in the driver with an invalid descriptor. I ran llama-bench.exe -fa 1 -p 512 -n 0 -m c:\models\llama-2-7b.Q4_0.gguf -sm tensor

Validation Error: [ VUID-VkDescriptorBufferInfo-offset-00340 ] | MessageID = 0xc23dafe5
vkUpdateDescriptorSets(): pDescriptorWrites[0].pBufferInfo[2].offset (18362368) is greater than or equal to buffer size (8388608).
The Vulkan spec states: offset must be less than the size of buffer (https://docs.vulkan.org/spec/latest/chapters/descriptorsets.html#VUID-VkDescriptorBufferInfo-offset-00340)

-		dst_buf	{buffer=shared_ptr {buffer={m_buffer=0x0000be00000000be {...} } device_memory={m_deviceMemory=0x0000bf00000000bf {...} } ...} [0x00000003 strong refs] [make_shared] ...}	vk_subbuffer
+		buffer	shared_ptr {buffer={m_buffer=0x0000be00000000be {...} } device_memory={m_deviceMemory=0x0000bf00000000bf {...} } ...} [0x00000003 strong refs] [make_shared]	std::shared_ptr<vk_buffer_struct>
		offset	0x0000000001183000	unsigned __int64
		size	0x0000000000800000	unsigned __int64
-		dst	0x0000018f98716fd0 {type=GGML_TYPE_F32 (0x00000000) buffer=0x0000018f98385c80 {iface={free_buffer=0x00007ffbc6647c50 {ggml-vulkan.dll!ggml_backend_vk_buffer_free_buffer(ggml_backend_buffer *)} ...} ...} ...}	ggml_tensor *
		type	GGML_TYPE_F32 (0x00000000)	ggml_type
+		buffer	0x0000018f98385c80 {iface={free_buffer=0x00007ffbc6647c50 {ggml-vulkan.dll!ggml_backend_vk_buffer_free_buffer(ggml_backend_buffer *)} ...} ...}	ggml_backend_buffer *
+		ne	0x0000018f98716fe0 {0x0000000000001000, 0x0000000000000200, 0x0000000000000001, 0x0000000000000001}	__int64[0x00000004]
+		nb	0x0000018f98717000 {0x0000000000000004, 0x0000000000004000, 0x0000000000800000, 0x0000000000800000}	unsigned __int64[0x00000004]
		op	GGML_OP_ADD (0x00000002)	ggml_op
+		op_params	0x0000018f98717024 {0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, ...}	int[0x00000010]
		flags	0x00000010	int
+		src	0x0000018f98717068 {0x00000191a1b1d1b0 {type=GGML_TYPE_F32 (0x00000000) buffer=0x0000018f97d4f130 {iface=...} ...}, ...}	ggml_tensor *[0x0000000a]
-		view_src	0x00000191a1b1d1b0 {type=GGML_TYPE_F32 (0x00000000) buffer=0x0000018f97d4f130 {iface={free_buffer=0x00007ffbc6647c50 {ggml-vulkan.dll!ggml_backend_vk_buffer_free_buffer(ggml_backend_buffer *)} ...} ...} ...}	ggml_tensor *
		type	GGML_TYPE_F32 (0x00000000)	ggml_type
+		buffer	0x0000018f97d4f130 {iface={free_buffer=0x00007ffbc6647c50 {ggml-vulkan.dll!ggml_backend_vk_buffer_free_buffer(ggml_backend_buffer *)} ...} ...}	ggml_backend_buffer *
+		ne	0x00000191a1b1d1c0 {0x0000000000001000, 0x0000000000000200, 0x0000000000000001, 0x0000000000000001}	__int64[0x00000004]
+		nb	0x00000191a1b1d1e0 {0x0000000000000004, 0x0000000000004000, 0x0000000000800000, 0x0000000000800000}	unsigned __int64[0x00000004]
		op	GGML_OP_MUL_MAT (0x0000001d)	ggml_op
+		op_params	0x00000191a1b1d204 {0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, ...}	int[0x00000010]
		flags	0x00000010	int
+		src	0x00000191a1b1d248 {0x0000018f9f29d940 {type=GGML_TYPE_Q4_0 (0x00000002) buffer=0x0000018f97d4f670 {...} ...}, ...}	ggml_tensor *[0x0000000a]
+		view_src	0x0000000000000000 <NULL>	ggml_tensor *
		view_offs	0x0000000000000000	unsigned __int64
		data	0x0000000001184000	void *
+		name	0x00000191a1b1d2b0 "attn_out-0"	char[0x00000040]
		extra	0x0000000000000000	void *
+		padding	0x00000191a1b1d2f8 ""	char[0x00000008]
		view_offs	0x0000000000000000	unsigned __int64
		data	0x0000000001184000	void *
+		name	0x0000018f987170d0 "attn_out-0 (view)"	char[0x00000040]
		extra	0x0000000000000000	void *
+		padding	0x0000018f98717118 ""	char[0x00000008]

It seems like tensor->data (and tensor->view_src->data) are too large. I haven't debugged further.

[jacekpoplawski]

works for me on 2x3090 for llama 3 8B and Mistral Nemo 12B

on Devstral I have OOM (expected because model size?)

ggml_backend_cuda_buffer_type_alloc_buffer: allocating 30720.00 MiB on device 0: cudaMalloc failed: out of memory
alloc_tensor_range: failed to allocate CUDA0 buffer of size 32212254720

Thread 1 "llama-server" received signal SIGABRT, Aborted.
Download failed: Invalid argument.  Continuing without source file ./nptl/./nptl/pthread_kill.c.
__pthread_kill_implementation (threadid=<optimized out>, signo=6, no_tid=0) at ./nptl/pthread_kill.c:44
warning: 44     ./nptl/pthread_kill.c: No such file or directory
(gdb) bt
#0  __pthread_kill_implementation (threadid=<optimized out>, signo=6, no_tid=0) at ./nptl/pthread_kill.c:44
#1  __pthread_kill_internal (threadid=<optimized out>, signo=6) at ./nptl/pthread_kill.c:89
#2  __GI___pthread_kill (threadid=<optimized out>, signo=signo@entry=6) at ./nptl/pthread_kill.c:100
#3  0x00007fffedc4579e in __GI_raise (sig=sig@entry=6) at ../sysdeps/posix/raise.c:26
#4  0x00007fffedc288cd in __GI_abort () at ./stdlib/abort.c:73
#5  0x00007ffff7719ec5 in ggml_abort (file=0x7ffff77bdaa8 "/home/jacek/git/llama.cpp/ggml/src/ggml-backend.cpp", line=119, fmt=0x7ffff77bd88f "GGML_ASSERT(%s) failed")
    at /home/jacek/git/llama.cpp/ggml/src/ggml.c:256
#6  0x00007ffff773565b in ggml_backend_buffer_get_size (buffer=0x0) at /home/jacek/git/llama.cpp/ggml/src/ggml-backend.cpp:119
#7  0x00007ffff7743b0e in ggml_backend_meta_alloc_ctx_tensors_from_buft (ctx=0x555557dc41f0, buft=0x55555b032b18) at /home/jacek/git/llama.cpp/ggml/src/ggml-backend-meta.cpp:638
#8  0x00007ffff7734a4a in ggml_backend_alloc_ctx_tensors_from_buft (ctx=0x555557dc41f0, buft=0x55555b032b18) at /home/jacek/git/llama.cpp/ggml/src/ggml-alloc.c:1245
#9  0x00007ffff6f3d6c2 in llama_kv_cache::llama_kv_cache (this=0x555556363770, model=..., type_k=GGML_TYPE_F16, type_v=GGML_TYPE_F16, v_trans=false, offload=true, unified=true, kv_size=393216,
    n_seq_max=4, n_pad=1, n_swa=0, swa_type=LLAMA_SWA_TYPE_NONE, filter=..., reuse=...) at /home/jacek/git/llama.cpp/src/llama-kv-cache.cpp:190
#10 0x00007ffff700faec in llama_model::create_memory (this=0x555556357ea0, params=..., cparams=...) at /home/jacek/git/llama.cpp/src/llama-model.cpp:7617
#11 0x00007ffff6ed008f in llama_context::llama_context (this=0x555557a65260, model=..., params=...) at /home/jacek/git/llama.cpp/src/llama-context.cpp:274
#12 0x00007ffff6edd53c in llama_init_from_model (model=0x555556357ea0, params=...) at /home/jacek/git/llama.cpp/src/llama-context.cpp:3046
#13 0x00005555558f4e24 in common_init_result::common_init_result (this=0x55555624dc70, params=...) at /home/jacek/git/llama.cpp/common/common.cpp:1183
#14 0x00005555558f50d0 in common_init_from_params (params=...) at /home/jacek/git/llama.cpp/common/common.cpp:1215
#15 0x0000555555710999 in server_context_impl::load_model (this=0x55555636eba0, params=...) at /home/jacek/git/llama.cpp/tools/server/server-context.cpp:625
#16 0x00005555556e9ac8 in server_context::load_model (this=0x7fffffff79b0, params=...) at /home/jacek/git/llama.cpp/tools/server/server-context.cpp:2856
#17 0x00005555556177bd in main (argc=7, argv=0x7fffffffe0c8) at /home/jacek/git/llama.cpp/tools/server/server.cpp:248
(gdb)

but on MInistral 14B too

ggml_backend_cuda_buffer_type_alloc_buffer: allocating 20480.00 MiB on device 0: cudaMalloc failed: out of memory
alloc_tensor_range: failed to allocate CUDA0 buffer of size 21474836480

Thread 1 "llama-server" received signal SIGABRT, Aborted.
Download failed: Invalid argument.  Continuing without source file ./nptl/./nptl/pthread_kill.c.
__pthread_kill_implementation (threadid=<optimized out>, signo=6, no_tid=0) at ./nptl/pthread_kill.c:44
warning: 44     ./nptl/pthread_kill.c: No such file or directory
(gdb) bt
#0  __pthread_kill_implementation (threadid=<optimized out>, signo=6, no_tid=0) at ./nptl/pthread_kill.c:44
#1  __pthread_kill_internal (threadid=<optimized out>, signo=6) at ./nptl/pthread_kill.c:89
#2  __GI___pthread_kill (threadid=<optimized out>, signo=signo@entry=6) at ./nptl/pthread_kill.c:100
#3  0x00007fffedc4579e in __GI_raise (sig=sig@entry=6) at ../sysdeps/posix/raise.c:26
#4  0x00007fffedc288cd in __GI_abort () at ./stdlib/abort.c:73
#5  0x00007ffff7719ec5 in ggml_abort (file=0x7ffff77bdaa8 "/home/jacek/git/llama.cpp/ggml/src/ggml-backend.cpp", line=119, fmt=0x7ffff77bd88f "GGML_ASSERT(%s) failed")
    at /home/jacek/git/llama.cpp/ggml/src/ggml.c:256
#6  0x00007ffff773565b in ggml_backend_buffer_get_size (buffer=0x0) at /home/jacek/git/llama.cpp/ggml/src/ggml-backend.cpp:119
#7  0x00007ffff7743b0e in ggml_backend_meta_alloc_ctx_tensors_from_buft (ctx=0x555557e339e0, buft=0x55555abb40b8) at /home/jacek/git/llama.cpp/ggml/src/ggml-backend-meta.cpp:638
#8  0x00007ffff7734a4a in ggml_backend_alloc_ctx_tensors_from_buft (ctx=0x555557e339e0, buft=0x55555abb40b8) at /home/jacek/git/llama.cpp/ggml/src/ggml-alloc.c:1245
#9  0x00007ffff6f3d6c2 in llama_kv_cache::llama_kv_cache (this=0x55555abb9470, model=..., type_k=GGML_TYPE_F16, type_v=GGML_TYPE_F16, v_trans=false, offload=true, unified=true, kv_size=262144,
    n_seq_max=4, n_pad=1, n_swa=0, swa_type=LLAMA_SWA_TYPE_NONE, filter=..., reuse=...) at /home/jacek/git/llama.cpp/src/llama-kv-cache.cpp:190
#10 0x00007ffff700faec in llama_model::create_memory (this=0x555556357e50, params=..., cparams=...) at /home/jacek/git/llama.cpp/src/llama-model.cpp:7617
#11 0x00007ffff6ed008f in llama_context::llama_context (this=0x555557a65140, model=..., params=...) at /home/jacek/git/llama.cpp/src/llama-context.cpp:274
#12 0x00007ffff6edd53c in llama_init_from_model (model=0x555556357e50, params=...) at /home/jacek/git/llama.cpp/src/llama-context.cpp:3046
#13 0x00005555558f4e24 in common_init_result::common_init_result (this=0x55555624dc70, params=...) at /home/jacek/git/llama.cpp/common/common.cpp:1183
#14 0x00005555558f50d0 in common_init_from_params (params=...) at /home/jacek/git/llama.cpp/common/common.cpp:1215
#15 0x0000555555710999 in server_context_impl::load_model (this=0x55555636eb30, params=...) at /home/jacek/git/llama.cpp/tools/server/server-context.cpp:625
#16 0x00005555556e9ac8 in server_context::load_model (this=0x7fffffff79b0, params=...) at /home/jacek/git/llama.cpp/tools/server/server-context.cpp:2856
#17 0x00005555556177bd in main (argc=7, argv=0x7fffffffe0c8) at /home/jacek/git/llama.cpp/tools/server/server.cpp:248
(gdb)

qwen 4B has a different issue

/home/jacek/git/llama.cpp/ggml/src/ggml-backend-meta.cpp:386: GGML_ASSERT(ne[split_dim] % n_simple_bufs == 0) failed

Thread 1 "llama-server" received signal SIGABRT, Aborted.
Download failed: Invalid argument.  Continuing without source file ./nptl/./nptl/pthread_kill.c.
__pthread_kill_implementation (threadid=<optimized out>, signo=6, no_tid=0) at ./nptl/pthread_kill.c:44
warning: 44     ./nptl/pthread_kill.c: No such file or directory
(gdb) bt
#0  __pthread_kill_implementation (threadid=<optimized out>, signo=6, no_tid=0) at ./nptl/pthread_kill.c:44
#1  __pthread_kill_internal (threadid=<optimized out>, signo=6) at ./nptl/pthread_kill.c:89
#2  __GI___pthread_kill (threadid=<optimized out>, signo=signo@entry=6) at ./nptl/pthread_kill.c:100
#3  0x00007fffedc4579e in __GI_raise (sig=sig@entry=6) at ../sysdeps/posix/raise.c:26
#4  0x00007fffedc288cd in __GI_abort () at ./stdlib/abort.c:73
#5  0x00007ffff7719ec5 in ggml_abort (file=0x7ffff77bedd8 "/home/jacek/git/llama.cpp/ggml/src/ggml-backend-meta.cpp", line=386, fmt=0x7ffff77beb9f "GGML_ASSERT(%s) failed")
    at /home/jacek/git/llama.cpp/ggml/src/ggml.c:256
#6  0x00007ffff7741e84 in ggml_backend_meta_buffer_init_tensor (buffer=0x55555a4421d0, tensor=0x55555a46df70) at /home/jacek/git/llama.cpp/ggml/src/ggml-backend-meta.cpp:386
#7  0x00007ffff7743a52 in ggml_backend_meta_alloc_ctx_tensors_from_buft (ctx=0x55555a4422f0, buft=0x55555a426188) at /home/jacek/git/llama.cpp/ggml/src/ggml-backend-meta.cpp:632
#8  0x00007ffff7734a4a in ggml_backend_alloc_ctx_tensors_from_buft (ctx=0x55555a4422f0, buft=0x55555a426188) at /home/jacek/git/llama.cpp/ggml/src/ggml-alloc.c:1245
#9  0x00007ffff6fff452 in llama_model::load_tensors (this=0x555556357e50, ml=...) at /home/jacek/git/llama.cpp/src/llama-model.cpp:7055
#10 0x00007ffff6e557fc in llama_model_load (fname="/mnt/models2/Qwen/Qwen_Qwen3-4B-Q8_0.gguf", splits=std::vector of length 0, capacity 0, model=..., params=...)
    at /home/jacek/git/llama.cpp/src/llama.cpp:876
#11 0x00007ffff6e56ce3 in llama_model_load_from_file_impl (path_model="/mnt/models2/Qwen/Qwen_Qwen3-4B-Q8_0.gguf", splits=std::vector of length 0, capacity 0, params=...)
    at /home/jacek/git/llama.cpp/src/llama.cpp:1069
#12 0x00007ffff6e56fe3 in llama_model_load_from_file (path_model=0x555556367610 "/mnt/models2/Qwen/Qwen_Qwen3-4B-Q8_0.gguf", params=...) at /home/jacek/git/llama.cpp/src/llama.cpp:1096
#13 0x00005555558f46c9 in common_init_result::common_init_result (this=0x55555624dc70, params=...) at /home/jacek/git/llama.cpp/common/common.cpp:1107
#14 0x00005555558f50d0 in common_init_from_params (params=...) at /home/jacek/git/llama.cpp/common/common.cpp:1215
#15 0x0000555555710999 in server_context_impl::load_model (this=0x55555636ead0, params=...) at /home/jacek/git/llama.cpp/tools/server/server-context.cpp:625
#16 0x00005555556e9ac8 in server_context::load_model (this=0x7fffffff79d0, params=...) at /home/jacek/git/llama.cpp/tools/server/server-context.cpp:2856
#17 0x00005555556177bd in main (argc=7, argv=0x7fffffffe0e8) at /home/jacek/git/llama.cpp/tools/server/server.cpp:248
(gdb)

[DocShotgun]

Interesting. If the CPU backend is able to be virtualized into multiple devices as described here, would it be possible to allow multiple NUMA nodes to be parallelized?

[gopinath87607]

will this pr solve if we use multiple rpc connected to gpu with cpu and when we use cpu moe flag ?

[JohannesGaessler]

@jacekpoplawski the combination of --split-mode tensor and llama_params_fit is not implemented so you'll have to set the context size manually if you didn't already.

@DocShotgun longer-term I intend to also enable this code for better NUMA support though I'm not yet sure what to do in terms of hardware for development. Originally I had intended to buy 1.5 TiB of DDR5 RAM and 2 EPYC CPUs but at the current prices that would be financially irresponsible of me to do.

@gopinath87607 I don't understand what you mean.

[FullstackSensei]

@DocShotgun longer-term I intend to also enable this code for better NUMA support though I'm not yet sure what to do in terms of hardware for development. Originally I had intended to buy 1.5 TiB of DDR5 RAM and 2 EPYC CPUs but at the current prices that would be financially irresponsible of me to do.

If you're looking for a DDR4 platform, I might be able to help with that, but unfortunately, not so much with RAM for the system. I'm also in Germany.

Wouldn't mind giving you access to my systems if you want. Have two dual Xeon systems one with P40s and the other with Mi50s.

Would be very happy to help either way.

[ggerganov]

Started doing some initial tests to get familiar with the changes. I'm using virtual Metal devices and things appear to be mostly working - e.g. seeing graph execution on both devices, llama-perplexity produces the same result with 2 devices.

However, the following command does not produce identical results on each run:

GGML_METAL_DEVICES=2 ./bin/llama-completion -m ~/models/llama-3.1-8b/ggml-model-f16.gguf -no-cnv -p "I believe the meaning of life is" -n 32 --top-k 1 -sm tensor

I think this either means that there could be a problem with the fallback for the missing backend API, or that I could have made an error in the implementation of the events and the async copy in Metal. Will investigate more.

For context, all of them are optional and can use existing operations as a fallback.

In the meantime, @JohannesGaessler do you confirm that the command above has deterministic results on your end? Also, if is it deterministic with the fallback calls?

[JohannesGaessler]

Generally speaking you cannot expect bit-for-bit identical results if you split the computation across multiple virtual devices. The order in which floats are summed up will be different which will in turn change the rounding error. If I run llama-perplexity using LLaMA 3 8b f16 I get 6.2560 for -sm layer and 6.2556 for -sm tensor. It's of course still possible that there are bugs on top of that, I would just say changes to the results are expected. Long-term I think we should test this new code in the same way we test ggml op fusion in test-backend-ops.

If you use -sm tensor with only a single GPU the executed ops should be the exact same and the result should be bit-for-bit identical to -sm layer.

[JohannesGaessler]

Sorry, I think I misread your post. If you are saying that the results are not deterministic with 2 virtual GPUs but they are with 1 GPU then that I think is indicative of a bug w.r.t. the synchronization.

[ggerganov]

Yes, I understand that 1GPU vs 2GPU will not be bit-for-bit identical. What I mean is that in my test, running the command with 2 GPUs several times produces non-deterministic results from one run to the other:

GGML_METAL_DEVICES=2 ./bin/llama-completion -m ~/models/llama-3.1-8b/ggml-model-f16.gguf -no-cnv -p "I believe the meaning of life is" -n 32 --top-k 1 -sm tensor

# run 1
I believe the meaning of life is to find your gift. The purpose of life is to give it away.
I believe that the meaning of life is to find your gift. The purpose of life

# run 2
I believe the meaning of life is to find your gift. The purpose of life is to give it away. To give it away, you have to find it. [end of text]

# run 3
I believe the meaning of life is to be happy. I believe that happiness is the only thing that matters. I believe that happiness is the only thing that matters. I believe that happiness is the

Sorry, I think I misread your post. If you are saying that the results are not deterministic with 2 virtual GPUs but they are with 1 GPU then that I think is indicative of a bug w.r.t. the synchronization.

Yes, seems like a synchronization issue. Was wondering if you observe it on your end with and without the fallback backend API. This will give me indication where to look for the issue.

[JohannesGaessler]

With the command you posted (minus the GGML_METAL_DEVICES=2) I get deterministic results on 2x RTX 4090, both with and without the fallback for shfl_tensor_async. The output of llama-perplexity is bit-for-bit identical with vs. without the fallback. So presumably either there is a bug with Metal synchronization or I made assumptions about the behavior of ggml backends that are correct for CUDA but not universally.

[pwilkin]

Ack for arg.cpp changes ;) BTW, shouldn't we be deprecating --split-mode row with this?

fixed

[gopinath87607]

hope one day this --split-mode tensor will work on over the rpc too btw nice work guys.

[pwilkin]

Congratulations on this :)

Wanted to try, but alas:
llama_init_from_model: simultaneous use of SPLIT_MODE_TENSOR and KV cache quantization not implemented #ohwell

[IMbackK]

hope one day this --split-mode tensor will work on over the rpc too btw nice work guys.

its going to be pretty brutal across consumer rpc interconnects

[FullstackSensei]

hope one day this --split-mode tensor will work on over the rpc too btw nice work guys.

Even a 40gb NIC will be too slow to handle this on any sizeable models.

[FullstackSensei]

Congratulations everyone! This was a gig effort!

[digitalscream]

Nice work, all - this was the only feature that kept me looking sideways at vLLM.

Now, how many beers do I need to buy the Vulkan team to get them to optimise for it? :D

[nawoa]

Would the 200G NCCL/RDMA link between DGX Spark nodes be compatible with this?

If I understand correctly this PR allows for arbitrary tensor splits. This might allow Sparks in a 3-node ring topology (100G interconnect between each unit) to actually be useful. Currently with vLLM the only options are 2/4/8.

I don't know what performance would look like in Llama vs vLLM but 384GB of RAM @ ~820 GB/s effective RAM bandwidth would be pretty nice. Saves the cost and latency of a CX7 switch.

Sorry if I'm off-base.

[Ph0rk0z]

bit of a hmm.. i can test perplexity of the gemma base at Q8_0 and everything is fine.

When I do it to q8_0_xl of the IT version it always craps itself out at iteration 260/576. I tried removing undervolt and some other things. In the case of using cuda devices 0,1,2,3 the GPU0 PCIE link is also taken down to 2x until I reboot despite there being no lane errors or any other messages. When I flip it around to 3,2,1,0 the test still crashes in the same place but I don't eat the link and I can keep re-running things.

1283.1084,[258]11271.1522,[259]11265.8000,[260]11286.2032,/home/supermicro/ai/llama.cpp.main/ggml/src/ggml-cuda/ggml-cuda.cu:97: CUDA error
CUDA error: an internal operation failed
  current device: 0, in function ggml_cuda_op_mul_mat_cublas at /home/supermicro/ai/llama.cpp.main/ggml/src/ggml-cuda/ggml-cuda.cu:1504
  cublasGemmEx(ctx.cublas_handle(id), CUBLAS_OP_T, CUBLAS_OP_N, row_diff, src1_ncols, ne10, &alpha_f16, src0_ptr, CUDA_R_16F, ne00, src1_ptr, CUDA_R_16F, ne10, &beta_f16, dst_f16.get(), CUDA_R_16F, ldc, CUBLAS_COMPUTE_16F, CUBLAS_GEMM_DEFAULT_TENSOR_OP)
[New LWP 136334]
[New LWP 136348]
[New LWP 136349]
[New LWP 136350]
[New LWP 136351]
[New LWP 136352]
[New LWP 136353]
[New LWP 136354]
[New LWP 136355]
[New LWP 136356]
[New LWP 136357]
[New LWP 136363]
[New LWP 136364]
[New LWP 136365]
[New LWP 136366]
[New LWP 136367]
[New LWP 136368]
[New LWP 136369]
[New LWP 136370]
[New LWP 136371]
[Thread debugging using libthread_db enabled]
Using host libthread_db library "/lib/x86_64-linux-gnu/libthread_db.so.1".
0x00007f1df14ea42f in __GI___wait4 (pid=161186, stat_loc=0x0, options=0, usage=0x0) at ../sysdeps/unix/sysv/linux/wait4.c:30
30	../sysdeps/unix/sysv/linux/wait4.c: No such file or directory.
#0  0x00007f1df14ea42f in __GI___wait4 (pid=161186, stat_loc=0x0, options=0, usage=0x0) at ../sysdeps/unix/sysv/linux/wait4.c:30
30	in ../sysdeps/unix/sysv/linux/wait4.c
#1  0x00007f1df215b40f in ggml_print_backtrace () from /home/supermicro/ai/llama.cpp.main/bin/libggml-base.so.0
#2  0x00007f1df215b582 in ggml_abort () from /home/supermicro/ai/llama.cpp.main/bin/libggml-base.so.0
#3  0x00007f1de13a3203 in ggml_cuda_error(char const*, char const*, char const*, int, char const*) () from /home/supermicro/ai/llama.cpp.main/bin/libggml-cuda.so.0
#4  0x00007f1de13ab94b in ggml_cuda_op_mul_mat_cublas(ggml_backend_cuda_context&, ggml_tensor const*, ggml_tensor const*, ggml_tensor*, char const*, float const*, char const*, float*, long, long, long, long, CUstream_st*) () from /home/supermicro/ai/llama.cpp.main/bin/libggml-cuda.so.0
#5  0x00007f1de13b12e2 in ggml_cuda_op_mul_mat(ggml_backend_cuda_context&, ggml_tensor const*, ggml_tensor const*, ggml_tensor*, void (*)(ggml_backend_cuda_context&, ggml_tensor const*, ggml_tensor const*, ggml_tensor*, char const*, float const*, char const*, float*, long, long, long, long, CUstream_st*), void (*)(float const*, int const*, void*, ggml_type, long, long, long, long, long, long, long, long, CUstream_st*)) () from /home/supermicro/ai/llama.cpp.main/bin/libggml-cuda.so.0
#6  0x00007f1de13b4f36 in ggml_cuda_compute_forward(ggml_backend_cuda_context&, ggml_tensor*) () from /home/supermicro/ai/llama.cpp.main/bin/libggml-cuda.so.0
#7  0x00007f1de13ba351 in ggml_cuda_graph_evaluate_and_capture(ggml_backend_cuda_context*, ggml_cgraph*, bool, bool, void const*) () from /home/supermicro/ai/llama.cpp.main/bin/libggml-cuda.so.0
#8  0x00007f1de13bc182 in ggml_backend_cuda_graph_compute(ggml_backend*, ggml_cgraph*) () from /home/supermicro/ai/llama.cpp.main/bin/libggml-cuda.so.0
#9  0x00007f1df218408b in ggml_backend_meta_graph_compute(ggml_backend*, ggml_cgraph*) () from /home/supermicro/ai/llama.cpp.main/bin/libggml-base.so.0
#10 0x00007f1df2179a51 in ggml_backend_sched_graph_compute_async () from /home/supermicro/ai/llama.cpp.main/bin/libggml-base.so.0
#11 0x00007f1df229cee2 in llama_context::graph_compute(ggml_cgraph*, bool) () from /home/supermicro/ai/llama.cpp.main/bin/libllama.so.0
#12 0x00007f1df229f50f in llama_context::process_ubatch(llama_ubatch const&, llm_graph_type, llama_memory_context_i*, ggml_status&) () from /home/supermicro/ai/llama.cpp.main/bin/libllama.so.0
#13 0x00007f1df22a7b29 in llama_context::decode(llama_batch const&) () from /home/supermicro/ai/llama.cpp.main/bin/libllama.so.0
#14 0x00007f1df22a95ec in llama_decode () from /home/supermicro/ai/llama.cpp.main/bin/libllama.so.0
#15 0x000055d56d435aba in perplexity(llama_context*, common_params const&, int) ()
#16 0x000055d56d425a9a in main ()
[Inferior 1 (process 136333) detached]
Aborted (core dumped)

A sweep bench up to 32k is fine and other kinds of inference appear to work. It's just perplexity. My other NCCL programs don't have this problem despite pushing more watts, etc. Trying to determine if it's a me problem or a code problem.