Mirrors/llama.cpp
1
0
Fork 0
mirror of https://github.com/ggerganov/llama.cpp.git synced 2025-01-27 20:43:07 +01:00
Code Issues Packages Projects Releases Wiki Activity

ggml-cuda : remove some superfluous comments for TQ2_0 tile loading

Browse Source
This commit is contained in:
Francis Couture-Harpin 2025-01-10 14:52:49 -05:00
parent 983aa09b5c
commit f5fddb6d24
1 changed files with 0 additions and 2 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
ggml/src/ggml-cuda/mmq.cuh
View File

@ -1848,13 +1848,11 @@ template <int mmq_y, int nwarps, bool need_check> static __device__ __forceinlin
#ifdef INT8_MMA_AVAILABLE #ifdef INT8_MMA_AVAILABLE
x_qs[i*MMQ_MMA_TILE_X_K_Q8_0 + k] = q; x_qs[i*MMQ_MMA_TILE_X_K_Q8_0 + k] = q;
#else #else
// NOTE: this might assume WARP_SIZE is >= 32
x_qs[i*(2*WARP_SIZE + 1) + k] = q; x_qs[i*(2*WARP_SIZE + 1) + k] = q;
#endif // INT8_MMA_AVAILABLE #endif // INT8_MMA_AVAILABLE
} }
} }
// TODO: does this work with WARP_SIZE != 32?
#pragma unroll #pragma unroll
for (int i0 = 0; i0 < mmq_y; i0 += nwarps * WARP_SIZE/(QI2_0/2)) { for (int i0 = 0; i0 < mmq_y; i0 += nwarps * WARP_SIZE/(QI2_0/2)) {
int i = i0 + threadIdx.y*(2*WARP_SIZE/QI2_0) + threadIdx.x/(QI2_0/2); int i = i0 + threadIdx.y*(2*WARP_SIZE/QI2_0) + threadIdx.x/(QI2_0/2);