speedup ROCm AMD Unified Memory Architecture

[Djip007]

GGML_HIP_UMA allow to use hipMallocManaged tu use UMA on AMD/HIP GPU.

I have a Ryzen 7940HS an made some test. Using UMA allow use much more memorie than reserved VRAM on the igpu and it is nice. It allow some speed up over CPU. But by default it use "Fine-grained" memorie that is "slow". If we can use "Coarse-grained" we can have more speed.

What I did for test is replace:

#define cudaMalloc hipMallocManaged
#define cudaMallocHost(ptr, size) hipHostMalloc(ptr, size)
#else
#define cudaMalloc hipMalloc
#define cudaMallocHost(ptr, size) hipHostMalloc(ptr, size, hipHostMallocDefault)
#endif

by:

#ifdef GGML_HIP_UMA
template<typename T>
static inline auto gpuAlloc(T** adr, size_t size) {
    auto res = hipMallocManaged(adr, size);
    if (res == hipSuccess) {
        return hipMemAdvise (*adr, size, hipMemAdviseSetCoarseGrain, 0);
    }
    return res;
}
#define cudaMalloc gpuAlloc
#define cudaMallocHost(ptr, size) hipHostMalloc(ptr, size)
#else
#define cudaMalloc hipMalloc
#define cudaMallocHost(ptr, size) hipHostMalloc(ptr, size, hipHostMallocDefault)
#endif

I test with llamafile 8.4 source code (sorry I made some other test with it like add option to add option with UMA...)
With "mistral-7b-instruct-v0.2.Q8_0.gguf" on a "Framwork 16 with 7940HS / no dGPU + fedora 40 (rocm6)" I have this result:

// - zen4 x 8 => CPU
llama_print_timings:        load time =    1383.08 ms
llama_print_timings:      sample time =       4.02 ms /   534 runs   (    0.01 ms per token, 132835.82 tokens per second)
llama_print_timings: prompt eval time =   29762.54 ms /  1466 tokens (   20.30 ms per token,    49.26 tokens per second)
llama_print_timings:        eval time =   82369.69 ms /   533 runs   (  154.54 ms per token,     6.47 tokens per second)
llama_print_timings:       total time =  112190.80 ms /  1999 tokens

// - gfx1103 / rocblas / HSA_OVERRIDE_GFX_VERSION=11.0.1 
llama_print_timings:        load time =    5391.01 ms
llama_print_timings:      sample time =       2.84 ms /   406 runs   (    0.01 ms per token, 143058.49 tokens per second)
llama_print_timings: prompt eval time =   14886.53 ms /  1466 tokens (   10.15 ms per token,    98.48 tokens per second)
llama_print_timings:        eval time =   67061.92 ms /   405 runs   (  165.58 ms per token,     6.04 tokens per second)
llama_print_timings:       total time =   82020.34 ms /  1871 tokens

// - gfx1103 / rocblas / HSA_OVERRIDE_GFX_VERSION=11.0.1 + "hipMemAdviseSetCoarseGrain"
llama_print_timings:        load time =    5470.93 ms
llama_print_timings:      sample time =       2.94 ms /   121 runs   (    0.02 ms per token, 41212.53 tokens per second)
llama_print_timings: prompt eval time =    8093.41 ms /  1466 tokens (    5.52 ms per token,   181.14 tokens per second)
llama_print_timings:        eval time =   12917.88 ms /   120 runs   (  107.65 ms per token,     9.29 tokens per second)
llama_print_timings:       total time =   21095.91 ms /  1586 tokens

As you can see I get x2 on GPU with "CoarseGrain" and x4 from CPU for prompt eval.
And event +40% for prompte processing ...

Note: to be fare I have crache from time to time on GPU... but that an other story and this gfx1103 is not fully supported (I don't know if crache is because of rocm 6.0 rebuild by Fedora or existe on other rocm ...

Ps: I need to use 'HSA_OVERRIDE_GFX_VERSION=11.0.1' env var... it is the faster than 1100 and 1102 until the 1103 is available.

ps2: sorry for my bad english ...

[Djip007]

I do not made a PR for it because as you see:

hipMemAdvise (*adr, size, hipMemAdviseSetCoarseGrain, 0);

There is a last parameter that is the "ID" of the GPU... I don't now how to get the "good" one if more than 1 GPU is possible and do not have such config to test...

https://rocm.docs.amd.com/projects/HIP/en/latest/doxygen/html/group___memory_m.html#ga5c8a3ea8a8702747588082ed39ea51bf

[slaren]

The device id is available in ggml_backend_cuda_buffer_type_alloc_buffer and ggml_cuda_pool::alloc, which is where the calls to cudaMalloc are made. It may make more sense to replace the calls to cudaMalloc there with a more generic function that takes a device id, such as cudaError_t ggml_cuda_alloc_device(void ** ptr, size_t size, int device).

[Djip007]

Thanks @slaren ! I have a look on is use more closely... a see 3 uses:

• device for ggml_cuda_pool::alloc
• buft_ctx->device for ggml_backend_cuda_buffer_type_alloc_buffer
• id for ggml_backend_cuda_split_buffer_init_tensor

Is it the "true" device ID or the id it have on "ggml_backend_cuda" ?

If it is that simple I can try to make a PR (not now it is time to go to sleep ;) )

[slaren]

It is the CUDA device id. I assume it is the same kind of device id that is used in hipMemAdvise, but I don't know that.

[Djip007]

https://rocm.docs.amd.com/projects/HIP/en/latest/doxygen/html/group___device.html#ga43c1e7f15925eeb762195ccb5e063eae

"hipSetDevice()" it say:

Many HIP APIs implicitly use the "default device" :

Any device memory subsequently allocated from this host thread (using hipMalloc) will be allocated on device.

So it make sens that it is the correct ID for hipMemAdvise in this case ...

👍

[Djip007]

may look like that...

static inline cudaError_t ggml_cuda_alloc_device(void ** ptr, size_t size, int device) {
#if defined(GGML_USE_HIPBLAS)
#if defined(GGML_HIP_UMA)
    auto res = hipMallocManaged(ptr, size);
    if (res == hipSuccess) {
        return hipMemAdvise (*ptr, size, hipMemAdviseSetCoarseGrain, device);
    }
    return res;
#else
    return hipMalloc(ptr, size);
#endif
#else
    return cudaMalloc(ptr, size);
#endif
}

or is it best practice to use CUDA_CHECK and return void?

[slaren]

Return the cudaError_t and let the caller handle it, it's not always ok to crash the application on an allocation failure.

[Djip007]

OK PR submit. Test on fc40/rocm6.0 on Ryzen 7940HS.

If any on can compare on other rocm / APU (Ryzen 6800HS / 7840U...) OS ...

My dream would be to see what it looks like with MI300A...

like report on merge request I get:

# build
make -j16 LLAMA_HIPBLAS=1 LLAMA_HIP_UMA=1 AMDGPU_TARGETS=gfx1101

# run
HSA_OVERRIDE_GFX_VERSION=11.0.1 ./main -m ~/LLM/mistral-7b-instruct-v0.2.Q8_0.gguf -ngl 999 --temp 0 -c 2048 -p "[INST] {ask resume of a long text...} [/INST]"

# before PR:
llama_print_timings:        load time =    3386,19 ms
llama_print_timings:      sample time =      12,14 ms /   609 runs   (    0,02 ms per token, 50177,14 tokens per second)
llama_print_timings: prompt eval time =   15051,81 ms /  1466 tokens (   10,27 ms per token,    97,40 tokens per second)
llama_print_timings:        eval time =  100420,23 ms /   608 runs   (  165,16 ms per token,     6,05 tokens per second)
llama_print_timings:       total time =  115547,81 ms /  2074 tokens

# after PR:
llama_print_timings:        load time =    2606,81 ms
llama_print_timings:      sample time =      12,19 ms /   609 runs   (    0,02 ms per token, 49946,69 tokens per second)
llama_print_timings: prompt eval time =    8120,08 ms /  1466 tokens (    5,54 ms per token,   180,54 tokens per second)
llama_print_timings:        eval time =   65652,40 ms /   608 runs   (  107,98 ms per token,     9,26 tokens per second)
llama_print_timings:       total time =   73841,90 ms /  2074 tokens

[Djip007]

Did some more test on same hardware.

But first an advise... use --no-mmap when use of LLAMA_HIP_UMA=1 it look to use memory more accurately, and we in all case need to load weights. Should we deactivate mmap (at least by default) when LLAMA_HIP_UMA is active on build?

some result:

//> Mixtral 7b Q4_K_M
// HSA_OVERRIDE_GFX_VERSION=11.0.1  ./main -m mistral-7b-instruct-v0.2.Q4_K_M.gguf -ngl 999 --no-mmap --temp 0 -c 2048 -p "[INST] ...
llama_print_timings:        load time =    4965,74 ms
llama_print_timings:      sample time =      10,01 ms /   387 runs   (    0,03 ms per token, 38653,62 tokens per second)
llama_print_timings: prompt eval time =    7884,96 ms /  1466 tokens (    5,38 ms per token,   185,92 tokens per second)
llama_print_timings:        eval time =   29657,97 ms /   386 runs   (   76,83 ms per token,    13,02 tokens per second)
llama_print_timings:       total time =   37612,63 ms /  1852 tokens

//> Mixtral 8x7b Q4_K_M
// HSA_OVERRIDE_GFX_VERSION=11.0.1  ./main -m mixtral-8x7b-instruct-v0.1.Q4_K_M.gguf -ngl 999 --no-mmap --temp 0 -c 2048 -p "[INST] ...
llama_print_timings:        load time =   12226.79 ms
llama_print_timings:      sample time =       1.03 ms /   138 runs   (    0.01 ms per token, 134502.92 tokens per second)
llama_print_timings: prompt eval time =   16152.11 ms /  1466 tokens (   11.02 ms per token,    90.76 tokens per second)
llama_print_timings:        eval time =   19030.05 ms /   137 runs   (  138.91 ms per token,     7.20 tokens per second)
llama_print_timings:       total time =   35221.27 ms /  1603 tokens

//> Mixtral 8x7b Q6_K
// HSA_OVERRIDE_GFX_VERSION=11.0.1  ./main -m mixtral-8x7b-instruct-v0.1.Q6_K.gguf -ngl 999 --no-mmap --temp 0 -c 2048 -p "[INST] ...
llama_print_timings:        load time =   30726,87 ms
llama_print_timings:      sample time =      27,06 ms /   388 runs   (    0,07 ms per token, 14339,04 tokens per second)
llama_print_timings: prompt eval time =   16200,58 ms /  1466 tokens (   11,05 ms per token,    90,49 tokens per second)
llama_print_timings:        eval time =   58623,74 ms /   387 runs   (  151,48 ms per token,     6,60 tokens per second)
llama_print_timings:       total time =   74941,65 ms /  1853 tokens

//> Meta-Llama-3-8B-Instruct.fp16.gguf
// HSA_OVERRIDE_GFX_VERSION=11.0.1  ./main -m Meta-Llama-3-8B-Instruct.fp16.gguf -ngl 999 --no-mmap --temp 0 -c 2048 -p "[INST] ...
llama_print_timings:        load time =    8016.27 ms
llama_print_timings:      sample time =       3.26 ms /   110 runs   (    0.03 ms per token, 33721.64 tokens per second)
llama_print_timings: prompt eval time =    6542.20 ms /  1328 tokens (    4.93 ms per token,   202.99 tokens per second)
llama_print_timings:        eval time =   22346.55 ms /   109 runs   (  205.01 ms per token,     4.88 tokens per second)
llama_print_timings:       total time =   28941.99 ms /  1437 tokens

[slaren]

I am not sure what issue you are having with mmap, but we cannot change the behavior of llama.cpp based on the configuration of a backend.

[Djip007]

The "issue"... not completely sure. I will try to explain how I see it (I don't know if that's "true")

Imagine: my computer have 64Go of RAM, 4Go is reserve for the GPU (VRAM) in bios. So 60Go of RAM is usable. OS/firefox... use 8Go (an example ;) ) so I have 52Go of RAM for the model.
The model is 48Go ...

When it is load with mmap active, and UMA, each weight is read from mmap and copie in a hipMallocManaged new memory. read from mmap involves loading in RAM the tensor and after copie it in "GPU_UMA". This means the tensor is 2 times in RAM... after 52/2=26Go of load no more memory is available, and the kernel need to find un-needed RAM. look it consider the memory that is the oldest accessed ... mmap "RAM" was used a short time ago... so Firefox RAM is Transfer to swap ... with more load it finally evicts the oldest mmap "RAM"...
With no mmap the weight look like it is directly copied on "GPU_UMA" (this is what I am least sure of...) so all the 48Go of tensor have enough space to be copied.

So don't really know, but mmap+swap+UMA ... puts a lot of pressure on Linux memory management...

For me use mmap with "full" GPU offloading is a bad idee ... (with normal "cudaMalloc" on device, imagine mmap on computer with 8G of RAM and a GPU of 48Go of VRAM (I know, it's not really a good choice ;) ) but if the model fit in VRAM it may be possible to made it work...

That say, I don't know if we simply need to give some advise/good practice ... or make llama.cpp more intelligent to chose "better" strategie ... like for exemple use mmap by default only if the weight will not be copied on "local backend" ... but simple to say... not to "code"

ps: I have many more question with llama.cpp memory management (like: with UMA, tensor are in RAM. Did we have to do something to avoid copie back in RAM if an OP is execute on CPU? ...) but I need to have closer look on that and open an other issue if I can't figure how it work...

[slaren]

Ok, I understand, thanks for the explanation. For the CPU and Metal backends, we have functions to create a backend buffer from a pointer. In this way, the backend can use the memory mapped file without additional copies. I think that may be possible to do this with CUDA/HIP for UMA devices using hip/cudaHostRegister. We would need to add a function similar to ggml_backend_cpu_buffer_from_ptr and ggml_backend_metal_buffer_from_ptr for the CUDA backend, and add support in llm_load_tensors.

[slaren]

I have made a test implementation in 518b752. It works for me on a 3090, although slowly.