vbatts/llama.cpp
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

Add some minimal optimizations for CDNA (#10498)

Browse source 
* Add some minimal optimizations for CDNA

* ggml_cuda: set launch bounds also for GCN as it helps there too
This commit is contained in:
uvos 2024-11-27 17:10:08 +01:00 • committed by GitHub
parent 46c69e0e75
commit 3ad5451f3b
No known key found for this signature in database
GPG key ID: B5690EEEBB952194
6 changed files with 36 additions and 8 deletions
Download patch file Download diff file Expand all files Collapse all files

4
ggml/src/ggml-cuda/mmq.cuh
View file

@ -2570,9 +2570,9 @@ static __device__ void mul_mat_q_process_tile(
template <ggml_type type, int mmq_x, int nwarps, bool need_check>
#if defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)
#if defined(RDNA3) || defined(RDNA2)
#if defined(RDNA3) || defined(RDNA2) || defined(CDNA) || defined(GCN)
__launch_bounds__(WARP_SIZE*nwarps, 2)
#endif // defined(RDNA3) || defined(RDNA2)
#endif // defined(RDNA3) || defined(RDNA2) || defined(CDNA) || defined(GCN)
#else
#if __CUDA_ARCH__ >= CC_VOLTA
__launch_bounds__(WARP_SIZE*nwarps, 1)