cuda : improve cuda pool efficiency using virtual memory

[slaren]

Reserves a large amount of contiguous virtual memory space, and maps new allocations into it when the pool needs to grow. In this way, no memory is wasted. The pool is effectively a stack now, so all deallocations must be done in reverse order of the allocations.

An example of probably one of the worst cases before, partial offloading with increasing prompt sizes ./llama-bench -n 0 -p 8,16,32,64,128,256,512 -ngl 30 -m models/7B/ggml-model-Q4_0.gguf, the pool size dropped from 710 MB to 460 MB.

[JohannesGaessler]

In my test implementation for peer access I was able to simply use CUDA_CHECK for the CUresult. I don't know whether that works universally though.

[slaren]

It didn't work for me, it is a different type and requires a different function to get the error description, cuGetErrorString.

[slaren]

The Makefile will probably need to be updated to add the possible paths for the CUDA driver library, the WSL path is a bit unconventional.

[ggerganov]

Tested on Ubuntu GTX 1660 and V100

The pool is effectively a stack now, so all deallocations must be done in reverse order of the allocations.

Should we RAII the allocations with a simple struct to guarantee this?

[slaren]

Should we RAII the allocations with a simple struct to guarantee this?

It could help in some cases, but many of the allocations are conditional, so we couldn't allocate the memory in the constructor, it would need to be done in a separate function, so it wouldn't guarantee that the order of the allocations is the same as the order of the declarations. It still is probably worth it just to remove the manual calls to pool_free, though.

[JohannesGaessler]

This PR breaks HIP compilation.

    ~/Pro/llama.cpp    sl/cuda-virt-pool *4 ?43  make clean && time make LLAMA_HIPBLAS=1 libllama.so quantize main perplexity llama-bench
I llama.cpp build info: 
I UNAME_S:   Linux
I UNAME_P:   unknown
I UNAME_M:   x86_64
I CFLAGS:    -I. -Icommon -D_XOPEN_SOURCE=600 -D_GNU_SOURCE -DNDEBUG  -std=c11   -fPIC -O3 -Wall -Wextra -Wpedantic -Wcast-qual -Wno-unused-function -Wshadow -Wstrict-prototypes -Wpointer-arith -Wmissing-prototypes -Werror=implicit-int -Werror=implicit-function-declaration -pthread -march=native -mtune=native -Wdouble-promotion 
I CXXFLAGS:  -I. -Icommon -D_XOPEN_SOURCE=600 -D_GNU_SOURCE -DNDEBUG  -std=c++11 -fPIC -O3 -Wall -Wextra -Wpedantic -Wcast-qual -Wno-unused-function -Wmissing-declarations -Wmissing-noreturn -pthread  -march=native -mtune=native -Wno-array-bounds -Wno-format-truncation -Wextra-semi
I NVCCFLAGS:  
I LDFLAGS:    
I CC:        cc (GCC) 13.2.1 20230801
I CXX:       g++ (GCC) 13.2.1 20230801

rm -vrf *.o tests/*.o *.so *.dll benchmark-matmult common/build-info.cpp *.dot *.gcno tests/*.gcno *.gcda tests/*.gcda *.gcov tests/*.gcov lcov-report gcovr-report main quantize quantize-stats perplexity embedding vdot q8dot train-text-from-scratch convert-llama2c-to-ggml simple batched batched-bench save-load-state server gguf llama-bench libllava.a llava-cli baby-llama beam-search speculative infill tokenize benchmark-matmult parallel finetune export-lora lookahead lookup tests/test-c.o metal tests/test-llama-grammar tests/test-grammar-parser tests/test-double-float tests/test-grad0 tests/test-opt tests/test-quantize-fns tests/test-quantize-perf tests/test-sampling tests/test-tokenizer-0-llama tests/test-tokenizer-0-falcon tests/test-tokenizer-1-llama tests/test-tokenizer-1-bpe tests/test-rope tests/test-backend-ops
removed 'build-info.o'
removed 'common.o'
removed 'console.o'
removed 'ggml-alloc.o'
removed 'ggml-backend.o'
removed 'ggml.o'
removed 'ggml-quants.o'
removed 'grammar-parser.o'
removed 'llama.o'
removed 'sampling.o'
removed 'common/build-info.cpp'
I llama.cpp build info: 
I UNAME_S:   Linux
I UNAME_P:   unknown
I UNAME_M:   x86_64
I CFLAGS:    -I. -Icommon -D_XOPEN_SOURCE=600 -D_GNU_SOURCE -DNDEBUG -DGGML_USE_HIPBLAS -DGGML_USE_CUBLAS  -std=c11   -fPIC -O3 -Wall -Wextra -Wpedantic -Wcast-qual -Wno-unused-function -Wshadow -Wstrict-prototypes -Wpointer-arith -Wmissing-prototypes -Werror=implicit-int -Werror=implicit-function-declaration -pthread -march=native -mtune=native -Wdouble-promotion 
I CXXFLAGS:  -I. -Icommon -D_XOPEN_SOURCE=600 -D_GNU_SOURCE -DNDEBUG -DGGML_USE_HIPBLAS -DGGML_USE_CUBLAS  -std=c++11 -fPIC -O3 -Wall -Wextra -Wpedantic -Wcast-qual -Wno-unused-function -Wmissing-declarations -Wmissing-noreturn -pthread  -march=native -mtune=native -Wno-array-bounds -Wno-format-truncation -Wextra-semi
I NVCCFLAGS:  
I LDFLAGS:   -L/opt/rocm/lib -Wl,-rpath=/opt/rocm/lib -lhipblas -lamdhip64 -lrocblas 
I CC:        cc (GCC) 13.2.1 20230801
I CXX:       g++ (GCC) 13.2.1 20230801

g++ -I. -Icommon -D_XOPEN_SOURCE=600 -D_GNU_SOURCE -DNDEBUG -DGGML_USE_HIPBLAS -DGGML_USE_CUBLAS  -std=c++11 -fPIC -O3 -Wall -Wextra -Wpedantic -Wcast-qual -Wno-unused-function -Wmissing-declarations -Wmissing-noreturn -pthread  -march=native -mtune=native -Wno-array-bounds -Wno-format-truncation -Wextra-semi -c llama.cpp -o llama.o
cc  -I. -Icommon -D_XOPEN_SOURCE=600 -D_GNU_SOURCE -DNDEBUG -DGGML_USE_HIPBLAS -DGGML_USE_CUBLAS  -std=c11   -fPIC -O3 -Wall -Wextra -Wpedantic -Wcast-qual -Wno-unused-function -Wshadow -Wstrict-prototypes -Wpointer-arith -Wmissing-prototypes -Werror=implicit-int -Werror=implicit-function-declaration -pthread -march=native -mtune=native -Wdouble-promotion    -c ggml.c -o ggml.o
/opt/rocm/bin/hipcc -I. -Icommon -D_XOPEN_SOURCE=600 -D_GNU_SOURCE -DNDEBUG -DGGML_USE_HIPBLAS -DGGML_USE_CUBLAS  -std=c++11 -fPIC -O3 -Wall -Wextra -Wpedantic -Wcast-qual -Wno-unused-function -Wmissing-declarations -Wmissing-noreturn -pthread  -march=native -mtune=native -Wno-array-bounds -Wno-format-truncation -Wextra-semi --offload-arch=gfx1030 -DGGML_CUDA_DMMV_X=32 -DGGML_CUDA_MMV_Y=1 -DK_QUANTS_PER_ITERATION=2 -x hip -c -o ggml-cuda.o ggml-cuda.cu
cc  -I. -Icommon -D_XOPEN_SOURCE=600 -D_GNU_SOURCE -DNDEBUG -DGGML_USE_HIPBLAS -DGGML_USE_CUBLAS  -std=c11   -fPIC -O3 -Wall -Wextra -Wpedantic -Wcast-qual -Wno-unused-function -Wshadow -Wstrict-prototypes -Wpointer-arith -Wmissing-prototypes -Werror=implicit-int -Werror=implicit-function-declaration -pthread -march=native -mtune=native -Wdouble-promotion    -c ggml-alloc.c -o ggml-alloc.o
cc  -I. -Icommon -D_XOPEN_SOURCE=600 -D_GNU_SOURCE -DNDEBUG -DGGML_USE_HIPBLAS -DGGML_USE_CUBLAS  -std=c11   -fPIC -O3 -Wall -Wextra -Wpedantic -Wcast-qual -Wno-unused-function -Wshadow -Wstrict-prototypes -Wpointer-arith -Wmissing-prototypes -Werror=implicit-int -Werror=implicit-function-declaration -pthread -march=native -mtune=native -Wdouble-promotion    -c ggml-backend.c -o ggml-backend.o
cc -I. -Icommon -D_XOPEN_SOURCE=600 -D_GNU_SOURCE -DNDEBUG -DGGML_USE_HIPBLAS -DGGML_USE_CUBLAS  -std=c11   -fPIC -O3 -Wall -Wextra -Wpedantic -Wcast-qual -Wno-unused-function -Wshadow -Wstrict-prototypes -Wpointer-arith -Wmissing-prototypes -Werror=implicit-int -Werror=implicit-function-declaration -pthread -march=native -mtune=native -Wdouble-promotion     -c ggml-quants.c -o ggml-quants.o
g++ -I. -Icommon -D_XOPEN_SOURCE=600 -D_GNU_SOURCE -DNDEBUG -DGGML_USE_HIPBLAS -DGGML_USE_CUBLAS  -std=c++11 -fPIC -O3 -Wall -Wextra -Wpedantic -Wcast-qual -Wno-unused-function -Wmissing-declarations -Wmissing-noreturn -pthread  -march=native -mtune=native -Wno-array-bounds -Wno-format-truncation -Wextra-semi -c common/common.cpp -o common.o
g++ -I. -Icommon -D_XOPEN_SOURCE=600 -D_GNU_SOURCE -DNDEBUG -DGGML_USE_HIPBLAS -DGGML_USE_CUBLAS  -std=c++11 -fPIC -O3 -Wall -Wextra -Wpedantic -Wcast-qual -Wno-unused-function -Wmissing-declarations -Wmissing-noreturn -pthread  -march=native -mtune=native -Wno-array-bounds -Wno-format-truncation -Wextra-semi -c common/sampling.cpp -o sampling.o
g++ -I. -Icommon -D_XOPEN_SOURCE=600 -D_GNU_SOURCE -DNDEBUG -DGGML_USE_HIPBLAS -DGGML_USE_CUBLAS  -std=c++11 -fPIC -O3 -Wall -Wextra -Wpedantic -Wcast-qual -Wno-unused-function -Wmissing-declarations -Wmissing-noreturn -pthread  -march=native -mtune=native -Wno-array-bounds -Wno-format-truncation -Wextra-semi -c common/grammar-parser.cpp -o grammar-parser.o
g++ -I. -Icommon -D_XOPEN_SOURCE=600 -D_GNU_SOURCE -DNDEBUG -DGGML_USE_HIPBLAS -DGGML_USE_CUBLAS  -std=c++11 -fPIC -O3 -Wall -Wextra -Wpedantic -Wcast-qual -Wno-unused-function -Wmissing-declarations -Wmissing-noreturn -pthread  -march=native -mtune=native -Wno-array-bounds -Wno-format-truncation -Wextra-semi -c common/console.cpp -o console.o
g++ -I. -Icommon -D_XOPEN_SOURCE=600 -D_GNU_SOURCE -DNDEBUG -DGGML_USE_HIPBLAS -DGGML_USE_CUBLAS  -std=c++11 -fPIC -O3 -Wall -Wextra -Wpedantic -Wcast-qual -Wno-unused-function -Wmissing-declarations -Wmissing-noreturn -pthread  -march=native -mtune=native -Wno-array-bounds -Wno-format-truncation -Wextra-semi -c common/build-info.cpp -o build-info.o
warning: unknown warning option '-Wno-format-truncation' [-Wunknown-warning-option]
ggml-cuda.cu:553:1: warning: function declared 'noreturn' should not return [-Winvalid-noreturn]
}
^
ggml-cuda.cu:580:63: warning: unused parameter 'a' [-Wunused-parameter]
static __device__ __forceinline__ float op_repeat(const float a, const float b) {
                                                              ^
ggml-cuda.cu:792:20: warning: comparison of integers of different signs: 'R' (aka 'unsigned int') and 'const int' [-Wsign-compare]
    if (blockIdx.z < ne02) { // src0
        ~~~~~~~~~~ ^ ~~~~
ggml-cuda.cu:838:56: warning: comparison of integers of different signs: 'R' (aka 'unsigned int') and 'const int' [-Wsign-compare]
    if (nidx < ne00 && blockIdx.y < ne01 && blockIdx.z < ne02) {
                                            ~~~~~~~~~~ ^ ~~~~
ggml-cuda.cu:838:35: warning: comparison of integers of different signs: 'R' (aka 'unsigned int') and 'const int' [-Wsign-compare]
    if (nidx < ne00 && blockIdx.y < ne01 && blockIdx.z < ne02) {
                       ~~~~~~~~~~ ^ ~~~~
ggml-cuda.cu:6805:13: error: use of undeclared identifier 'g_device_vmm'
            g_device_vmm[id] = !!deviceSupportsVmm;
            ^
ggml-cuda.cu:6809:124: error: use of undeclared identifier 'g_device_vmm'
            fprintf(stderr, "  Device %d: %s, compute capability %d.%d, VMM: %s\n", id, prop.name, prop.major, prop.minor, g_device_vmm[id] ? "yes" : "no");
                                                                                                                           ^
ggml-cuda.cu:8012:19: warning: unused variable 'nrows0' [-Wunused-variable]
    const int64_t nrows0 = ggml_nrows(src0);
                  ^
7 warnings and 2 errors generated when compiling for gfx1030.
make: *** [Makefile:486: ggml-cuda.o] Error 1
make: *** Waiting for unfinished jobs....
/usr/bin/make --jobs=32 LLAMA_HIPBLAS=1 libllama.so quantize main perplexity   48,42s user 1,32s system 261% cpu 18,988 total

More generally, I would like it if the correct order of allocations/deallocations was actually enforced. You could make it so that instead of just a size ggml_cuda_pool_malloc returns a struct that contains both a size and an index. A static variable in ggml_cuda_pool_malloc keeps track of the next index expected for deallocation. If it does not align an error is raised. (This PR should be fine even without such an enforcement mechanism though.)

[slaren]

I would like it if the correct order of allocations/deallocations was actually enforced

This is already done, pool_free checks that the address being released is at the end of the pool.

[slaren]

Added the RAII struct and changed most of the code to use it, but this is a very error-prone process so another review would be good.

[slaren]

I think it was a mistake opting for a list of macros instead of hipify.

[slaren]

ggml_cuda_op_mul_mat is still directly using ggml_cuda_pool_malloc but I assume you're aware of that.

Using it in ggml_cuda_op_mul_mat is not so simple because just replacing the pointers and sizes with an array of cuda_pool_alloc wouldn't work, because the order of destruction wouldn't be the same as the order of allocations (reversed). It may work if the allocations are grouped per-device in a struct or tuple, but it requires more changes.

I'm assuming that the order in which the destructors is called is the reverse order in which the constructors were called so the correct order should be guaranteed.

Yes.

[city96]

This PR seems to break LLaMA2 70B for me. I also tested 13B (F16 precision) but that one worked fine, so I'm guessing it's GQA related?

CUDA error: cudaMemcpyAsync(src1_ddq_i, src1_ddq_i_source, src1_ncols*src1_padded_col_size*q8_1_ts/q8_1_bs, cudaMemcpyDeviceToDevice, stream): invalid argument
  in function ggml_cuda_op_mul_mat at ggml-cuda.cu:8196
GGML_ASSERT: ggml-cuda.cu:239: !"CUDA error"

(Build command is make clean && make -j16 LLAMA_CUBLAS=1 LLAMA_CUDA_FORCE_MMQ=1 LLAMA_CUDA_PEER_MAX_BATCH_SIZE=0. CUDA 11.7 on ubuntu 22.04.3. Tesla V100S 32GB + P40 24GB in a dual CPU server, so the cards are on different NUMA nodes. LMK if you need any more info or if I should open a bug report about this instead.)

Server on (#708e179)

./server --host 192.168.4.6 -c 4096 -t 1 -ngl 128 -b 128 -m /mnt/models/llama2-70b-q6_K.bin -ts 15.6,12
ggml_init_cublas: GGML_CUDA_FORCE_MMQ:   yes
ggml_init_cublas: CUDA_USE_TENSOR_CORES: no
ggml_init_cublas: found 2 CUDA devices:
  Device 0: Tesla V100S-PCIE-32GB, compute capability 7.0
  Device 1: Tesla P40, compute capability 6.1
{"timestamp":1703433377,"level":"INFO","function":"main","line":2711,"message":"build info","build":1696,"commit":"708e179"}
{"timestamp":1703433377,"level":"INFO","function":"main","line":2714,"message":"system info","n_threads":1,"n_threads_batch":-1,"total_threads":32,"system_info":"AVX = 1 | AVX2 = 0 | AVX512 = 0 | AVX512_VBMI = 0 | AVX512_VNNI = 0 | FMA = 0 | NEON = 0 | ARM_FMA = 0 | F16C = 1 | FP16_VA = 0 | WASM_SIMD = 0 | BLAS = 1 | SSE3 = 1 | SSSE3 = 1 | VSX = 0 | "}
llama_model_loader: loaded meta data with 16 key-value pairs and 723 tensors from /mnt/models/llama2-70b-q6_K.bin (version GGUF V3 (latest))
llama_model_loader: Dumping metadata keys/values. Note: KV overrides do not apply in this output.
llama_model_loader: - kv   0:                       general.architecture str              = llama
llama_model_loader: - kv   1:                               general.name str              = LLaMA v2
llama_model_loader: - kv   2:                       llama.context_length u32              = 4096
llama_model_loader: - kv   3:                     llama.embedding_length u32              = 8192
llama_model_loader: - kv   4:                          llama.block_count u32              = 80
llama_model_loader: - kv   5:                  llama.feed_forward_length u32              = 28672
llama_model_loader: - kv   6:                 llama.rope.dimension_count u32              = 128
llama_model_loader: - kv   7:                 llama.attention.head_count u32              = 64
llama_model_loader: - kv   8:              llama.attention.head_count_kv u32              = 8
llama_model_loader: - kv   9:     llama.attention.layer_norm_rms_epsilon f32              = 0.000010
llama_model_loader: - kv  10:                          general.file_type u32              = 18
llama_model_loader: - kv  11:                       tokenizer.ggml.model str              = llama
llama_model_loader: - kv  12:                      tokenizer.ggml.tokens arr[str,32000]   = ["<unk>", "<s>", "</s>", "<0x00>", "<...
llama_model_loader: - kv  13:                      tokenizer.ggml.scores arr[f32,32000]   = [0.000000, 0.000000, 0.000000, 0.0000...
llama_model_loader: - kv  14:                  tokenizer.ggml.token_type arr[i32,32000]   = [2, 3, 3, 6, 6, 6, 6, 6, 6, 6, 6, 6, ...
llama_model_loader: - kv  15:               general.quantization_version u32              = 2
llama_model_loader: - type  f32:  161 tensors
llama_model_loader: - type q6_K:  562 tensors
llm_load_vocab: special tokens definition check successful ( 259/32000 ).
llm_load_print_meta: format           = GGUF V3 (latest)
llm_load_print_meta: arch             = llama
llm_load_print_meta: vocab type       = SPM
llm_load_print_meta: n_vocab          = 32000
llm_load_print_meta: n_merges         = 0
llm_load_print_meta: n_ctx_train      = 4096
llm_load_print_meta: n_embd           = 8192
llm_load_print_meta: n_head           = 64
llm_load_print_meta: n_head_kv        = 8
llm_load_print_meta: n_layer          = 80
llm_load_print_meta: n_rot            = 128
llm_load_print_meta: n_gqa            = 8
llm_load_print_meta: f_norm_eps       = 0.0e+00
llm_load_print_meta: f_norm_rms_eps   = 1.0e-05
llm_load_print_meta: f_clamp_kqv      = 0.0e+00
llm_load_print_meta: f_max_alibi_bias = 0.0e+00
llm_load_print_meta: n_ff             = 28672
llm_load_print_meta: n_expert         = 0
llm_load_print_meta: n_expert_used    = 0
llm_load_print_meta: rope scaling     = linear
llm_load_print_meta: freq_base_train  = 10000.0
llm_load_print_meta: freq_scale_train = 1
llm_load_print_meta: n_yarn_orig_ctx  = 4096
llm_load_print_meta: rope_finetuned   = unknown
llm_load_print_meta: model type       = 70B
llm_load_print_meta: model ftype      = Q6_K
llm_load_print_meta: model params     = 68.98 B
llm_load_print_meta: model size       = 52.70 GiB (6.56 BPW) 
llm_load_print_meta: general.name     = LLaMA v2
llm_load_print_meta: BOS token        = 1 '<s>'
llm_load_print_meta: EOS token        = 2 '</s>'
llm_load_print_meta: UNK token        = 0 '<unk>'
llm_load_print_meta: LF token         = 13 '<0x0A>'
llm_load_tensors: ggml ctx size       =    0.28 MiB
llm_load_tensors: using CUDA for GPU acceleration
llm_load_tensors: system memory used  =  205.35 MiB
llm_load_tensors: VRAM used           = 53760.11 MiB
llm_load_tensors: offloading 80 repeating layers to GPU
llm_load_tensors: offloading non-repeating layers to GPU
llm_load_tensors: offloaded 81/81 layers to GPU
....................................................................................................
llama_new_context_with_model: n_ctx      = 4096
llama_new_context_with_model: freq_base  = 10000.0
llama_new_context_with_model: freq_scale = 1
llama_kv_cache_init: VRAM kv self = 1280.00 MB
llama_new_context_with_model: KV self size  = 1280.00 MiB, K (f16):  640.00 MiB, V (f16):  640.00 MiB
llama_build_graph: non-view tensors processed: 1684/1684
llama_new_context_with_model: compute buffer total size = 145.19 MiB
llama_new_context_with_model: VRAM scratch buffer: 142.00 MiB
llama_new_context_with_model: total VRAM used: 55182.11 MiB (model: 53760.11 MiB, context: 1422.00 MiB)
Available slots:
 -> Slot 0 - max context: 4096

llama server listening at http://192.168.4.6:8080

{"timestamp":1703433421,"level":"INFO","function":"main","line":3141,"message":"HTTP server listening","port":"8080","hostname":"192.168.4.6"}
all slots are idle and system prompt is empty, clear the KV cache
{"timestamp":1703433443,"level":"INFO","function":"log_server_request","line":2652,"message":"request","remote_addr":"192.168.1.156","remote_port":45466,"status":200,"method":"GET","path":"/","params":{}}
{"timestamp":1703433443,"level":"INFO","function":"log_server_request","line":2652,"message":"request","remote_addr":"192.168.1.156","remote_port":45466,"status":200,"method":"GET","path":"/index.js","params":{}}
{"timestamp":1703433443,"level":"INFO","function":"log_server_request","line":2652,"message":"request","remote_addr":"192.168.1.156","remote_port":45470,"status":200,"method":"GET","path":"/completion.js","params":{}}
{"timestamp":1703433443,"level":"INFO","function":"log_server_request","line":2652,"message":"request","remote_addr":"192.168.1.156","remote_port":45482,"status":200,"method":"GET","path":"/json-schema-to-grammar.mjs","params":{}}
slot 0 is processing [task id: 0]
slot 0 : in cache: 0 tokens | to process: 63 tokens
slot 0 : kv cache rm - [0, end)

print_timings: prompt eval time =    1041.54 ms /    63 tokens (   16.53 ms per token,    60.49 tokens per second)
print_timings:        eval time =    6993.77 ms /    41 runs   (  170.58 ms per token,     5.86 tokens per second)
print_timings:       total time =    8035.31 ms
slot 0 released (105 tokens in cache)
{"timestamp":1703433463,"level":"INFO","function":"log_server_request","line":2652,"message":"request","remote_addr":"192.168.1.156","remote_port":43100,"status":200,"method":"POST","path":"/completion","params":{}}
slot 0 released (105 tokens in cache)

Server on (#5bf3953)

./server --host 192.168.4.6 -c 4096 -t 1 -ngl 128 -b 128 -m /mnt/models/llama2-70b-q6_K.bin -ts 15.6,12
ggml_init_cublas: GGML_CUDA_FORCE_MMQ:   yes
ggml_init_cublas: CUDA_USE_TENSOR_CORES: no
ggml_init_cublas: found 2 CUDA devices:
  Device 0: Tesla V100S-PCIE-32GB, compute capability 7.0, VMM: yes
  Device 1: Tesla P40, compute capability 6.1, VMM: yes
{"timestamp":1703433717,"level":"INFO","function":"main","line":2711,"message":"build info","build":1697,"commit":"5bf3953"}
{"timestamp":1703433717,"level":"INFO","function":"main","line":2714,"message":"system info","n_threads":1,"n_threads_batch":-1,"total_threads":32,"system_info":"AVX = 1 | AVX2 = 0 | AVX512 = 0 | AVX512_VBMI = 0 | AVX512_VNNI = 0 | FMA = 0 | NEON = 0 | ARM_FMA = 0 | F16C = 1 | FP16_VA = 0 | WASM_SIMD = 0 | BLAS = 1 | SSE3 = 1 | SSSE3 = 1 | VSX = 0 | "}
llama_model_loader: loaded meta data with 16 key-value pairs and 723 tensors from /mnt/models/llama2-70b-q6_K.bin (version GGUF V3 (latest))
llama_model_loader: Dumping metadata keys/values. Note: KV overrides do not apply in this output.
llama_model_loader: - kv   0:                       general.architecture str              = llama
llama_model_loader: - kv   1:                               general.name str              = LLaMA v2
llama_model_loader: - kv   2:                       llama.context_length u32              = 4096
llama_model_loader: - kv   3:                     llama.embedding_length u32              = 8192
llama_model_loader: - kv   4:                          llama.block_count u32              = 80
llama_model_loader: - kv   5:                  llama.feed_forward_length u32              = 28672
llama_model_loader: - kv   6:                 llama.rope.dimension_count u32              = 128
llama_model_loader: - kv   7:                 llama.attention.head_count u32              = 64
llama_model_loader: - kv   8:              llama.attention.head_count_kv u32              = 8
llama_model_loader: - kv   9:     llama.attention.layer_norm_rms_epsilon f32              = 0.000010
llama_model_loader: - kv  10:                          general.file_type u32              = 18
llama_model_loader: - kv  11:                       tokenizer.ggml.model str              = llama
llama_model_loader: - kv  12:                      tokenizer.ggml.tokens arr[str,32000]   = ["<unk>", "<s>", "</s>", "<0x00>", "<...
llama_model_loader: - kv  13:                      tokenizer.ggml.scores arr[f32,32000]   = [0.000000, 0.000000, 0.000000, 0.0000...
llama_model_loader: - kv  14:                  tokenizer.ggml.token_type arr[i32,32000]   = [2, 3, 3, 6, 6, 6, 6, 6, 6, 6, 6, 6, ...
llama_model_loader: - kv  15:               general.quantization_version u32              = 2
llama_model_loader: - type  f32:  161 tensors
llama_model_loader: - type q6_K:  562 tensors
llm_load_vocab: special tokens definition check successful ( 259/32000 ).
llm_load_print_meta: format           = GGUF V3 (latest)
llm_load_print_meta: arch             = llama
llm_load_print_meta: vocab type       = SPM
llm_load_print_meta: n_vocab          = 32000
llm_load_print_meta: n_merges         = 0
llm_load_print_meta: n_ctx_train      = 4096
llm_load_print_meta: n_embd           = 8192
llm_load_print_meta: n_head           = 64
llm_load_print_meta: n_head_kv        = 8
llm_load_print_meta: n_layer          = 80
llm_load_print_meta: n_rot            = 128
llm_load_print_meta: n_gqa            = 8
llm_load_print_meta: f_norm_eps       = 0.0e+00
llm_load_print_meta: f_norm_rms_eps   = 1.0e-05
llm_load_print_meta: f_clamp_kqv      = 0.0e+00
llm_load_print_meta: f_max_alibi_bias = 0.0e+00
llm_load_print_meta: n_ff             = 28672
llm_load_print_meta: n_expert         = 0
llm_load_print_meta: n_expert_used    = 0
llm_load_print_meta: rope scaling     = linear
llm_load_print_meta: freq_base_train  = 10000.0
llm_load_print_meta: freq_scale_train = 1
llm_load_print_meta: n_yarn_orig_ctx  = 4096
llm_load_print_meta: rope_finetuned   = unknown
llm_load_print_meta: model type       = 70B
llm_load_print_meta: model ftype      = Q6_K
llm_load_print_meta: model params     = 68.98 B
llm_load_print_meta: model size       = 52.70 GiB (6.56 BPW) 
llm_load_print_meta: general.name     = LLaMA v2
llm_load_print_meta: BOS token        = 1 '<s>'
llm_load_print_meta: EOS token        = 2 '</s>'
llm_load_print_meta: UNK token        = 0 '<unk>'
llm_load_print_meta: LF token         = 13 '<0x0A>'
llm_load_tensors: ggml ctx size       =    0.28 MiB
llm_load_tensors: using CUDA for GPU acceleration
llm_load_tensors: system memory used  =  205.35 MiB
llm_load_tensors: VRAM used           = 53760.11 MiB
llm_load_tensors: offloading 80 repeating layers to GPU
llm_load_tensors: offloading non-repeating layers to GPU
llm_load_tensors: offloaded 81/81 layers to GPU
....................................................................................................
llama_new_context_with_model: n_ctx      = 4096
llama_new_context_with_model: freq_base  = 10000.0
llama_new_context_with_model: freq_scale = 1
llama_kv_cache_init: VRAM kv self = 1280.00 MB
llama_new_context_with_model: KV self size  = 1280.00 MiB, K (f16):  640.00 MiB, V (f16):  640.00 MiB
llama_build_graph: non-view tensors processed: 1684/1684
llama_new_context_with_model: compute buffer total size = 145.19 MiB
llama_new_context_with_model: VRAM scratch buffer: 142.00 MiB
llama_new_context_with_model: total VRAM used: 55182.11 MiB (model: 53760.11 MiB, context: 1422.00 MiB)
CUDA error: cudaMemcpyAsync(src1_ddq_i, src1_ddq_i_source, src1_ncols*src1_padded_col_size*q8_1_ts/q8_1_bs, cudaMemcpyDeviceToDevice, stream): invalid argument
  in function ggml_cuda_op_mul_mat at ggml-cuda.cu:8196
GGML_ASSERT: ggml-cuda.cu:239: !"CUDA error"
Aborted (core dumped)

[slaren]

@city96 can you test this change?

diff --git a/ggml-cuda.cu b/ggml-cuda.cu
index ac3b3c14..3f7e6d1a 100644
--- a/ggml-cuda.cu
+++ b/ggml-cuda.cu
@@ -8193,8 +8193,8 @@ static void ggml_cuda_op_mul_mat(
                     if (id != g_main_device) {
                         if (convert_src1_to_q8_1) {
                             char * src1_ddq_i_source = src1_ddq[g_main_device] + src1_ddq_i_offset;
-                            CUDA_CHECK(cudaMemcpyAsync(src1_ddq_i, src1_ddq_i_source, src1_ncols*src1_padded_col_size*q8_1_ts/q8_1_bs,
-                                                    cudaMemcpyDeviceToDevice, stream));
+                            CUDA_CHECK(cudaMemcpyPeerAsync(src1_ddq_i, id, src1_ddq_i_source, g_main_device,
+                                                            src1_ncols*src1_padded_col_size*q8_1_ts/q8_1_bs, stream));
                         } else {
                             float * src1_ddf_i_source = (float *) src1_extra->data_device[g_main_device];
                             src1_ddf_i_source += (i0*ne11 + src1_col_0) * ne10;

[city96]

@slaren That seems to fix it (applied directly to 5bf3953)

[slaren]

I think it is failing in a memcpy between devices with quantized models. Changing it to use cudaMemcpyPeerAsync fixes the issue. However the cudaMemcpyAsync still works when not using mmq:

                // copy src0, src1 to device if necessary
                if (src1->backend == GGML_BACKEND_GPU && src1_is_contiguous) {
                    if (id != g_main_device) {
                        if (convert_src1_to_q8_1) {
                            char * src1_ddq_i_source = src1_ddq[g_main_device] + src1_ddq_i_offset;
                            CUDA_CHECK(cudaMemcpyPeerAsync(src1_ddq_i, id, src1_ddq_i_source, g_main_device,
                                                            src1_ncols*src1_padded_col_size*q8_1_ts/q8_1_bs, stream));
                        } else {
                            float * src1_ddf_i_source = (float *) src1_extra->data_device[g_main_device];
                            src1_ddf_i_source += (i0*ne11 + src1_col_0) * ne10;
                            CUDA_CHECK(cudaMemcpyAsync(src1_ddf_i, src1_ddf_i_source, src1_ncols*ne10*sizeof(float),
                                                    cudaMemcpyDeviceToDevice, stream));
                        }
                    }

I think the difference is that the src1 source is only allocated on the pool when using mmq. So I think that the issue happens only when copying between pools of different devices. Does this sound right? Are there any other copies between different devices that should be changed? @JohannesGaessler

Fix: #4620

@city96 thanks for testing.

[JohannesGaessler]

@city96 can you also test the version I just pushed to #4602 ? It may be the same issue that I found where (I think) peer access status changes during a data transfer.

[city96]

@JohannesGaessler

Both llama 70B and mixtral currently fail with the following error:

CUDA error: cudaMemcpyAsync(src1_ddq_i, src1_ddq_i_source, src1_ncols*src1_padded_col_size*q8_1_ts/q8_1_bs, cudaMemcpyDeviceToDevice, stream): invalid argument
  in function ggml_cuda_op_mul_mat at ggml-cuda.cu:8200
GGML_ASSERT: ggml-cuda.cu:240: !"CUDA error"

[cuda-fix-peer-sync commit e692c2d]

I have peer access disabled with LLAMA_CUDA_PEER_MAX_BATCH_SIZE=0 due to the 2 GPUs being connected to different CPUs (nvidia-smi topo -m reports SYS [Connection traversing PCIe as well as the SMP interconnect between NUMA nodes (e.g., QPI/UPI)]) - if that matters.

Happy to test anything else!

[SomeoneSerge]

By the way LLAMA_EXTRA_LIBS ends up in target_link_libraries(... PUBLIC ...) which probably isn't what was intended here?