3 ms·
Thank you for sharing! Would've loved to peek at your CUDA code but looks like it's gone now. Looking forward to seeing the full version when it's finished.
by reasonableklout 2y ago
Thank you for sharing! Would've loved to peek at your CUDA code but looks like it's gone now. Looking forward to seeing the full version when it's finished.
- ryao 2y agoAsk and you shall receive: https://github.com/ryao/llama3.c/commit/dc27f2b1c6884c72bd4bd3d8409a94bd12fdcaaf https://github.com/ryao/llama3.c/commit/dc27f2b1c6884c72bd4b... It is still a WIP, but the CUDA kernels are unlikely to see much improvement for a while as I had gone over them with a fine toothed comb for performance in the past few days. The C code is more likely where I will be making improvements. For example, I need to switch to a structure to stop abusing an array of void pointers for the cudaMemcpy() aggregration trick for the cublasGemmBatchedEx() calls. The current version is performing either on par to or slightly better than llama.cpp's F16 codepath on my 3090 Ti, and I have not even implemented CUDA graphs yet. Both get around 53 tokens per second, although on really good runs, I have seen my code do 54 tokens per second and I have not seen that from llama.cpp so far.
- reasonableklout 2y agoVery cool! I have been bedridden and unable to do much work the last few days but it will be cool to dig into this.
- ryao 2y agoI hope you feel better soon. If you want to get in touch in the future after we have both stoped checking this thread, you can find my email on one of my commits on GitHub.
- reasonableklout 2y agoThanks, will do!