char-fork/llama.cpp
git clone https://git.t4t.associates/char-fork/llama.cpp
e42214804
char/aliceai-t5-moe
1#pragma once 2 3#include "vecdotq.cuh" 4 5#include "mmq.cuh" 6 7template <ggml_type type, int J, bool fallback> static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q1_0( 8 const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { 9 constexpr int warp_size = ggml_cuda_get_physical_warp_size(); 10 constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; 11 constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); 12 constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); 13 14#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 15 int * x_qs = (int *) x_tile; 16 float * x_df = (float *) (x_qs + 2*MMQ_TILE_NE_K); 17#else 18 constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_Q8_0, I); 19 int * x_qs = (int *) x_tile; 20 float * x_df = (float *) (x_qs + txs.qs); 21#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 22 23 constexpr int blocks_per_iter = MMQ_ITER_K / QK1_0; 24 constexpr int threads_per_row = blocks_per_iter * QI1_0; 25 constexpr int nrows = warp_size / threads_per_row; 26 constexpr int scale_entries_per_block = QK1_0 / QK8_1; 27 constexpr int scale_entries_per_row = blocks_per_iter * scale_entries_per_block; 28 29 const int txi = threadIdx.x % threads_per_row; 30 const int kbx = txi / QI1_0; 31 const int kqsx = txi % QI1_0; 32 33#pragma unroll 34 for (int i0 = 0; i0 < I; i0 += nrows*nwarps) { 35 int i = i0 + threadIdx.y*nrows + threadIdx.x/threads_per_row; 36 37 if (fallback) { 38 i = min(i, i_max); 39 } 40 41 const block_q1_0 * bxi = (const block_q1_0 *) x + kbx0 + i*stride + kbx; 42 const int16_t * qxi = (const int16_t *) bxi->qs + kqsx * 2; 43 44 const int dst_offset = kbx*(scale_entries_per_block*QI8_0) + kqsx*QI8_0; 45#pragma unroll 46 for (int j = 0; j < 2; ++j) { 47 const int q = qxi[j]; 48 49 // unpack crumbs into nibble indices 50 const int n0 = __byte_perm(0x11100100, 0x11100100, q >> 0); // [0, 1, 4, 5] [ 8, 9, 12, 13] 51 const int n1 = __byte_perm(0x11100100, 0x11100100, q >> 2); // [2, 3, 6, 7] [10, 11, 14, 15] 52 // unpack nibbles into byte values 53 const int s0 = __byte_perm(0x01FF, 0x01FF, n0 >> 0); 54 const int s1 = __byte_perm(0x01FF, 0x01FF, n1 >> 0); 55 const int s2 = __byte_perm(0x01FF, 0x01FF, n0 >> 16); 56 const int s3 = __byte_perm(0x01FF, 0x01FF, n1 >> 16); 57 // unshuffle values 58 const int v0 = __byte_perm(s0, s1, 0x5410); 59 const int v1 = __byte_perm(s0, s1, 0x7632); 60 const int v2 = __byte_perm(s2, s3, 0x5410); 61 const int v3 = __byte_perm(s2, s3, 0x7632); 62 63#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 64 x_qs[i*sram_stride + dst_offset + j*4+0] = v0; 65 x_qs[i*sram_stride + dst_offset + j*4+1] = v1; 66 x_qs[i*sram_stride + dst_offset + j*4+2] = v2; 67 x_qs[i*sram_stride + dst_offset + j*4+3] = v3; 68#else 69 x_qs[i*(2*MMQ_TILE_NE_K + 1) + dst_offset + j*4+0] = v0; 70 x_qs[i*(2*MMQ_TILE_NE_K + 1) + dst_offset + j*4+1] = v1; 71 x_qs[i*(2*MMQ_TILE_NE_K + 1) + dst_offset + j*4+2] = v2; 72 x_qs[i*(2*MMQ_TILE_NE_K + 1) + dst_offset + j*4+3] = v3; 73#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 74 } 75 } 76 77 const int ksx = threadIdx.x % scale_entries_per_row; 78 const int scale_block = ksx / scale_entries_per_block; 79 80#pragma unroll 81 for (int i0 = 0; i0 < I; i0 += nwarps) { 82 int i = i0 + threadIdx.y; 83 84 if (fallback) { 85 i = min(i, i_max); 86 } 87 88 const block_q1_0 * bxi = (const block_q1_0 *) x + kbx0 + i*stride + scale_block; 89 90#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 91 x_df[i*sram_stride + ksx] = bxi->d; 92#else 93 x_df[i*(2*MMQ_TILE_NE_K/QI8_0) + i/(QI8_0/2) + ksx] = bxi->d; 94#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 95 } 96} 97 98template <ggml_type type, int J, bool fallback> static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q2_0( 99 const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { 100 constexpr int warp_size = ggml_cuda_get_physical_warp_size(); 101 constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; 102 constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); 103 constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); 104 105#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 106 int * x_qs = (int *) x_tile; 107 float * x_df = (float *) (x_qs + 2*MMQ_TILE_NE_K); 108#else 109 constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_Q8_0, I); 110 int * x_qs = (int *) x_tile; 111 float * x_df = (float *) (x_qs + txs.qs); 112#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 113 114 constexpr int blocks_per_iter = MMQ_ITER_K / QK2_0; 115 constexpr int threads_per_row = blocks_per_iter * QI2_0; 116 constexpr int nrows = warp_size / threads_per_row; 117 constexpr int scale_entries_per_block = QK2_0 / QK8_1; 118 constexpr int scale_entries_per_row = blocks_per_iter * scale_entries_per_block; 119 120 const int txi = threadIdx.x % threads_per_row; 121 const int kbx = txi / QI2_0; 122 const int kqsx = txi % QI2_0; 123 124#pragma unroll 125 for (int i0 = 0; i0 < I; i0 += nrows*nwarps) { 126 int i = i0 + threadIdx.y*nrows + threadIdx.x/threads_per_row; 127 128 if (fallback) { 129 i = min(i, i_max); 130 } 131 132 const block_q2_0 * bxi = (const block_q2_0 *) x + kbx0 + i*stride + kbx; 133 const int16_t * qxi = (const int16_t *) bxi->qs + kqsx * 4; 134 135 const int dst_offset = kbx*(scale_entries_per_block*QI8_0) + kqsx*QI8_0; 136 137#pragma unroll 138 for (int j = 0; j < 4; ++j) { 139 const int q = qxi[j]; 140 141#if defined(GGML_USE_HIP) 142 const uint32_t qx_indices = (q & 0x03) | ((q & 0x0C) << 6) | ((q & 0x30) << 12) | ((q & 0xC0) << 18); 143 const uint32_t qy_bits = q >> 8; 144 const uint32_t qy_indices = (qy_bits & 0x03) | ((qy_bits & 0x0C) << 6) | ((qy_bits & 0x30) << 12) | ((qy_bits & 0xC0) << 18); 145 const int qx = __builtin_amdgcn_perm(0x020100FF, 0x020100FF, qx_indices); 146 const int qy = __builtin_amdgcn_perm(0x020100FF, 0x020100FF, qy_indices); 147#else 148 // unpack even and odd crumbs into byte values 149 const int qe = __byte_perm(0x020100FF, 0x020100FF, q >> 0); 150 const int qo = __byte_perm(0x020100FF, 0x020100FF, q >> 2); 151 // unshuffle values 152 const int qx = __byte_perm(qe, qo, 0x5140); 153 const int qy = __byte_perm(qe, qo, 0x7362); 154#endif // defined(GGML_USE_HIP) 155 156#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 157 x_qs[i*sram_stride + dst_offset + j*2+0] = qx; 158 x_qs[i*sram_stride + dst_offset + j*2+1] = qy; 159#else 160 x_qs[i*(2*MMQ_TILE_NE_K + 1) + dst_offset + j*2+0] = qx; 161 x_qs[i*(2*MMQ_TILE_NE_K + 1) + dst_offset + j*2+1] = qy; 162#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 163 } 164 } 165 166 const int ksx = threadIdx.x % scale_entries_per_row; 167 const int scale_block = ksx / scale_entries_per_block; 168 169#pragma unroll 170 for (int i0 = 0; i0 < I; i0 += nwarps) { 171 int i = i0 + threadIdx.y; 172 173 if (fallback) { 174 i = min(i, i_max); 175 } 176 177 const block_q2_0 * bxi = (const block_q2_0 *) x + kbx0 + i*stride + scale_block; 178 179#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 180 x_df[i*sram_stride + ksx] = bxi->d; 181#else 182 x_df[i*(2*MMQ_TILE_NE_K/QI8_0) + i/(QI8_0/2) + ksx] = bxi->d; 183#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 184 } 185} 186 187template <ggml_type type, int J, bool fallback> static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q4_0( 188 const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { 189 constexpr int warp_size = ggml_cuda_get_physical_warp_size(); 190 constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; 191 constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); 192 constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); 193 194#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 195 int * x_qs = (int *) x_tile; 196 float * x_df = (float *) (x_qs + 2*MMQ_TILE_NE_K); 197#else 198 constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_Q4_0, I); 199 int * x_qs = (int *) x_tile; 200 float * x_df = (float *) (x_qs + txs.qs); 201#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 202 203 constexpr int threads_per_row = MMQ_ITER_K / (4 * QR4_0); 204 constexpr int nrows = warp_size / threads_per_row; 205 const int txi = warp_size > threads_per_row ? threadIdx.x % threads_per_row : threadIdx.x; 206 const int kbx = txi / QI4_0; 207 const int kqsx = txi % QI4_0; 208 209#pragma unroll 210 for (int i0 = 0; i0 < I; i0 += nrows*nwarps) { 211 int i = i0 + (nrows == 1 ? threadIdx.y : threadIdx.y*nrows + threadIdx.x/threads_per_row); 212 213 if (fallback) { 214 i = min(i, i_max); 215 } 216 217 const block_q4_0 * bxi = (const block_q4_0 *) x + kbx0 + i*stride + kbx; 218 const int qs0 = get_int_b2(bxi->qs, kqsx); 219 220#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 221 x_qs[i*sram_stride + kbx*(2*QI4_0) + kqsx + 0] = __vsubss4((qs0 >> 0) & 0x0F0F0F0F, 0x08080808); 222 x_qs[i*sram_stride + kbx*(2*QI4_0) + kqsx + QI4_0] = __vsubss4((qs0 >> 4) & 0x0F0F0F0F, 0x08080808); 223#else 224 x_qs[i*(MMQ_TILE_NE_K + 1) + txi] = qs0; 225#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) 226 } 227 228 constexpr int blocks_per_tile_x_row = MMQ_TILE_NE_K / QI4_0; 229 constexpr int rows_per_warp = warp_size / blocks_per_tile_x_row; 230 const int kbxd = threadIdx.x % blocks_per_tile_x_row; 231 232#pragma unroll 233 for (int i0 = 0; i0 < I; i0 += nwarps * rows_per_warp) { 234 int i = i0 + threadIdx.y * rows_per_warp + threadIdx.x / blocks_per_tile_x_row; 235 236 if (fallback) { 237 i = min(i, i_max); 238 } 239 240 const block_q4_0 * bxi = (const block_q4_0 *) x + kbx0 + i*stride + kbxd; 241 242#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 243 x_df[i*sram_stride + kbxd] = bxi->d; 244#else 245 x_df[i*(MMQ_TILE_NE_K/QI4_0) + i/QI4_0 + kbxd] = bxi->d; 246#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 247 } 248} 249 250template <ggml_type type, int J, bool fallback> static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q4_1( 251 const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { 252 constexpr int warp_size = ggml_cuda_get_physical_warp_size(); 253 constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; 254 constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); 255 constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); 256 257#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 258 int * x_qs = (int *) x_tile; 259 half2 * x_dm = (half2 *) (x_qs + 2*MMQ_TILE_NE_K); 260#else 261 constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_Q4_1, I); 262 int * x_qs = (int *) x_tile; 263 half2 * x_dm = (half2 *) (x_qs + txs.qs); 264#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 265 266 constexpr int threads_per_row = MMQ_ITER_K / (4 * QR4_1); 267 constexpr int nrows = warp_size / threads_per_row; 268 const int txi = warp_size > threads_per_row ? threadIdx.x % threads_per_row : threadIdx.x; 269 const int kbx = txi / QI4_1; 270 const int kqsx = txi % QI4_1; 271 272#pragma unroll 273 for (int i0 = 0; i0 < I; i0 += nrows*nwarps) { 274 int i = i0 + (nrows == 1 ? threadIdx.y : threadIdx.y*nrows + threadIdx.x/threads_per_row); 275 276 if (fallback) { 277 i = min(i, i_max); 278 } 279 280 const block_q4_1 * bxi = (const block_q4_1 *) x + kbx0 + i*stride + kbx; 281 const int qs0 = get_int_b4(bxi->qs, kqsx); 282 283#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 284 x_qs[i*sram_stride + kbx*(2*QI4_1) + kqsx + 0] = (qs0 >> 0) & 0x0F0F0F0F; 285 x_qs[i*sram_stride + kbx*(2*QI4_1) + kqsx + QI4_1] = (qs0 >> 4) & 0x0F0F0F0F; 286#else 287 x_qs[i*(MMQ_TILE_NE_K + 1) + txi] = qs0; 288#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 289 } 290 291 constexpr int blocks_per_tile_x_row = MMQ_TILE_NE_K / QI4_1; 292 constexpr int rows_per_warp = warp_size / blocks_per_tile_x_row; 293 const int kbxd = threadIdx.x % blocks_per_tile_x_row; 294 295#pragma unroll 296 for (int i0 = 0; i0 < I; i0 += nwarps * rows_per_warp) { 297 int i = i0 + threadIdx.y * rows_per_warp + threadIdx.x / blocks_per_tile_x_row; 298 299 if (fallback) { 300 i = min(i, i_max); 301 } 302 303 const block_q4_1 * bxi = (const block_q4_1 *) x + kbx0 + i*stride + kbxd; 304 305#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 306 x_dm[i*sram_stride + kbxd] = bxi->dm; 307#else 308 x_dm[i*(MMQ_TILE_NE_K/QI4_1) + i/QI4_1 + kbxd] = bxi->dm; 309#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 310 } 311} 312 313template <ggml_type type, int J, bool fallback> static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q5_0( 314 const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { 315 constexpr int warp_size = ggml_cuda_get_physical_warp_size(); 316 constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; 317 constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); 318 constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); 319 320#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 321 int * x_qs = (int *) x_tile; 322 float * x_df = (float *) (x_qs + MMQ_TILE_NE_K*2); 323#else 324 constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_Q5_0, I); 325 int * x_qs = (int *) x_tile; 326 float * x_df = (float *) (x_qs + txs.qs); 327#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 328 329 constexpr int threads_per_row = MMQ_ITER_K / (4 * QR5_0); 330 constexpr int nrows = warp_size / threads_per_row; 331 const int txi = warp_size > threads_per_row ? threadIdx.x % threads_per_row : threadIdx.x; 332 const int kbx = txi / QI5_0; 333 const int kqsx = txi % QI5_0; 334 335#pragma unroll 336 for (int i0 = 0; i0 < I; i0 += nrows*nwarps) { 337 int i = i0 + (nrows == 1 ? threadIdx.y : threadIdx.y*nrows + threadIdx.x/threads_per_row); 338 339 if (fallback) { 340 i = min(i, i_max); 341 } 342 343 const block_q5_0 * bxi = (const block_q5_0 *) x + kbx0 + i*stride + kbx; 344 345 const int ql = get_int_b2(bxi->qs, kqsx); 346 const int qh = get_int_b2(bxi->qh, 0) >> (4 * kqsx); 347 348 int qs0 = (ql >> 0) & 0x0F0F0F0F; 349 qs0 |= (qh << 4) & 0x00000010; // 0 -> 4 350 qs0 |= (qh << 11) & 0x00001000; // 1 -> 12 351 qs0 |= (qh << 18) & 0x00100000; // 2 -> 20 352 qs0 |= (qh << 25) & 0x10000000; // 3 -> 28 353 qs0 = __vsubss4(qs0, 0x10101010); // subtract 16 354 355 int qs1 = (ql >> 4) & 0x0F0F0F0F; 356 qs1 |= (qh >> 12) & 0x00000010; // 16 -> 4 357 qs1 |= (qh >> 5) & 0x00001000; // 17 -> 12 358 qs1 |= (qh << 2) & 0x00100000; // 18 -> 20 359 qs1 |= (qh << 9) & 0x10000000; // 19 -> 28 360 qs1 = __vsubss4(qs1, 0x10101010); // subtract 16 361 362#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 363 x_qs[i*sram_stride + kbx*(2*QI5_0) + kqsx + 0] = qs0; 364 x_qs[i*sram_stride + kbx*(2*QI5_0) + kqsx + QI5_0] = qs1; 365#else 366 x_qs[i*(2*MMQ_TILE_NE_K + 1) + kbx*(2*QI5_0) + kqsx + 0] = qs0; 367 x_qs[i*(2*MMQ_TILE_NE_K + 1) + kbx*(2*QI5_0) + kqsx + QI5_0] = qs1; 368#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 369 } 370 371 constexpr int blocks_per_tile_x_row = MMQ_TILE_NE_K / QI5_0; 372 constexpr int rows_per_warp = warp_size / blocks_per_tile_x_row; 373 const int kbxd = threadIdx.x % blocks_per_tile_x_row; 374 375#pragma unroll 376 for (int i0 = 0; i0 < I; i0 += nwarps * rows_per_warp) { 377 int i = i0 + threadIdx.y * rows_per_warp + threadIdx.x / blocks_per_tile_x_row; 378 379 if (fallback) { 380 i = min(i, i_max); 381 } 382 383 const block_q5_0 * bxi = (const block_q5_0 *) x + kbx0 + i*stride + kbxd; 384 385#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 386 x_df[i*sram_stride + kbxd] = bxi->d; 387#else 388 x_df[i*(MMQ_TILE_NE_K/QI5_0) + i/QI5_0 + kbxd] = bxi->d; 389#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 390 } 391} 392 393template <ggml_type type, int J, bool fallback> static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q5_1( 394 const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { 395 constexpr int warp_size = ggml_cuda_get_physical_warp_size(); 396 constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; 397 constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); 398 constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); 399 400#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 401 int * x_qs = (int *) x_tile; 402 half2 * x_dm = (half2 *) (x_qs + 2*MMQ_TILE_NE_K); 403#else 404 constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_Q5_1, I); 405 int * x_qs = (int *) x_tile; 406 half2 * x_dm = (half2 *) (x_qs + txs.qs); 407#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 408 409 constexpr int threads_per_row = MMQ_ITER_K / (4 * QR5_1); 410 constexpr int nrows = warp_size / threads_per_row; 411 const int txi = warp_size > threads_per_row ? threadIdx.x % threads_per_row : threadIdx.x; 412 const int kbx = txi / QI5_1; 413 const int kqsx = txi % QI5_1; 414 415#pragma unroll 416 for (int i0 = 0; i0 < I; i0 += nrows*nwarps) { 417 int i = i0 + (nrows == 1 ? threadIdx.y : threadIdx.y*nrows + threadIdx.x/threads_per_row); 418 419 if (fallback) { 420 i = min(i, i_max); 421 } 422 423 const block_q5_1 * bxi = (const block_q5_1 *) x + kbx0 + i*stride + kbx; 424 425 const int ql = get_int_b4(bxi->qs, kqsx); 426 const int qh = get_int_b4(bxi->qh, 0) >> (4 * kqsx); 427 428 int qs0 = (ql >> 0) & 0x0F0F0F0F; 429 qs0 |= (qh << 4) & 0x00000010; // 0 -> 4 430 qs0 |= (qh << 11) & 0x00001000; // 1 -> 12 431 qs0 |= (qh << 18) & 0x00100000; // 2 -> 20 432 qs0 |= (qh << 25) & 0x10000000; // 3 -> 28 433 434 int qs1 = (ql >> 4) & 0x0F0F0F0F; 435 qs1 |= (qh >> 12) & 0x00000010; // 16 -> 4 436 qs1 |= (qh >> 5) & 0x00001000; // 17 -> 12 437 qs1 |= (qh << 2) & 0x00100000; // 18 -> 20 438 qs1 |= (qh << 9) & 0x10000000; // 19 -> 28 439 440#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 441 x_qs[i*sram_stride + kbx*(2*QI5_1) + kqsx + 0] = qs0; 442 x_qs[i*sram_stride + kbx*(2*QI5_1) + kqsx + QI5_1] = qs1; 443#else 444 x_qs[i*(2*MMQ_TILE_NE_K + 1) + kbx*(2*QI5_1) + kqsx + 0] = qs0; 445 x_qs[i*(2*MMQ_TILE_NE_K + 1) + kbx*(2*QI5_1) + kqsx + QI5_1] = qs1; 446#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 447 } 448 449 constexpr int blocks_per_tile_x_row = MMQ_TILE_NE_K / QI5_1; 450 constexpr int rows_per_warp = warp_size / blocks_per_tile_x_row; 451 const int kbxd = threadIdx.x % blocks_per_tile_x_row; 452 453#pragma unroll 454 for (int i0 = 0; i0 < I; i0 += nwarps * rows_per_warp) { 455 int i = i0 + threadIdx.y * rows_per_warp + threadIdx.x / blocks_per_tile_x_row; 456 457 if (fallback) { 458 i = min(i, i_max); 459 } 460 461 const block_q5_1 * bxi = (const block_q5_1 *) x + kbx0 + i*stride + kbxd; 462 463#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 464 x_dm[i*sram_stride + kbxd] = bxi->dm; 465#else 466 x_dm[i*(MMQ_TILE_NE_K/QI5_1) + i/QI5_1 + kbxd] = bxi->dm; 467#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 468 } 469} 470 471template <ggml_type type, int J, bool fallback> static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q8_0( 472 const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { 473 constexpr int warp_size = ggml_cuda_get_physical_warp_size(); 474 constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; 475 constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); 476 constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); 477 478#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 479 int * x_qs = (int *) x_tile; 480 float * x_df = (float *) (x_tile + 2*MMQ_TILE_NE_K); 481#else 482 constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_Q8_0, I); 483 int * x_qs = (int *) x_tile; 484 float * x_df = (float *) (x_qs + txs.qs); 485#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 486 487 // MMQ_ITER_K / (4 * QR8_0) == 64 required. but NV has only 32 threads per warp 488 constexpr int threads_per_row = 32; 489 constexpr int nrows = warp_size / threads_per_row; 490 const int txi = warp_size > threads_per_row ? threadIdx.x % threads_per_row : threadIdx.x; 491 const int kbx = txi / QI8_0; 492 const int kqsx = txi % QI8_0; 493 494#pragma unroll 495 for (int i0 = 0; i0 < I; i0 += nrows*nwarps) { 496 int i = i0 + (nrows == 1 ? threadIdx.y : threadIdx.y*nrows + threadIdx.x/threads_per_row); 497 498 if (fallback) { 499 i = min(i, i_max); 500 } 501 502 const block_q8_0 * bxi = (const block_q8_0 *) x + kbx0 + i*stride + kbx; 503 504#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 505 x_qs[i*sram_stride + 0 + txi] = get_int_b2(bxi[0].qs, kqsx); 506 x_qs[i*sram_stride + MMQ_TILE_NE_K + txi] = get_int_b2(bxi[MMQ_TILE_NE_K/QI8_0].qs, kqsx); 507#else 508 x_qs[i*(2*MMQ_TILE_NE_K + 1) + 0 + txi] = get_int_b2(bxi[0].qs, kqsx); 509 x_qs[i*(2*MMQ_TILE_NE_K + 1) + MMQ_TILE_NE_K + txi] = get_int_b2(bxi[MMQ_TILE_NE_K/QI8_0].qs, kqsx); 510#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 511 } 512 513 constexpr int blocks_per_tile_x_row = 2*MMQ_TILE_NE_K / QI8_0; 514 constexpr int rows_per_warp = warp_size / blocks_per_tile_x_row; 515 const int kbxd = threadIdx.x % blocks_per_tile_x_row; 516 517#pragma unroll 518 for (int i0 = 0; i0 < I; i0 += nwarps * rows_per_warp) { 519 int i = i0 + threadIdx.y * rows_per_warp + threadIdx.x / blocks_per_tile_x_row; 520 521 if (fallback) { 522 i = min(i, i_max); 523 } 524 525 const block_q8_0 * bxi = (const block_q8_0 *) x + kbx0 + i*stride + kbxd; 526 527#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 528 x_df[i*sram_stride + kbxd] = bxi->d; 529#else 530 x_df[i*(2*MMQ_TILE_NE_K/QI8_0) + i/(QI8_0/2) + kbxd] = bxi->d; 531#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 532 } 533} 534 535// --------------------------------------------------------------------------------------------- 536 537template <ggml_type type, int J, bool fallback> static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q2_K( 538 const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { 539 constexpr int warp_size = ggml_cuda_get_physical_warp_size(); 540 constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; 541 constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); 542 constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); 543 544#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 545 int * x_qs = (int *) x_tile; 546 half2 * x_dm = (half2 *) (x_qs + 2*MMQ_TILE_NE_K); 547#else 548 constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_Q2_K, I); 549 int * x_qs = (int *) x_tile; 550 half2 * x_dm = (half2 *) (x_qs + txs.qs); 551#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 552 553 constexpr int threads_per_row = MMQ_ITER_K / (4 * QR2_K); 554 constexpr int nrows = ggml_cuda_get_physical_warp_size() / threads_per_row; 555 const int kqsx = threadIdx.x % threads_per_row; 556 557#pragma unroll 558 for (int i0 = 0; i0 < I; i0 += nrows*nwarps) { 559 int i = i0 + threadIdx.y*nrows + threadIdx.x/threads_per_row; 560 561 if (fallback) { 562 i = min(i, i_max); 563 } 564 565 const block_q2_K * bxi = (const block_q2_K *) x + kbx0 + i*stride; 566 567 const int x_ql_0 = get_int_b2(bxi->qs, kqsx); 568 569#pragma unroll 570 for (int l = 0; l < QR2_K; ++l) { 571 const int k = (kqsx/8)*32 + l*8 + kqsx % 8; 572 573 const int x_qs_k = (x_ql_0 >> (2*l)) & 0x03030303; 574 575#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 576 x_qs[i*sram_stride + k] = x_qs_k; 577#else 578 x_qs[i*(2*MMQ_TILE_NE_K + 1) + k] = x_qs_k; 579#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 580 } 581 582 const int sc_m = bxi->scales[kqsx]; 583#ifdef FAST_FP16_AVAILABLE 584 const half2 x_dm_ik = __hmul2(bxi->dm, make_half2(sc_m & 0x0F, sc_m >> 4)); 585#else 586 const float2 bxi_dmf = __half22float2(bxi->dm); 587 const half2 x_dm_ik = make_half2(bxi_dmf.x*(sc_m & 0x0F), bxi_dmf.y*(sc_m >> 4)); 588#endif // FAST_FP16_AVAILABLE 589 590#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 591 x_dm[i*sram_stride + kqsx] = x_dm_ik; 592#else 593 x_dm[i*(MMQ_TILE_NE_K + 1) + kqsx] = x_dm_ik; 594#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 595 } 596} 597 598template <ggml_type type, int J, bool fallback> static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q3_K( 599 const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { 600 constexpr int warp_size = ggml_cuda_get_physical_warp_size(); 601 constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; 602 constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); 603 constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); 604 605#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 606 int * x_qs = (int *) x_tile; 607 float * x_df = (float *) (x_qs + MMQ_TILE_NE_K*2); 608#else 609 constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_Q3_K, I); 610 int * x_qs = (int *) x_tile; 611 float * x_df = (float *) (x_qs + txs.qs); 612 int * x_sc = (int *) (x_df + txs.dm); 613#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) 614 615 constexpr int threads_per_row = MMQ_ITER_K / (4 * QR3_K); 616 constexpr int nrows = warp_size / threads_per_row; 617 const int kqsx = threadIdx.x % threads_per_row; 618 619#pragma unroll 620 for (int i0 = 0; i0 < I; i0 += nrows*nwarps) { 621 int i = i0 + threadIdx.y*nrows + threadIdx.x/threads_per_row; 622 623 if (fallback) { 624 i = min(i, i_max); 625 } 626 627 const block_q3_K * bxi = (const block_q3_K *) x + kbx0 + i*stride; 628 629 const int x_ql_0 = get_int_b2(bxi->qs, kqsx); 630 const int x_qh_0 = get_int_b2(bxi->hmask, kqsx % (QI3_K/2)) >> (4 * (kqsx / (QI3_K/2))); 631 632#pragma unroll 633 for (int l = 0; l < QR3_K; ++l) { 634 const int k = (kqsx/8)*32 + l*8 + kqsx % 8; 635 636 const int x_ql_k = (x_ql_0 >> (2*l)) & 0x03030303; 637 const int x_qh_k = ((x_qh_0 >> l) << 2) & 0x04040404; 638 639 const int x_qs_k = __vsubss4(x_ql_k | x_qh_k, 0x04040404); 640 641#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 642 x_qs[i*sram_stride + k] = x_qs_k; 643#else 644 x_qs[i*(2*MMQ_TILE_NE_K + 1) + k] = x_qs_k; 645#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 646 } 647 } 648 649 constexpr int rows_per_warp = warp_size / 4; 650#pragma unroll 651 for (int i0 = 0; i0 < I; i0 += nwarps*rows_per_warp) { 652 int i = i0 + threadIdx.y*rows_per_warp + threadIdx.x/4; 653 654 if (fallback) { 655 i = min(i, i_max); 656 } 657 658 const block_q3_K * bxi = (const block_q3_K *) x + kbx0 + i*stride; 659 660 const int ksc = threadIdx.x % 4; 661 662 const int ksc_low = ksc % (QI3_K/8); 663 const int shift_low = 4 * (ksc / (QI3_K/8)); 664 const int sc_low = (get_int_b2(bxi->scales, ksc_low) >> shift_low) & 0x0F0F0F0F; 665 666 const int ksc_high = QI3_K/8; 667 const int shift_high = 2 * ksc; 668 const int sc_high = ((get_int_b2(bxi->scales, ksc_high) >> shift_high) << 4) & 0x30303030; 669 670 const int sc = __vsubss4(sc_low | sc_high, 0x20202020); 671 672#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 673 const int8_t * sc8 = (const int8_t *) ≻ 674 const float d = bxi->d; 675 676#pragma unroll 677 for (int l = 0; l < int(sizeof(int)); ++l) { 678 x_df[i*sram_stride + sizeof(int)*ksc + l] = d*sc8[l]; 679 } 680#else 681 x_sc[i*(MMQ_TILE_NE_K/8) + i/8 + ksc] = sc; 682#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 683 } 684 685#if !(defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)) 686#pragma unroll 687 for (int i0 = 0; i0 < I; i0 += nwarps*warp_size) { 688 int i = (i0 + threadIdx.y*warp_size + threadIdx.x) % I; 689 690 if (fallback) { 691 i = min(i, i_max); 692 } 693 694 const block_q3_K * bxi = (const block_q3_K *) x + kbx0 + i*stride; 695 696 x_df[i] = bxi->d; 697 } 698#endif // !(defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE)) || defined(AMD_WMMA_AVAILABLE) 699} 700 701static __device__ __forceinline__ int unpack_scales_q45_K(const int * scales, const int ksc) { 702 // scale arrangement after the following two lines: 703 // - ksc == 0: sc0, sc1, sc2, sc3 704 // - ksc == 1: sc4, sc5, sc6, sc7 705 // - ksc == 2: m0, m1, m2, m3 706 // - ksc == 3: m4, m5, m6, m7 707 return ((scales[(ksc%2) + (ksc!=0)] >> (4 * (ksc & (ksc/2)))) & 0x0F0F0F0F) | // lower 4 bits 708 ((scales[ksc/2] >> (2 * (ksc % 2))) & 0x30303030); // upper 2 bits 709} 710 711template <ggml_type type, int J, bool fallback> static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q4_K( 712 const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { 713 constexpr int warp_size = ggml_cuda_get_physical_warp_size(); 714 constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; 715 constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); 716 constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); 717 718#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 719 int * x_qs = (int *) x_tile; 720 half2 * x_dm = (half2 *) (x_qs + 2*MMQ_TILE_NE_K); 721#else 722 constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_Q4_K, I); 723 int * x_qs = (int *) x_tile; 724 half2 * x_dm = (half2 *) (x_qs + txs.qs); 725 int * x_sc = (int *) (x_dm + txs.dm); 726#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 727 728 constexpr int threads_per_row = MMQ_ITER_K / (4 * QR4_K); 729 constexpr int nrows = warp_size / threads_per_row; 730 const int txi = warp_size > threads_per_row ? threadIdx.x % threads_per_row : threadIdx.x; 731 732#pragma unroll 733 for (int i0 = 0; i0 < I; i0 += nrows*nwarps) { 734 int i = i0 + (nrows == 1 ? threadIdx.y : threadIdx.y*nrows + threadIdx.x/threads_per_row); 735 736 if (fallback) { 737 i = min(i, i_max); 738 } 739 740 const block_q4_K * bxi = (const block_q4_K *) x + kbx0 + i*stride; 741 const int qs0 = get_int_b4(bxi->qs, txi); 742 743#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 744 x_qs[i*sram_stride + 16*(txi/8) + txi % 8 + 0] = (qs0 >> 0) & 0x0F0F0F0F; 745 x_qs[i*sram_stride + 16*(txi/8) + txi % 8 + 8] = (qs0 >> 4) & 0x0F0F0F0F; 746#else 747 x_qs[i*(MMQ_TILE_NE_K + 1) + txi] = qs0; 748#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 749 } 750 751#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 752 constexpr int rows_per_warp = warp_size / 2; 753#pragma unroll 754 for (int i0 = 0; i0 < I; i0 += nwarps*rows_per_warp) { 755#if defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 756 // Need if on AMD instead of % because warp_size == 64 757 // This causes double work and throughput loss (MI300X) 758 // H100 loses about 100 t/s with 'if' condition over '%' 759 int i = i0 + threadIdx.y*rows_per_warp + threadIdx.x/2; 760 if (i < I) { 761#else 762 int i = (i0 + threadIdx.y*rows_per_warp + threadIdx.x/2) % I; 763 { 764#endif // defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 765 if (fallback) { 766 i = min(i, i_max); 767 } 768 769 const block_q4_K * bxi = (const block_q4_K *) x + kbx0 + i*stride; 770 771 const int * scales = (const int *) bxi->scales; 772 const int ksc = threadIdx.x % 2; 773 774 const int sc32 = unpack_scales_q45_K(scales, ksc + 0); 775 const int m32 = unpack_scales_q45_K(scales, ksc + 2); 776 777 const uint8_t * sc8 = (const uint8_t *) &sc32; 778 const uint8_t * m8 = (const uint8_t *) &m32; 779 780 const half2 dm = bxi->dm * make_half2(1.0f, -1.0f); 781 782 #pragma unroll 783 for (int l = 0; l < sizeof(int); ++l) { 784 x_dm[i*sram_stride + sizeof(int)*ksc + l] = dm*make_half2(sc8[l], m8[l]); 785 } 786 } 787 } 788#else 789#pragma unroll 790 for (int i0 = 0; i0 < I; i0 += nwarps*warp_size) { 791 int i = (i0 + threadIdx.y*warp_size + threadIdx.x) % I; 792 793 if (fallback) { 794 i = min(i, i_max); 795 } 796 797 const block_q4_K * bxi = (const block_q4_K *) x + kbx0 + i*stride; 798 799 x_dm[i] = bxi->dm; 800 } 801 constexpr int rows_per_warp = warp_size / 4; 802#pragma unroll 803 for (int i0 = 0; i0 < I; i0 += nwarps*rows_per_warp) { 804 int i = (i0 + threadIdx.y*rows_per_warp + threadIdx.x/(MMQ_TILE_NE_K/8)) % I; 805 806 if (fallback) { 807 i = min(i, i_max); 808 } 809 810 const block_q4_K * bxi = (const block_q4_K *) x + kbx0 + i*stride + (threadIdx.x % (MMQ_TILE_NE_K/8)) / (QI4_K/8); 811 812 const int * scales = (const int *) bxi->scales; 813 814 const int ksc = threadIdx.x % (MMQ_TILE_NE_K/8); 815 const int scales8 = unpack_scales_q45_K(scales, ksc); 816 817 x_sc[i*(MMQ_TILE_NE_K/8) + i/8 + ksc] = scales8; 818 } 819#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 820} 821 822template <ggml_type type, int J, bool fallback> static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q5_K( 823 const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { 824 constexpr int warp_size = ggml_cuda_get_physical_warp_size(); 825 constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; 826 constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); 827 constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); 828 829#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 830 int * x_qs = (int *) x_tile; 831 half2 * x_dm = (half2 *) (x_qs + MMQ_TILE_NE_K*2); 832#else 833 constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_Q5_K, I); 834 int * x_qs = (int *) x_tile; 835 half2 * x_dm = (half2 *) (x_qs + txs.qs); 836 int * x_sc = (int *) (x_dm + txs.dm); 837#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) 838 839 constexpr int threads_per_row = MMQ_ITER_K / (4 * QR5_K); 840 constexpr int nrows = warp_size / threads_per_row; 841 const int txi = warp_size > threads_per_row ? threadIdx.x % threads_per_row : threadIdx.x; 842 843#pragma unroll 844 for (int i0 = 0; i0 < I; i0 += nrows*nwarps) { 845 int i = i0 + (nrows == 1 ? threadIdx.y : threadIdx.y*nrows + threadIdx.x/threads_per_row); 846 847 if (fallback) { 848 i = min(i, i_max); 849 } 850 851 const block_q5_K * bxi = (const block_q5_K *) x + kbx0 + i*stride; 852 const int ky = QR5_K*txi; 853 854 const int ql = get_int_b4(bxi->qs, txi); 855 const int ql0 = (ql >> 0) & 0x0F0F0F0F; 856 const int ql1 = (ql >> 4) & 0x0F0F0F0F; 857 858 const int qh = get_int_b4(bxi->qh, txi % (QI5_K/4)); 859 const int qh0 = ((qh >> (2 * (txi / (QI5_K/4)) + 0)) << 4) & 0x10101010; 860 const int qh1 = ((qh >> (2 * (txi / (QI5_K/4)) + 1)) << 4) & 0x10101010; 861 862 const int kq0 = ky - ky % (QI5_K/2) + txi % (QI5_K/4) + 0; 863 const int kq1 = ky - ky % (QI5_K/2) + txi % (QI5_K/4) + QI5_K/4; 864 865#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 866 x_qs[i*sram_stride + kq0] = ql0 | qh0; 867 x_qs[i*sram_stride + kq1] = ql1 | qh1; 868#else 869 x_qs[i*(2*MMQ_TILE_NE_K + 1) + kq0] = ql0 | qh0; 870 x_qs[i*(2*MMQ_TILE_NE_K + 1) + kq1] = ql1 | qh1; 871#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 872 } 873 874#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 875 constexpr int rows_per_warp = warp_size / 2; 876#pragma unroll 877 for (int i0 = 0; i0 < I; i0 += nwarps*rows_per_warp) { 878#if defined(AMD_MFMA_AVAILABLE) 879 // Need if on AMD instead of % because warp_size == 64 880 // This causes double work and throughput loss (MI300X) 881 // H100 loses about 100 t/s with 'if' condition over '%' 882 int i = i0 + threadIdx.y*rows_per_warp + threadIdx.x/2; 883 if (i < I) { 884#else 885 int i = (i0 + threadIdx.y*rows_per_warp + threadIdx.x/2) % I; 886 { 887#endif // defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 888 if (fallback) { 889 i = min(i, i_max); 890 } 891 892 const block_q5_K * bxi = (const block_q5_K *) x + kbx0 + i*stride; 893 894 const int * scales = (const int *) bxi->scales; 895 const int ksc = threadIdx.x % 2; 896 897 const int sc32 = unpack_scales_q45_K(scales, ksc + 0); 898 const int m32 = unpack_scales_q45_K(scales, ksc + 2); 899 900 const uint8_t * sc8 = (const uint8_t *) &sc32; 901 const uint8_t * m8 = (const uint8_t *) &m32; 902 903 const half2 dm = bxi->dm * make_half2(1.0f, -1.0f); 904 905#pragma unroll 906 for (int l = 0; l < int(sizeof(int)); ++l) { 907 x_dm[i*sram_stride + sizeof(int)*ksc + l] = dm*make_half2(sc8[l], m8[l]); 908 } 909 } 910 } 911#else 912#pragma unroll 913 for (int i0 = 0; i0 < I; i0 += nwarps*warp_size) { 914 int i = (i0 + threadIdx.y*warp_size + threadIdx.x) % I; 915 916 if (fallback) { 917 i = min(i, i_max); 918 } 919 920 const block_q5_K * bxi = (const block_q5_K *) x + kbx0 + i*stride; 921 922 x_dm[i] = bxi->dm; 923 } 924 925 constexpr int rows_per_warp = warp_size / 4; 926#pragma unroll 927 for (int i0 = 0; i0 < I; i0 += nwarps*rows_per_warp) { 928 int i = (i0 + threadIdx.y*rows_per_warp + threadIdx.x/(MMQ_TILE_NE_K/8)) % I; 929 930 if (fallback) { 931 i = min(i, i_max); 932 } 933 934 const block_q5_K * bxi = (const block_q5_K *) x + kbx0 + i*stride; 935 936 const int * scales = (const int *) bxi->scales; 937 938 const int ksc = threadIdx.x % (MMQ_TILE_NE_K/8); 939 const int scales8 = unpack_scales_q45_K(scales, ksc); 940 941 x_sc[i*(MMQ_TILE_NE_K/8) + i/8 + ksc] = scales8; 942 } 943#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 944} 945 946template <ggml_type type, int J, bool fallback> static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q6_K( 947 const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { 948 constexpr int warp_size = ggml_cuda_get_physical_warp_size(); 949 constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; 950 constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); 951 constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); 952 953#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 954 int * x_qs = (int *) x_tile; 955 float * x_df = (float *) (x_qs + MMQ_TILE_NE_K*2); 956 int * x_sc = (int *) (x_df + MMQ_TILE_NE_K/QI6_K); 957#else 958 constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_Q6_K, I); 959 int * x_qs = (int *) x_tile; 960 float * x_df = (float *) (x_qs + txs.qs); 961 int * x_sc = (int *) (x_df + txs.dm); 962#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 963 964 constexpr int threads_per_row = MMQ_ITER_K / (4 * QR6_K); 965 constexpr int nrows = warp_size / threads_per_row; 966 const int txi = warp_size > threads_per_row ? threadIdx.x % threads_per_row : threadIdx.x; 967 968#pragma unroll 969 for (int i0 = 0; i0 < I; i0 += nrows*nwarps) { 970 int i = i0 + (nrows == 1 ? threadIdx.y : threadIdx.y*nrows + threadIdx.x/threads_per_row); 971 972 if (fallback) { 973 i = min(i, i_max); 974 } 975 976 const block_q6_K * bxi = (const block_q6_K *) x + kbx0 + i*stride; 977 978 const int ql = get_int_b2(bxi->ql, txi); 979 const int ql0 = (ql >> 0) & 0x0F0F0F0F; 980 const int ql1 = (ql >> 4) & 0x0F0F0F0F; 981 982 const int qh = get_int_b2(bxi->qh, (QI6_K/4) * (txi / (QI6_K/2)) + txi % (QI6_K/4)); 983 const int qh0 = ((qh >> ((txi & 0x08) >> 2)) << 4) & 0x30303030; 984 const int qh1 = (qh >> ((txi & 0x08) >> 2)) & 0x30303030; 985 986 const int kq0 = 2*txi - txi % (QI6_K/2) + 0; 987 const int kq1 = 2*txi - txi % (QI6_K/2) + QI6_K/2; 988 989#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 990 x_qs[i*sram_stride + kq0] = __vsubss4(ql0 | qh0, 0x20202020); 991 x_qs[i*sram_stride + kq1] = __vsubss4(ql1 | qh1, 0x20202020); 992#else 993 x_qs[i*(2*MMQ_TILE_NE_K + 1) + kq0] = __vsubss4(ql0 | qh0, 0x20202020); 994 x_qs[i*(2*MMQ_TILE_NE_K + 1) + kq1] = __vsubss4(ql1 | qh1, 0x20202020); 995#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 996 } 997 998#pragma unroll 999 for (int i0 = 0; i0 < I; i0 += nwarps*warp_size) { 1000 int i = (i0 + threadIdx.y*warp_size + threadIdx.x) % I; 1001 1002 if (fallback) { 1003 i = min(i, i_max); 1004 } 1005 1006 const block_q6_K * bxi = (const block_q6_K *) x + kbx0 + i*stride; 1007 1008#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1009 x_df[i*sram_stride] = bxi->d; 1010#else 1011 x_df[i*(MMQ_TILE_NE_K/QI6_K) + i/QI6_K] = bxi->d; 1012#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1013 } 1014 1015 constexpr int rows_per_warp = warp_size / 4; 1016#pragma unroll 1017 for (int i0 = 0; i0 < I; i0 += nwarps*rows_per_warp) { 1018 int i = (i0 + threadIdx.y*rows_per_warp + threadIdx.x/(MMQ_TILE_NE_K/8)) % I; 1019 1020 if (fallback) { 1021 i = min(i, i_max); 1022 } 1023 1024 const block_q6_K * bxi = (const block_q6_K *) x + kbx0 + i*stride + (threadIdx.x % (MMQ_TILE_NE_K/8)) / 4; 1025 1026#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1027 x_sc[i*sram_stride + threadIdx.x%4] = get_int_b2(bxi->scales, threadIdx.x % (MMQ_TILE_NE_K/8)); 1028#else 1029 x_sc[i*(MMQ_TILE_NE_K/8) + i/8 + threadIdx.x%(MMQ_TILE_NE_K/8)] = get_int_b2(bxi->scales, threadIdx.x%(QI6_K/8)); 1030#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1031 } 1032} 1033 1034// --------------------------------------------------------------------------------------------- 1035 1036template <ggml_type type, int J, bool fallback> static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_iq1_s( 1037 const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { 1038 constexpr int warp_size = ggml_cuda_get_physical_warp_size(); 1039 constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; 1040 constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); 1041 constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); 1042 1043#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1044 int * x_qs = (int *) x_tile; 1045 half2 * x_ds = (half2 *) (x_qs + MMQ_TILE_NE_K*2); 1046#else 1047 constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_IQ3_S, I); 1048 int * x_qs = (int *) x_tile; 1049 half2 * x_ds = (half2 *) (x_qs + txs.qs); 1050#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1051 1052 constexpr int threads_per_row = MMQ_ITER_K / (4 * QR1_S); 1053 constexpr int nrows = warp_size / threads_per_row; 1054 const int kqsx = threadIdx.x % threads_per_row; 1055 1056#pragma unroll 1057 for (int i0 = 0; i0 < I; i0 += nwarps * nrows) { 1058 int i = i0 + threadIdx.y*nrows + threadIdx.x/threads_per_row; 1059 1060 if (fallback) { 1061 i = min(i, i_max); 1062 } 1063 1064 const block_iq1_s * bxi = (const block_iq1_s *) x + kbx0 + i*stride; 1065 1066 const int qs_packed = get_int_b2(bxi->qs, kqsx); 1067 const uint8_t * qs = (const uint8_t *) &qs_packed; 1068 1069 const int qh = bxi->qh[kqsx]; 1070 1071 #pragma unroll 1072 for (int l = 0; l < QR1_S/2; ++l) { 1073 const int grid = iq1s_grid_gpu[qs[l] | (((qh >> (3*l)) & 0x07) << 8)]; 1074 1075 const int grid0 = (grid >> 0) & 0x0F0F0F0F; 1076 const int grid1 = (grid >> 4) & 0x0F0F0F0F; 1077 1078#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1079 x_qs[i*sram_stride + 8*kqsx + (2*l+0)] = grid0; 1080 x_qs[i*sram_stride + 8*kqsx + (2*l+1)] = grid1; 1081#else 1082 x_qs[i*(2*MMQ_TILE_NE_K + 1) + 8*kqsx + (2*l+0)] = grid0; 1083 x_qs[i*(2*MMQ_TILE_NE_K + 1) + 8*kqsx + (2*l+1)] = grid1; 1084#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1085 } 1086 1087 const float d1q = __half2float(bxi->d) * (((qh >> 11) & 0x0E) + 1); 1088 const float delta = -1.0f + IQ1S_DELTA - (qh & 0x8000) * (2.0f*IQ1S_DELTA/0x8000); 1089 1090#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1091 x_ds[i*sram_stride + kqsx] = make_half2(d1q, d1q*delta); 1092#else 1093 x_ds[i*(MMQ_TILE_NE_K/4) + i/4 + kqsx] = make_half2(d1q, d1q*delta); 1094#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1095 } 1096} 1097 1098template <ggml_type type, int J, bool fallback> static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_iq2_xxs( 1099 const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { 1100 constexpr int warp_size = ggml_cuda_get_physical_warp_size(); 1101 constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; 1102 constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); 1103 constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); 1104 1105#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1106 int * x_qs = (int *) x_tile; 1107 float * x_df = (float *) (x_qs + MMQ_TILE_NE_K*2); 1108#else 1109 constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_IQ2_XXS, I); 1110 int * x_qs = (int *) x_tile; 1111 float * x_df = (float *) (x_qs + txs.qs); 1112#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1113 1114 constexpr int threads_per_row = (MMQ_ITER_K / (4 * QR2_XXS)) / 2; 1115 constexpr int nrows = warp_size / threads_per_row; 1116 const int kqsx = warp_size > threads_per_row ? threadIdx.x % threads_per_row : threadIdx.x; 1117 1118#pragma unroll 1119 for (int i0 = 0; i0 < I; i0 += nwarps * nrows) { 1120 int i = i0 + threadIdx.y*nrows + threadIdx.x/threads_per_row; 1121 1122 if (fallback) { 1123 i = min(i, i_max); 1124 } 1125 1126 const block_iq2_xxs * bxi = (const block_iq2_xxs *) x + kbx0 + i*stride; 1127 1128 const int q2 = get_int_b2(bxi->qs, 2*kqsx+0); 1129 const uint8_t * aux8 = (const uint8_t *) &q2; 1130 const uint32_t aux32 = get_int_b2(bxi->qs, 2*kqsx+1); 1131 1132#pragma unroll 1133 for (int l = 0; l < QR2_XXS; ++l) { 1134 const uint2 grid_pos = ((const uint2*)iq2xxs_grid)[aux8[l]]; 1135 const uint32_t signs = unpack_ksigns(aux32 >> (7 * l)); 1136 1137 const int signs0 = __vcmpne4(signs & 0x08040201, 0); 1138 const int grid0 = __vsub4(grid_pos.x ^ signs0, signs0); 1139 1140 const int signs1 = __vcmpne4(signs & 0x80402010, 0); 1141 const int grid1 = __vsub4(grid_pos.y ^ signs1, signs1); 1142 1143#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1144 x_qs[i*sram_stride + 8*kqsx + (2*l + 0)] = grid0; 1145 x_qs[i*sram_stride + 8*kqsx + (2*l + 1)] = grid1; 1146#else 1147 x_qs[i*(2*MMQ_TILE_NE_K + 1) + 8*kqsx + (2*l + 0)] = grid0; 1148 x_qs[i*(2*MMQ_TILE_NE_K + 1) + 8*kqsx + (2*l + 1)] = grid1; 1149#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1150 } 1151 1152 const int ls = aux32 >> 27 | 1; // (scale * 2 + 1) 1153 const float d = bxi->d; 1154#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1155 x_df[i*sram_stride + kqsx] = d * ls / 8; // (d * scale + d / 2) / 4 1156#else 1157 x_df[i*(MMQ_TILE_NE_K/4) + i/4 + kqsx] = d * ls / 8; // (d * scale + d / 2) / 4 1158#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1159 } 1160} 1161 1162template <ggml_type type, int J, bool fallback> static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_iq2_xs( 1163 const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { 1164 constexpr int warp_size = ggml_cuda_get_physical_warp_size(); 1165 constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; 1166 constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); 1167 constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); 1168 1169#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1170 int * x_qs = (int *) x_tile; 1171 float * x_df = (float *) (x_qs + MMQ_TILE_NE_K*2); 1172#else 1173 constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_IQ2_XS, I); 1174 int * x_qs = (int *) x_tile; 1175 float * x_df = (float *) (x_qs + txs.qs); 1176#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1177 1178 constexpr int threads_per_row = (MMQ_ITER_K / (4 * QR2_XS)) / 2; 1179 constexpr int nrows = warp_size / threads_per_row; 1180 const int kqsx = threadIdx.x % threads_per_row; 1181 1182#pragma unroll 1183 for (int i0 = 0; i0 < I; i0 += nwarps * nrows) { 1184 int i = i0 + threadIdx.y*nrows + threadIdx.x/threads_per_row; 1185 1186 if (fallback) { 1187 i = min(i, i_max); 1188 } 1189 1190 const block_iq2_xs * bxi = (const block_iq2_xs *) x + kbx0 + i*stride; 1191 1192 const int2 q2_packed = make_int2(get_int_b2(bxi->qs, 2*kqsx+0), get_int_b2(bxi->qs, 2*kqsx+1)); 1193 const uint16_t * q2 = (const uint16_t *) &q2_packed; 1194 1195 #pragma unroll 1196 for (int l = 0; l < QR2_XS; ++l) { 1197 const uint2 grid_pos = ((const uint2*)iq2xs_grid)[q2[l] & 0x1FF]; 1198 const uint32_t signs = unpack_ksigns(q2[l] >> 9); 1199 1200 const int signs0 = __vcmpne4(signs & 0x08040201, 0); 1201 const int grid_l = __vsub4(grid_pos.x ^ signs0, signs0); 1202 1203 const int signs1 = __vcmpne4(signs & 0x80402010, 0); 1204 const int grid_h = __vsub4(grid_pos.y ^ signs1, signs1); 1205 1206#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1207 x_qs[i*sram_stride + 8*kqsx + (2*l + 0)] = grid_l; 1208 x_qs[i*sram_stride + 8*kqsx + (2*l + 1)] = grid_h; 1209#else 1210 x_qs[i*(2*MMQ_TILE_NE_K + 1) + 8*kqsx + (2*l + 0)] = grid_l; 1211 x_qs[i*(2*MMQ_TILE_NE_K + 1) + 8*kqsx + (2*l + 1)] = grid_h; 1212#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1213 } 1214 1215 const int ls = bxi->scales[kqsx]; 1216 const float d = bxi->d; 1217#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1218 x_df[i*sram_stride + 2*kqsx+0] = ((ls & 0x0F)*d + d/2)/4; 1219 x_df[i*sram_stride + 2*kqsx+1] = ((ls >> 4)*d + d/2)/4; 1220#else 1221 x_df[i*(2*MMQ_TILE_NE_K*2/QI8_0) + i/(QI8_0/4) + 2*kqsx+0] = ((ls & 0x0F)*d + d/2)/4; 1222 x_df[i*(2*MMQ_TILE_NE_K*2/QI8_0) + i/(QI8_0/4) + 2*kqsx+1] = ((ls >> 4)*d + d/2)/4; 1223#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1224 } 1225} 1226 1227template <ggml_type type, int J, bool fallback> static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_iq2_s( 1228 const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { 1229 constexpr int warp_size = ggml_cuda_get_physical_warp_size(); 1230 constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; 1231 constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); 1232 constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); 1233 1234#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1235 int * x_qs = (int *) x_tile; 1236 float * x_df = (float *) (x_qs + MMQ_TILE_NE_K*2); 1237#else 1238 constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_IQ2_S, I); 1239 int * x_qs = (int *) x_tile; 1240 float * x_df = (float *) (x_qs + txs.qs); 1241#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1242 constexpr int threads_per_row = (MMQ_ITER_K / (4 * QR2_S)) / 2; 1243 constexpr int nrows = warp_size / threads_per_row; 1244 const int kqsx = threadIdx.x % threads_per_row; 1245 1246#pragma unroll 1247 for (int i0 = 0; i0 < I; i0 += nwarps * nrows) { 1248 int i = i0 + threadIdx.y*nrows + threadIdx.x/threads_per_row; 1249 1250 if (fallback) { 1251 i = min(i, i_max); 1252 } 1253 1254 const block_iq2_s * bxi = (const block_iq2_s *) x + kbx0 + i*stride; 1255 1256 const int qs_packed = get_int_b2(bxi->qs, kqsx); 1257 const uint8_t * qs = (const uint8_t *) &qs_packed; 1258 1259 const int qh = bxi->qh[kqsx]; 1260 1261 const int signs_packed_32 = get_int_b2(bxi->qs, QK_K/32 + kqsx); 1262 const uint8_t * signs_packed_8 = (const uint8_t *) &signs_packed_32; 1263 1264#pragma unroll 1265 for (int l = 0; l < QR2_S; ++l) { 1266 const int * grid_pos = (const int *)(iq2s_grid + (qs[l] | ((qh << (8-2*l)) & 0x300))); 1267 1268 const int signs0 = __vcmpne4(((signs_packed_8[l] & 0x03) << 7) | ((signs_packed_8[l] & 0x0C) << 21), 0x00000000); 1269 const int signs1 = __vcmpne4(((signs_packed_8[l] & 0x30) << 3) | ((signs_packed_8[l] & 0xC0) << 17), 0x00000000); 1270 1271 const int grid_l = __vsub4(grid_pos[0] ^ signs0, signs0); 1272 const int grid_h = __vsub4(grid_pos[1] ^ signs1, signs1); 1273 1274#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1275 x_qs[i*sram_stride + 8*kqsx + (2*l + 0)] = grid_l; 1276 x_qs[i*sram_stride + 8*kqsx + (2*l + 1)] = grid_h; 1277#else 1278 x_qs[i*(2*MMQ_TILE_NE_K + 1) + 8*kqsx + (2*l + 0)] = grid_l; 1279 x_qs[i*(2*MMQ_TILE_NE_K + 1) + 8*kqsx + (2*l + 1)] = grid_h; 1280#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1281 } 1282 1283 const int ls = bxi->scales[kqsx]; 1284 const float d = bxi->d; 1285#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1286 x_df[i*sram_stride + 2*kqsx+0] = ((ls & 0x0F)*d + d/2)/4; 1287 x_df[i*sram_stride + 2*kqsx+1] = ((ls >> 4)*d + d/2)/4; 1288#else 1289 x_df[i*(2*MMQ_TILE_NE_K*2/QI8_0) + i/(QI8_0/4) + 2*kqsx+0] = ((ls & 0x0F)*d + d/2)/4; 1290 x_df[i*(2*MMQ_TILE_NE_K*2/QI8_0) + i/(QI8_0/4) + 2*kqsx+1] = ((ls >> 4)*d + d/2)/4; 1291#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1292 } 1293} 1294 1295template <ggml_type type, int J, bool fallback> static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_iq3_xxs( 1296 const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { 1297 constexpr int warp_size = ggml_cuda_get_physical_warp_size(); 1298 constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; 1299 constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); 1300 constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); 1301 1302#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1303 int * x_qs = (int *) x_tile; 1304 float * x_df = (float *) (x_qs + MMQ_TILE_NE_K*2); 1305#else 1306 constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_IQ3_XXS, I); 1307 int * x_qs = (int *) x_tile; 1308 float * x_df = (float *) (x_qs + txs.qs); 1309#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1310 1311 constexpr int threads_per_row = (MMQ_ITER_K / (4 * QR3_XXS)) / 2; 1312 constexpr int nrows = warp_size / threads_per_row; 1313 const int kqsx = threadIdx.x % threads_per_row; 1314 1315#pragma unroll 1316 for (int i0 = 0; i0 < I; i0 += nwarps * nrows) { 1317 int i = i0 + threadIdx.y*nrows + threadIdx.x/threads_per_row; 1318 1319 if (fallback) { 1320 i = min(i, i_max); 1321 } 1322 1323 const block_iq3_xxs * bxi = (const block_iq3_xxs *) x + kbx0 + i*stride; 1324 1325 const int2 q3_packed = make_int2(get_int_b2(bxi->qs, 2*kqsx+0), get_int_b2(bxi->qs, 2*kqsx+1)); 1326 const uint8_t * q3 = (const uint8_t *) &q3_packed; 1327 const uint32_t aux32 = get_int_b2(bxi->qs, QK_K/16 + kqsx); 1328 1329#pragma unroll 1330 for (int l = 0; l < QR3_XXS; ++l) { 1331 const int2 grid_pos = make_int2(iq3xxs_grid[q3[2*l+0]], iq3xxs_grid[q3[2*l+1]]); 1332 const uint32_t signs = unpack_ksigns(aux32 >> (7*l)); 1333 1334 const int signs0 = __vcmpne4(signs & 0x08040201, 0); 1335 const int grid_l = __vsub4(grid_pos.x ^ signs0, signs0); 1336 1337 const int signs1 = __vcmpne4(signs & 0x80402010, 0); 1338 const int grid_h = __vsub4(grid_pos.y ^ signs1, signs1); 1339 1340#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1341 x_qs[i*sram_stride + 8*kqsx + (2*l + 0)] = grid_l; 1342 x_qs[i*sram_stride + 8*kqsx + (2*l + 1)] = grid_h; 1343#else 1344 x_qs[i*(2*MMQ_TILE_NE_K + 1) + 8*kqsx + (2*l + 0)] = grid_l; 1345 x_qs[i*(2*MMQ_TILE_NE_K + 1) + 8*kqsx + (2*l + 1)] = grid_h; 1346#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1347 } 1348 1349 const int ls = aux32 >> 28; 1350 const float d = bxi->d; 1351#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1352 x_df[i*sram_stride + kqsx] = (ls*d + d/2)/2; 1353#else 1354 x_df[i*(MMQ_TILE_NE_K/4) + i/4 + kqsx] = (ls*d + d/2)/2; 1355#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1356 } 1357} 1358 1359template <ggml_type type, int J, bool fallback> static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_iq3_s( 1360 const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { 1361 constexpr int warp_size = ggml_cuda_get_physical_warp_size(); 1362 constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; 1363 constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); 1364 constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); 1365 1366#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1367 int * x_qs = (int *) x_tile; 1368 float * x_df = (float *) (x_qs + MMQ_TILE_NE_K*2); 1369#else 1370 constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_IQ3_S, I); 1371 int * x_qs = (int *) x_tile; 1372 float * x_df = (float *) (x_qs + txs.qs); 1373#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1374 1375 constexpr int threads_per_row = (MMQ_ITER_K / (4 * QR3_S)) / 2; 1376 constexpr int nrows = warp_size / threads_per_row; 1377 const int kqsx = threadIdx.x % threads_per_row; 1378 1379#pragma unroll 1380 for (int i0 = 0; i0 < I; i0 += nwarps * nrows) { 1381 int i = i0 + threadIdx.y*nrows + threadIdx.x/threads_per_row; 1382 1383 if (fallback) { 1384 i = min(i, i_max); 1385 } 1386 1387 const block_iq3_s * bxi = (const block_iq3_s *) x + kbx0 + i*stride; 1388 1389 const int2 qs_packed = make_int2(get_int_b2(bxi->qs, 2*kqsx+0), get_int_b2(bxi->qs, 2*kqsx+1)); 1390 const uint8_t * qs = (const uint8_t *) &qs_packed; 1391 1392 const int qh = bxi->qh[kqsx]; 1393 1394 const int signs_packed_32 = get_int_b2(bxi->signs, kqsx); 1395 const uint8_t * signs_packed_8 = (const uint8_t *) &signs_packed_32; 1396 1397#pragma unroll 1398 for (int l = 0; l < QR3_S; ++l) { 1399 const int2 grid_pos = make_int2( 1400 iq3s_grid[qs[2*l+0] | ((qh << (8 - 2*l)) & 0x100)], 1401 iq3s_grid[qs[2*l+1] | ((qh << (7 - 2*l)) & 0x100)]); 1402 1403 const int signs0 = __vcmpne4(((signs_packed_8[l] & 0x03) << 7) | ((signs_packed_8[l] & 0x0C) << 21), 0x00000000); 1404 const int signs1 = __vcmpne4(((signs_packed_8[l] & 0x30) << 3) | ((signs_packed_8[l] & 0xC0) << 17), 0x00000000); 1405 1406 const int grid_l = __vsub4(grid_pos.x ^ signs0, signs0); 1407 const int grid_h = __vsub4(grid_pos.y ^ signs1, signs1); 1408 1409#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1410 x_qs[i*sram_stride + 8*kqsx + (2*l+0)] = grid_l; 1411 x_qs[i*sram_stride + 8*kqsx + (2*l+1)] = grid_h; 1412#else 1413 x_qs[i*(2*MMQ_TILE_NE_K + 1) + 8*kqsx + (2*l+0)] = grid_l; 1414 x_qs[i*(2*MMQ_TILE_NE_K + 1) + 8*kqsx + (2*l+1)] = grid_h; 1415#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1416 } 1417 1418 const int ls = 1 + 2*((bxi->scales[kqsx/2] >> (((2*kqsx) << 1) & 0x04)) & 0x0F); 1419 const float d = bxi->d; 1420#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1421 x_df[i*sram_stride + kqsx] = ls*d; 1422#else 1423 x_df[i*(MMQ_TILE_NE_K/4) + i/4 + kqsx] = ls*d; 1424#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1425 } 1426} 1427 1428template <ggml_type type, int J, bool fallback> static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_iq4_xs( 1429 const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { 1430 constexpr int warp_size = ggml_cuda_get_physical_warp_size(); 1431 constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; 1432 constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); 1433 constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); 1434 1435#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1436 int * x_qs = (int *) x_tile; 1437 float * x_df = (float *) (x_qs + MMQ_TILE_NE_K*2); 1438#else 1439 constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_IQ4_XS, I); 1440 int * x_qs = (int *) x_tile; 1441 float * x_df = (float *) (x_qs + txs.qs); 1442#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1443 1444 constexpr int threads_per_row = MMQ_ITER_K / (4 * QR4_XS); 1445 constexpr int nrows = warp_size / threads_per_row; 1446 const int kqsx = threadIdx.x % threads_per_row; 1447 1448#pragma unroll 1449 for (int i0 = 0; i0 < I; i0 += nrows*nwarps) { 1450 int i = i0 + (nrows == 1 ? threadIdx.y : threadIdx.y*nrows + threadIdx.x/threads_per_row); 1451 1452 if (fallback) { 1453 i = min(i, i_max); 1454 } 1455 1456 const block_iq4_xs * bxi = (const block_iq4_xs *) x + kbx0 + i*stride; 1457 1458 const int aux_q4 = get_int_b4(bxi->qs, kqsx); 1459 const int2 v = get_int_from_table_16(aux_q4, kvalues_iq4nl); 1460 const int k0 = 8 * (kqsx / 4) + kqsx % 4; 1461 1462#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1463 x_qs[i*sram_stride + k0 + 0] = v.x; 1464 x_qs[i*sram_stride + k0 + 4] = v.y; 1465#else 1466 x_qs[i*(2*MMQ_TILE_NE_K + 1) + k0 + 0] = v.x; 1467 x_qs[i*(2*MMQ_TILE_NE_K + 1) + k0 + 4] = v.y; 1468#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1469 } 1470 1471 constexpr int rows_per_warp = warp_size / 8; 1472#pragma unroll 1473 for (int i0 = 0; i0 < I; i0 += nwarps * rows_per_warp) { 1474 int i = i0 + threadIdx.y * rows_per_warp + threadIdx.x / (MMQ_TILE_NE_K/4); 1475 1476 if (fallback) { 1477 i = min(i, i_max); 1478 } 1479 1480 const block_iq4_xs * bxi = (const block_iq4_xs *) x + kbx0 + i*stride; 1481 1482 const float d = __half2float(bxi->d); 1483 1484 const int ls = ((bxi->scales_l[(threadIdx.x % 8)/2] >> (4*(threadIdx.x % 2))) & 0x0F) 1485 | (((bxi->scales_h >> (2*(threadIdx.x % 8))) & 0x03) << 4); 1486 1487#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1488 x_df[i*sram_stride + threadIdx.x % 8] = d * (ls - 32); 1489#else 1490 x_df[i*(MMQ_TILE_NE_K/4) + i/4 + threadIdx.x % 8] = d * (ls - 32); 1491#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1492 } 1493} 1494 1495template <ggml_type type, int J, bool fallback> static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_iq4_nl( 1496 const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { 1497 constexpr int warp_size = ggml_cuda_get_physical_warp_size(); 1498 constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; 1499 constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); 1500 constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); 1501 1502#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1503 int * x_qs = (int *) x_tile; 1504 float * x_df = (float *) (x_qs + MMQ_TILE_NE_K*2); 1505#else 1506 constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_IQ4_NL, I); 1507 int * x_qs = (int *) x_tile; 1508 float * x_df = (float *) (x_qs + txs.qs); 1509#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1510 1511 constexpr int threads_per_row = MMQ_ITER_K / (4 * QR4_NL); 1512 constexpr int nrows = warp_size / threads_per_row; 1513 const int txi = warp_size > threads_per_row ? threadIdx.x % threads_per_row : threadIdx.x; 1514 const int kbx = txi / QI4_NL; 1515 const int kqsx = txi % QI4_NL; 1516 1517#pragma unroll 1518 for (int i0 = 0; i0 < I; i0 += nrows*nwarps) { 1519 int i = i0 + (nrows == 1 ? threadIdx.y : threadIdx.y*nrows + threadIdx.x/threads_per_row); 1520 1521 if (fallback) { 1522 i = min(i, i_max); 1523 } 1524 1525 const block_iq4_nl * bxi = (const block_iq4_nl *) x + kbx0 + i*stride + kbx; 1526 1527 const int aux_q4 = get_int_b2(bxi->qs, kqsx); 1528 const int2 v = get_int_from_table_16(aux_q4, kvalues_iq4nl); 1529 const int k0 = kbx * (2 * QI4_NL) + kqsx; 1530 1531#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1532 x_qs[i*sram_stride + k0 + 0] = v.x; 1533 x_qs[i*sram_stride + k0 + QI4_NL] = v.y; 1534#else 1535 x_qs[i*(2*MMQ_TILE_NE_K + 1) + k0 + 0] = v.x; 1536 x_qs[i*(2*MMQ_TILE_NE_K + 1) + k0 + QI4_NL] = v.y; 1537#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1538 } 1539 1540 constexpr int blocks_per_tile_x_row = MMQ_TILE_NE_K / QI4_NL; 1541 constexpr int rows_per_warp = warp_size / blocks_per_tile_x_row; 1542 const int kbxd = threadIdx.x % blocks_per_tile_x_row; 1543 1544#pragma unroll 1545 for (int i0 = 0; i0 < I; i0 += nwarps * rows_per_warp) { 1546 int i = i0 + threadIdx.y * rows_per_warp + threadIdx.x / blocks_per_tile_x_row; 1547 1548 if (fallback) { 1549 i = min(i, i_max); 1550 } 1551 1552 const block_iq4_nl * bxi = (const block_iq4_nl *) x + kbx0 + i*stride + kbxd; 1553 1554#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1555 x_df[i*sram_stride + kbxd] = __half2float(bxi->d); 1556#else 1557 x_df[i*(MMQ_TILE_NE_K/QI4_NL) + i/QI4_NL + kbxd] = __half2float(bxi->d); 1558#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1559 } 1560} 1561 1562// --------------------------------------------------------------------------------------------- 1563 1564template <ggml_type type, int J, bool fallback> static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_mxfp4( 1565 const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { 1566 constexpr int warp_size = ggml_cuda_get_physical_warp_size(); 1567 constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; 1568 constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); 1569 constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); 1570 1571#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1572 int * x_qs = (int *) x_tile; 1573 float * x_df = (float *) (x_qs + MMQ_TILE_NE_K*2); 1574#else 1575 constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_MXFP4, I); 1576 int * x_qs = (int *) x_tile; 1577 float * x_df = (float *) (x_qs + txs.qs); 1578#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1579 1580 constexpr int threads_per_row = MMQ_ITER_K / (4 * QR_MXFP4); 1581 constexpr int nrows = warp_size / threads_per_row; 1582 const int txi = warp_size > threads_per_row ? threadIdx.x % threads_per_row : threadIdx.x; 1583 const int kbx = txi / QI_MXFP4; 1584 const int kqsx = txi % QI_MXFP4; 1585 1586#pragma unroll 1587 for (int i0 = 0; i0 < I; i0 += nrows*nwarps) { 1588 int i = i0 + (nrows == 1 ? threadIdx.y : threadIdx.y*nrows + threadIdx.x/threads_per_row); 1589 1590 if (fallback) { 1591 i = min(i, i_max); 1592 } 1593 1594 const block_mxfp4 * bxi = (const block_mxfp4 *) x + kbx0 + i*stride + kbx; 1595 1596 const int aux_q4 = get_int_b1(bxi->qs, kqsx); 1597 const int2 v = get_int_from_table_16(aux_q4, kvalues_mxfp4); 1598 const int k0 = kbx * (2 * QI_MXFP4) + kqsx; 1599 1600#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1601 x_qs[i*sram_stride + k0 + 0] = v.x; 1602 x_qs[i*sram_stride + k0 + QI_MXFP4] = v.y; 1603#else 1604 x_qs[i*(2*MMQ_TILE_NE_K + 1) + k0 + 0] = v.x; 1605 x_qs[i*(2*MMQ_TILE_NE_K + 1) + k0 + QI_MXFP4] = v.y; 1606#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1607 } 1608 1609 constexpr int blocks_per_tile_x_row = MMQ_TILE_NE_K / QI_MXFP4; 1610 constexpr int rows_per_warp = warp_size / blocks_per_tile_x_row; 1611 const int kbxd = threadIdx.x % blocks_per_tile_x_row; 1612 1613#pragma unroll 1614 for (int i0 = 0; i0 < I; i0 += nwarps * rows_per_warp) { 1615 int i = i0 + threadIdx.y * rows_per_warp + threadIdx.x / blocks_per_tile_x_row; 1616 1617 if (fallback) { 1618 i = min(i, i_max); 1619 } 1620 1621 const block_mxfp4 * bxi = (const block_mxfp4 *) x + kbx0 + i*stride + kbxd; 1622 1623#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1624 x_df[i*sram_stride + kbxd] = ggml_cuda_e8m0_to_fp32(bxi->e)*0.5f; 1625#else 1626 x_df[i*(MMQ_TILE_NE_K/QI_MXFP4) + i/QI_MXFP4 + kbxd] = ggml_cuda_e8m0_to_fp32(bxi->e)*0.5f; 1627#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1628 } 1629} 1630 1631template <ggml_type type, int J, bool fallback> static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_mxfp4_fp4( 1632 const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { 1633 constexpr int warp_size = ggml_cuda_get_physical_warp_size(); 1634 constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; 1635 constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); 1636 constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); 1637 1638 int * x_qs = (int *) x_tile; 1639 uint32_t * x_sc = (uint32_t *) (x_qs + 2 * MMQ_TILE_NE_K); 1640 1641 const int txi = threadIdx.x; 1642 1643 constexpr int iter_k = ggml_cuda_mmq_get_K_vram(type, J, fallback); 1644 1645 constexpr int threads_per_row = iter_k / QK_MXFP4; // each thread processes 1 block 1646 constexpr int rows_per_warp = warp_size / threads_per_row; 1647 const int kbx = txi % threads_per_row; 1648 const int row_in_warp = txi / threads_per_row; 1649 1650#pragma unroll 1651 for (int i0 = 0; i0 < I; i0 += rows_per_warp * nwarps) { 1652 int i = i0 + threadIdx.y * rows_per_warp + row_in_warp; 1653 1654 if constexpr (fallback) { 1655 i = min(i, i_max); 1656 } 1657 1658 const block_mxfp4 * bxi = (const block_mxfp4 *) x + kbx0 + i * stride + kbx; 1659 1660 // quantize_mxfp4_mmq permutes nibbles to match the quantized format 1661 const int k0 = kbx * 4; 1662 memcpy(x_qs + i*sram_stride + k0, bxi->qs, 16); 1663 1664 // Load E8M0 scales: pack 2 consecutive scales into one uint32 1665 if (kbx % 2 == 0) { 1666 uint32_t e = bxi->e; 1667 e |= ((bxi + 1)->e << 8); 1668 x_sc[i*sram_stride + kbx / 2] = e; 1669 } 1670 } 1671} 1672 1673template <ggml_type type, int J, bool fallback> static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_nvfp4( 1674 const char * __restrict__ x, int * __restrict__ x_tile, const int kb0, const int i_max, const int stride) { 1675 constexpr int warp_size = ggml_cuda_get_physical_warp_size(); 1676 constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; 1677 constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); 1678 constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); 1679 1680#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1681 int * x_qs = (int *) x_tile; 1682 float * x_df = (float *) (x_qs + MMQ_TILE_NE_K*2); 1683#else 1684 constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_NVFP4, I); 1685 int * x_qs = (int *) x_tile; 1686 float * x_df = (float *) (x_qs + txs.qs); 1687#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1688 1689 constexpr int threads_per_row = MMQ_ITER_K / QK_NVFP4; 1690 constexpr int rows_per_warp = warp_size / threads_per_row; 1691 const int kbx = threadIdx.x % threads_per_row; 1692 const int row_in_warp = threadIdx.x / threads_per_row; 1693 1694#pragma unroll 1695 for (int i0 = 0; i0 < I; i0 += rows_per_warp * nwarps) { 1696 int i = i0 + threadIdx.y * rows_per_warp + row_in_warp; 1697 1698 if constexpr (fallback) { 1699 i = min(i, i_max); 1700 } 1701 1702 const block_nvfp4 * bxi = (const block_nvfp4 *) x + kb0 + i * stride + kbx; 1703 const uint32_t * __restrict__ src_qs = reinterpret_cast<const uint32_t *>(bxi->qs); 1704 const int kqs = 16 * kbx; 1705 const int ksc = 4 * kbx; 1706 1707#pragma unroll 1708 for (int sub = 0; sub < QK_NVFP4 / QK_NVFP4_SUB; ++sub) { 1709 const int2 q0 = get_int_from_table_16(src_qs[2 * sub + 0], kvalues_mxfp4); 1710 const int2 q1 = get_int_from_table_16(src_qs[2 * sub + 1], kvalues_mxfp4); 1711 1712#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1713 x_qs[i*sram_stride + kqs + 4 * sub + 0] = q0.x; 1714 x_qs[i*sram_stride + kqs + 4 * sub + 1] = q1.x; 1715 x_qs[i*sram_stride + kqs + 4 * sub + 2] = q0.y; 1716 x_qs[i*sram_stride + kqs + 4 * sub + 3] = q1.y; 1717 x_df[i*sram_stride + ksc + sub] = ggml_cuda_ue4m3_to_fp32(bxi->d[sub]); 1718#else 1719 x_qs[i * (2 * MMQ_TILE_NE_K + 1) + kqs + 4 * sub + 0] = q0.x; 1720 x_qs[i * (2 * MMQ_TILE_NE_K + 1) + kqs + 4 * sub + 1] = q1.x; 1721 x_qs[i * (2 * MMQ_TILE_NE_K + 1) + kqs + 4 * sub + 2] = q0.y; 1722 x_qs[i * (2 * MMQ_TILE_NE_K + 1) + kqs + 4 * sub + 3] = q1.y; 1723 x_df[i * (2 * MMQ_TILE_NE_K * 2 / QI_NVFP4) + i / (QK_NVFP4_SUB / QI_NVFP4) + ksc + sub] = ggml_cuda_ue4m3_to_fp32(bxi->d[sub]); 1724#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) 1725 } 1726 } 1727} 1728 1729template <ggml_type type, int J, bool fallback> static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_nvfp4_nvfp4( 1730 const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { 1731 constexpr int warp_size = ggml_cuda_get_physical_warp_size(); 1732 constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; 1733 constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); 1734 constexpr int iter_k = ggml_cuda_mmq_get_K_vram(type, J, fallback); 1735 constexpr int threads_per_row = iter_k / QK_NVFP4; // each thread processes 1 block 1736 constexpr int rows_per_warp = warp_size / threads_per_row; 1737 constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); 1738 1739 uint32_t * x_u32 = (uint32_t *) x_tile; 1740 1741 const int txi = threadIdx.x; 1742 const int kbx = txi % threads_per_row; 1743 const int row_in_warp = txi / threads_per_row; 1744 1745 const block_nvfp4 * bxi_base = (const block_nvfp4 *) x + kbx0 + kbx; 1746 uint32_t * x_u32_scale = x_u32 + 64 + kbx; 1747 1748#pragma unroll 1749 for (int i0 = 0; i0 < I; i0 += rows_per_warp * nwarps) { 1750 int i = i0 + threadIdx.y * rows_per_warp + row_in_warp; 1751 1752 if constexpr (fallback) { 1753 i = min(i, i_max); 1754 } 1755 1756 const block_nvfp4 * bxi = bxi_base + i * stride; 1757 1758 const uint32_t * src_qs = reinterpret_cast<const uint32_t *>(bxi->qs); 1759 1760#pragma unroll 1761 for (int sub = 0; sub < QK_NVFP4 / QK_NVFP4_SUB; ++sub) { 1762 x_u32[i*sram_stride + 8*kbx + 2 * sub + 0] = src_qs[2 * sub + 0]; 1763 x_u32[i*sram_stride + 8*kbx + 2 * sub + 1] = src_qs[2 * sub + 1]; 1764 } 1765 1766 x_u32_scale[i*sram_stride] = get_int_b4(bxi->d, 0); 1767 } 1768}