Has anyone tried ZLUDA?

[shanecranor]

Curious how the performance of this compares to ZLUDA, however I couldn't get ZLUDA to work for kobaldcpp on my machine

ggml_init_cublas: found 1 CUDA devices:
CUDA error: invalid argument
current device: 0, in function ggml_init_cublas at d:\a\koboldcpp\koboldcpp\ggml-cuda.cu:7961
cuDeviceGetAttribute(&device_vmm, CU_DEVICE_ATTRIBUTE_VIRTUAL_MEMORY_MANAGEMENT_SUPPORTED, device)
GGML_ASSERT: d:\a\koboldcpp\koboldcpp\ggml-cuda.cu:241: !"CUDA error"

Not sure if this is a me problem or if this requires features that ZLUDA doesn't support.

[Negatrev]

Just in case, I'm pretty sure the zluda docs say you need to be running ROCm 5.7, not ROCm 6. Are you running the older version?

[EmanuelOverride]

With ROCm 5.7, running the benchmark on the oldcpu version, VRAM gets allocated, everything gets ready to roll, then:

CUDA error: invalid device function
current device: 0, in function launch_fattn at d:\a\koboldcpp\koboldcpp\ggml\src\ggml-cuda\fattn-common.cuh:687
cudaGetLastError()
d:\a\koboldcpp\koboldcpp\ggml\src\ggml-cuda.cu:103: CUDA error

103 seems to be a "validity test" error:

ERROR: CUDA error (code 103, UnknownMember) Introduced in CUDA 11.2, this error return indicates that at least one of these tests has failed and the validity of either the runtime or the driver could not be established

So presumably, something along the lines of a "-skip-cuda-test" command might resolve it? With sd during tasks that combine a bunch of different functions it sometimes takes a few repeated attempts for zluda to generate a workable zluda.db file, all the while reporting errors and aborting prematurely, then once the file stops growing it's generally set to work from then on. Perfect strat for a luddite like myself, just keep banging my head against the wall until it "magically" works. ;)

EDIT: I forgot an important step that resolved the error -- the cublas.dll that's provided with Zluda should be inserted instead of the cublas64_11.dll provided with Kobold. This isn't always necessary with image generation, to make it work that is, but it seems to be crucial to get an application to actually recognize your GPU as compatible. So, with a bit of fiddling around, Zluda-CUDA is indeed functional (although I'm running everything at bare minimum at the moment). I'll do a bit of further testing, to see if any particular options break it.

However, there's already lots that could be improved on my end to make the process smoother, if only I could figure out how:

• Injecting cublas via zluda, since the option appears to be there but on testing seemed to do nothing, would allow me to run the kobold.exe rather than koboldcpp.py. (Update: Nope, apparently both zluda --cublas injection and direct replacement of file necessary. It says "inject instead of system dll" so I guess I misunderstood the function and a separate default cublas file on my system is incompatible?)
• Can't get gui to work with koboldcpp.py, having to launch with skipgui which prohibits additional launch commands? So no benchmark?
• Ultimately, zluda-cuda could be added as another selectable option in the gui if the additional dll could be added without conflict with regular cuda. Already from the short test I did, it seems to perform quite well, and more options is more good right?

[EmanuelOverride]

Ok, a quick test later and all the settings work, as in the application starts up and the model generates, except flashattention which errors out with aforementioned error 103. Note, this may be due to my GPU which is not supported by HIP for Win (I think?). From past experience, even if I managed to launch with flashattention, bypassing the initial test, it would probably crash or freeze during generation anyway, but I don't know enough about what's what to confidently say your experience will be the same. On my RX 480, Zluda is about the fastest alternative there is, but screws up memory operations -- 20 restarts later, I've gone from 200 idle VRAM to 900. In other words, it clutters up the place and any functionality that deals with memory optimization may or may not fail spectacularly, at least in my experience.

Quick summary: Replaced Kobold's cublas64_11.dll with Zluda's cublas.dll (renamed), ran "zluda --cublas path/cublas64_11.dll -- python koboldcpp.py --skiplauncher", selected a config with everything on except flashattention and it works.

EDIT: MMQ also breaks it for me,

Processing Prompt [BLAS] (256 / 3748 tokens)CUDA error: invalid device function
current device: 0, in function launch_mul_mat_q at d:\a\koboldcpp\koboldcpp\ggml\src\ggml-cuda\template-instances../mmq.cuh:2771
cudaFuncSetAttribute(mul_mat_q<type, mmq_x, 8, false>, cudaFuncAttributeMaxDynamicSharedMemorySize, shmem)
d:\a\koboldcpp\koboldcpp\ggml\src\ggml-cuda.cu:103: CUDA error

EDIT2: Ran for about an hour, everything stable, great gpu utilization, faster than Vulkan but I can't do a proper benchmark. Looks like the biggest performance gains are in processing, roughly +10T/s and +/- same generation speed. Until I figure out ROCm with AVX1 my personal performance chart is Zluda-Cuda>Vulkan>ZludaOpenCL>OpenCL.

[brknsoul]

If you edit your .kcpps config file and set "benchmark" : true, then you can run benchmarks;

Processing Prompt [BLAS] (3996 / 3996 tokens)
Generating (100 / 100 tokens)
CtxLimit:4096/4096, Amt:100/100, Process:16.17s (4.0ms/T = 247.08T/s), Generate:13.12s (131.2ms/T = 7.62T/s), Total:29.29s (3.41T/s)
Benchmark Completed - v1.72 Results:
======
Flags: NoAVX2=False Threads=12 HighPriority=False NoBlas=False Cublas_Args=['normal', '0'] Tensor_Split=None BlasThreads=12 BlasBatchSize=512 FlashAttention=False KvCache=0
Timestamp: 2024-08-03 06:45:48.327973+00:00
Backend: koboldcpp_cublas.dll
Layers: 200
Model: openhermes-2.5-mistral-7b.Q4_K_M
MaxCtx: 4096
GenAmount: 100
-----
ProcessingTime: 16.173s
ProcessingSpeed: 247.08T/s
GenerationTime: 13.116s
GenerationSpeed: 7.62T/s
TotalTime: 29.289s
Output: 11111
-----

Weird.. the benchmark is much slower than actual generation:
Process:2.09s (3.8ms/T = 265.58T/s), Generate:4.29s (33.2ms/T = 30.10T/s), Total:6.37s (20.24T/s)

[EmanuelOverride]

Good call, hadn't thought of that:

Processing Prompt [BLAS] (3996 / 3996 tokens)
Generating (100 / 100 tokens)
CtxLimit:4096/4096, Amt:100/100, Process:34.12s (8.5ms/T = 117.12T/s), Generate:31.79s (317.9ms/T = 3.15T/s), Total:65.91s (1.52T/s)
Benchmark Completed - v1.71.1 Results:

Flags: NoAVX2=False(actually True) Threads=4 HighPriority=False NoBlas=False Cublas_Args=['normal', '0'] Tensor_Split=None BlasThreads=4 BlasBatchSize=512 FlashAttention=False KvCache=0
Timestamp: 2024-08-03 10:31:06.650690+00:00
Backend: koboldcpp_cublas.dll
Layers: 35
Model: dolphin-2.6-mistral-7b.Q5_K_M
MaxCtx: 4096
GenAmount: 100

ProcessingTime: 34.120s
ProcessingSpeed: 117.12T/s
GenerationTime: 31.793s
GenerationSpeed: 3.15T/s
TotalTime: 65.913s
Output: 11111

Compared to my Vulkan, same settings:

ProcessingTime: 51.741s
ProcessingSpeed: 77.23T/s
GenerationTime: 26.123s
GenerationSpeed: 3.83T/s
TotalTime: 77.864s
Output: 11111

Zluda does like to continuously add to AppData\Local\ZLUDA\ComputeCache\zluda.db whenever any element of the ongoing operation seems unfamiliar or subject to performance adjustment, which delays whatever CUDA process you're running, at least the very first time you run it. But, there may also be unresolved elements left over that need several passes, so maybe that could explain it? (EDIT: In addition to that, I believe Zluda draws from zluda.db on a need to know basis, which causes the first draw-turn to be slower, but then maintains that data for subsequent operations which means short-term performance improvement over time. In other words, even if your compute cache has been finalized for the type of operations that you're running, each new instance of that operation may start out slightly slower before reaching peak performance. Basically more grounded version of my analogy below.)

Other than that, it's my general experience that bench is always slower than actual gen on average, like bench gives the worst-case scenario result that may or may not crop up during actual operation -- like revving your engine combined with a cold start and comparing the way it performs in that scenario to a long and fast drive on a highway where all the diverse elements of operation have time to settle in place. Maybe accurate analogy?

Lastly, I've become curious about https://github.com/likelovewant/ROCmLibs-for-gfx1103-AMD780M-APU in relation to my quest for ROCm AVX1 -- if it's anything, it should have some overlap with Zluda performance as well. Won't have time to properly try it out any time soon though.

EDIT2: I'm pressed for time but still wanted to try it out. :) With previous link in mind, running with rocblas & library from OverrideWithVega 10.7 for gfx803 is stable, gives performance improvement basically within margin of error:

ProcessingTime: 33.873s
ProcessingSpeed: 117.97T/s
GenerationTime: 31.810s
GenerationSpeed: 3.14T/s
TotalTime: 65.683s
Output: 11111

Allowed Zluda to generate new Compute Cache, it shrank in half of previous size. So I guess it's safe to assume that these are something like hand-made pre-calculations that Zluda would otherwise have to figure out? What isn't reflected in the benchmark stats is the lower memory overhead and faster startup. Looks promising enough, considering updating hip to 6.1.2 and testing it against all the various front-ends I've accumulated.

If you're trying it out, could you keep an eye on memory and memory optimization? Does your setup also accumulate junk memory in vram over time (several instance runs)? Do you also get "malloc: failed to allocate" as standard? I believe, not 100%, I just lost 2 malloc failures that I would otherwise have when running with this new library, but maybe that's wishful thinking on my part. Need time to do proper comparative tests, but it's been my experience that Zluda struggles with mem-ops.

[brknsoul]

Yeah, the same happens with SD, the first generation of a new install without a zluda db takes about 10-15mins to build the db.

---

Odd, it seems if I manually run a benchmark with >3000 tokens (an input of 3072 "1s") the output is significantly slower in Zluda (generate: 10.43T/s) than in Vulkan (generate: 21.39T/s)

But if use only 512 1s, I get
Zluda: 31.49T/s
Vulkan: 35.29T/s

[EmanuelOverride]

Installed HIP 6.1 to test gpu-specific libraries for ROCm, Zluda doesn't support 6.1 out of the box, asks for 5.7 specific files, tried building own Zluda, errors out -- save that experiment for when I have the time.

[EmanuelOverride]

Newer versions of Zluda don't work with KoboldCPP due to --cublas command being removed. With 11cc, replacing cublas64_11.dll without injecting causes koboldcpp_cublas.dll to fail (file not found), while doing the same and injecting with --cublas works fine.