char-fork/llama.cpp

git clone https://git.t4t.associates/char-fork/llama.cpp

Pasha KhosraviCUDA: add Q2_0 support (#25707)9b2a08881

char/aliceai-t5-moe
brancheschar/aliceai-t5-moeall refs →
25.4 KiB705 linesraw
1#include "convert.cuh"
2#include "dequantize.cuh"
3
4#include <cstdint>
5
6#define CUDA_Q8_0_NE_ALIGN 2048
7
8template <int qk, int qr, dequantize_kernel_t dequantize_kernel, typename dst_t>
9static __global__ void dequantize_block(const void * __restrict__ vx, dst_t * __restrict__ y,
10        const int64_t ne00, const int64_t ne01,
11        const int64_t ne0203, const uint3 ne02,
12        const int64_t s01, const int64_t s02, const int64_t s03) {
13    const int64_t i00 = 2 * (int64_t(blockDim.x)*blockIdx.x + threadIdx.x);
14
15    if (i00 >= ne00) {
16        return;
17    }
18
19    for (int64_t i01 = blockIdx.y; i01 < ne01; i01 += gridDim.y) {
20        for (int64_t i0203 = blockIdx.z; i0203 < ne0203; i0203 += gridDim.z) {
21            const uint2 dm = fast_div_modulo((uint32_t)i0203, ne02);
22            const int64_t i02 = dm.y;
23            const int64_t i03 = dm.x;
24
25            const int64_t ibx0 = i03*s03 + i02*s02 + i01*s01;
26
27            const int64_t ib = ibx0 + i00/qk; // block index
28            const int64_t iqs = (i00%qk)/qr; // quant index
29            const int64_t iybs = i00 - i00%qk; // y block start index
30            const int64_t y_offset = qr == 1 ? 1 : qk/2;
31
32            // dequantize
33            float2 v;
34            dequantize_kernel(vx, ib, iqs, v);
35
36            const int64_t iy0 = (i0203*ne01 + i01)*ne00 + iybs + iqs;
37            y[iy0 + 0]        = ggml_cuda_cast<dst_t>(v.x);
38            y[iy0 + y_offset] = ggml_cuda_cast<dst_t>(v.y);
39        }
40    }
41}
42
43template <bool need_check>
44static __global__ void dequantize_block_q8_0_f16(const void * __restrict__ vx, half * __restrict__ y, const int64_t k) {
45#if __CUDA_ARCH__ >= GGML_CUDA_CC_PASCAL
46    constexpr int nint = CUDA_Q8_0_NE_ALIGN/sizeof(int) + WARP_SIZE;
47
48    const int64_t   i0 = CUDA_Q8_0_NE_ALIGN*blockIdx.x;
49    const int * x0 = ((int *) vx) + blockIdx.x * nint;
50    half2 * y2 = (half2 *) (y + i0);
51
52    __shared__ int vals[nint];
53
54#pragma unroll
55    for (int ix0 = 0; ix0 < nint; ix0 += WARP_SIZE) {
56        if (need_check && i0*sizeof(block_q8_0)/QK8_0 + sizeof(int)*(ix0 + threadIdx.x) >= k*sizeof(block_q8_0)/QK8_0) {
57            break;
58        }
59
60        const int ix = ix0 + threadIdx.x;
61        vals[ix] = x0[ix];
62    }
63
64    __syncthreads();
65
66#pragma unroll
67    for (int iy = 0; iy < CUDA_Q8_0_NE_ALIGN; iy += 2*WARP_SIZE) {
68        if (need_check && i0 + iy + 2*threadIdx.x >= k) {
69            return;
70        }
71
72        const half * b0 = ((const half  *) vals) + (sizeof(block_q8_0)/sizeof(half)) * ((iy + 2*threadIdx.x)/QK8_0);
73        const half    d = *b0;
74        const char2  qs = ((const char2 *) (b0 + 1))[threadIdx.x % (QK8_0/2)];
75
76        y2[iy/2 + threadIdx.x] = __hmul2(make_half2(qs.x, qs.y), __half2half2(d));
77    }
78#else
79    GGML_UNUSED_VARS(vx, y, k);
80    NO_DEVICE_CODE;
81#endif // __CUDA_ARCH__ >= GGML_CUDA_CC_PASCAL
82}
83
84template<typename dst_t>
85static __global__ void dequantize_block_q4_0(const void * __restrict__ vx, dst_t * __restrict__ yy, int nb32) {
86
87    const int64_t i = blockIdx.x;
88
89    // assume 32 threads
90    const int64_t tid = threadIdx.x;
91    const int64_t il  = tid/8;
92    const int64_t ir  = tid%8;
93    const int64_t ib = 8*i + ir;
94    if (ib >= nb32) {
95        return;
96    }
97
98    dst_t * y = yy + 256*i + 32*ir + 4*il;
99
100    const block_q4_0 * x = (const block_q4_0 *)vx + ib;
101    const float d = __half2float(x->d);
102    const float dm = -8*d;
103
104    const uint8_t * q = x->qs + 4*il;
105
106    for (int l = 0; l < 4; ++l) {
107        y[l+ 0] = ggml_cuda_cast<dst_t>(d * (q[l] & 0xF) + dm);
108        y[l+16] = ggml_cuda_cast<dst_t>(d * (q[l] >>  4) + dm);
109    }
110}
111
112template<typename dst_t>
113static __global__ void dequantize_block_q4_1(const void * __restrict__ vx, dst_t * __restrict__ yy, int nb32) {
114
115    const int64_t i = blockIdx.x;
116
117    // assume 32 threads
118    const int64_t tid = threadIdx.x;
119    const int64_t il  = tid/8;
120    const int64_t ir  = tid%8;
121    const int64_t ib = 8*i + ir;
122    if (ib >= nb32) {
123        return;
124    }
125
126    dst_t * y = yy + 256*i + 32*ir + 4*il;
127
128    const block_q4_1 * x = (const block_q4_1 *)vx + ib;
129    const float2 d = __half22float2(x->dm);
130
131    const uint8_t * q = x->qs + 4*il;
132
133    for (int l = 0; l < 4; ++l) {
134        y[l+ 0] = ggml_cuda_cast<dst_t>(d.x * (q[l] & 0xF) + d.y);
135        y[l+16] = ggml_cuda_cast<dst_t>(d.x * (q[l] >>  4) + d.y);
136    }
137}
138
139//================================== k-quants
140
141template<typename dst_t>
142static __global__ void dequantize_block_q2_K(const void * __restrict__ vx, dst_t * __restrict__ yy) {
143    const int64_t i = blockIdx.x;
144
145    dequantize_q2_K(vx, i, yy + i*QK_K, threadIdx.x);
146}
147
148template<typename dst_t>
149static __global__ void dequantize_block_q3_K(const void * __restrict__ vx, dst_t * __restrict__ yy) {
150    const int64_t i = blockIdx.x;
151
152    dequantize_q3_K(vx, i, yy + i*QK_K, threadIdx.x);
153}
154
155template<typename dst_t>
156static __global__ void dequantize_block_q4_K(const void * __restrict__ vx, dst_t * __restrict__ yy) {
157    const int64_t i = blockIdx.x;
158
159    dequantize_q4_K(vx, i, yy + i*QK_K, threadIdx.x);
160}
161
162template<typename dst_t>
163static __global__ void dequantize_block_q5_K(const void * __restrict__ vx, dst_t * __restrict__ yy) {
164    const int64_t i = blockIdx.x;
165
166    dequantize_q5_K(vx, i, yy + i*QK_K, threadIdx.x);
167}
168
169template<typename dst_t>
170static __global__ void dequantize_block_q6_K(const void * __restrict__ vx, dst_t * __restrict__ yy) {
171    const int64_t i = blockIdx.x;
172
173    dequantize_q6_K(vx, i, yy + i*QK_K, threadIdx.x);
174}
175
176template<typename dst_t>
177static __global__ void dequantize_block_iq2_xxs(const void * __restrict__ vx, dst_t * __restrict__ yy) {
178    const int64_t i = blockIdx.x;
179
180    dequantize_iq2_xxs(vx, i, yy + i*QK_K, threadIdx.x);
181}
182
183template<typename dst_t>
184static __global__ void dequantize_block_iq2_xs(const void * __restrict__ vx, dst_t * __restrict__ yy) {
185    const int64_t i = blockIdx.x;
186
187    dequantize_iq2_xs(vx, i, yy + i*QK_K, threadIdx.x);
188}
189
190template<typename dst_t>
191static __global__ void dequantize_block_iq2_s(const void * __restrict__ vx, dst_t * __restrict__ yy) {
192    const int64_t i = blockIdx.x;
193
194    dequantize_iq2_s(vx, i, yy + i*QK_K, threadIdx.x);
195}
196
197template<typename dst_t>
198static __global__ void dequantize_block_iq3_xxs(const void * __restrict__ vx, dst_t * __restrict__ yy) {
199    const int64_t i = blockIdx.x;
200
201    dequantize_iq3_xxs(vx, i, yy + i*QK_K, threadIdx.x);
202}
203
204template<typename dst_t>
205static __global__ void dequantize_block_iq3_s(const void * __restrict__ vx, dst_t * __restrict__ yy) {
206    const int64_t i = blockIdx.x;
207
208    dequantize_iq3_s(vx, i, yy + i*QK_K, threadIdx.x);
209}
210
211template<typename dst_t>
212static __global__ void dequantize_block_iq1_s(const void * __restrict__ vx, dst_t * __restrict__ yy) {
213    const int64_t i = blockIdx.x;
214
215    dequantize_iq1_s(vx, i, yy + i*QK_K, threadIdx.x);
216}
217
218template<typename dst_t>
219static __global__ void dequantize_block_iq1_m(const void * __restrict__ vx, dst_t * __restrict__ yy) {
220    const int64_t i = blockIdx.x;
221
222    dequantize_iq1_m(vx, i, yy + i*QK_K, threadIdx.x);
223}
224
225template<typename dst_t>
226static __global__ void dequantize_block_iq4_nl(const void * __restrict__ vx, dst_t * __restrict__ yy) {
227    const int64_t i = blockIdx.x;
228
229    dequantize_iq4_nl(vx, i, yy + i*QK_K, threadIdx.x);
230}
231
232template<typename dst_t>
233static __global__ void dequantize_block_iq4_xs(const void * __restrict__ vx, dst_t * __restrict__ yy) {
234    const int64_t i = blockIdx.x;
235
236    dequantize_iq4_xs(vx, i, yy + i*QK_K, threadIdx.x);
237}
238
239template<typename dst_t>
240static __global__ void dequantize_block_mxfp4(const void * __restrict__ vx, dst_t * __restrict__ yy) {
241    const int64_t i = blockIdx.x;
242
243    dequantize_mxfp4(vx, i, yy + i*QK_K, threadIdx.x);
244}
245
246template <int qk, int qr, dequantize_kernel_t dequantize_kernel, typename dst_t>
247static void dequantize_block_cuda(const void * vx, dst_t * y,
248        const int64_t ne00, const int64_t ne01, const int64_t ne02, const int64_t ne03,
249        const int64_t s01, const int64_t s02, const int64_t s03, cudaStream_t stream) {
250    const int64_t ne0203 = ne02*ne03;
251    const uint3 ne02_fdv = init_fastdiv_values(ne02);
252    const dim3 num_blocks((ne00 + 2*CUDA_DEQUANTIZE_BLOCK_SIZE - 1) / (2*CUDA_DEQUANTIZE_BLOCK_SIZE), (int)std::min(ne01, (int64_t)65535), (int)std::min(ne0203, (int64_t)65535));
253    dequantize_block<qk, qr, dequantize_kernel><<<num_blocks, CUDA_DEQUANTIZE_BLOCK_SIZE, 0, stream>>>
254        (vx, y, ne00, ne01, ne0203, ne02_fdv, s01, s02, s03);
255}
256
257template <int qk, int qr, dequantize_kernel_t dequantize_kernel, typename dst_t>
258static void dequantize_block_cont_cuda(const void * __restrict__ vx, dst_t * __restrict__ y, const int64_t k, cudaStream_t stream) {
259    dequantize_block_cuda<qk, qr, dequantize_kernel, dst_t>(vx, y, k, 1, 1, 1, k/qk, k/qk, k/qk, stream);
260}
261
262static void dequantize_block_q8_0_f16_cuda(const void * __restrict__ vx, half * __restrict__ y, const int64_t k, cudaStream_t stream) {
263    const int num_blocks = (k + CUDA_Q8_0_NE_ALIGN - 1) / CUDA_Q8_0_NE_ALIGN;
264    if (k % CUDA_Q8_0_NE_ALIGN == 0) {
265        const bool need_check = false;
266        dequantize_block_q8_0_f16<need_check><<<num_blocks, WARP_SIZE, 0, stream>>>(vx, y, k);
267    } else {
268        const bool need_check = true;
269        dequantize_block_q8_0_f16<need_check><<<num_blocks, WARP_SIZE, 0, stream>>>(vx, y, k);
270    }
271}
272
273template<typename dst_t>
274static void dequantize_row_q2_K_cuda(const void * vx, dst_t * y, const int64_t k, cudaStream_t stream) {
275    const int nb = k / QK_K;
276    dequantize_block_q2_K<<<nb, 64, 0, stream>>>(vx, y);
277}
278
279template<typename dst_t>
280static void dequantize_row_q3_K_cuda(const void * vx, dst_t * y, const int64_t k, cudaStream_t stream) {
281    const int nb = k / QK_K;
282    dequantize_block_q3_K<<<nb, 64, 0, stream>>>(vx, y);
283}
284
285template<typename dst_t>
286static void dequantize_row_q4_0_cuda(const void * vx, dst_t * y, const int64_t k, cudaStream_t stream) {
287    const int nb32 = k / 32;
288    const int nb = (k + 255) / 256;
289    dequantize_block_q4_0<<<nb, 32, 0, stream>>>(vx, y, nb32);
290}
291
292template<typename dst_t>
293static void dequantize_row_q4_1_cuda(const void * vx, dst_t * y, const int64_t k, cudaStream_t stream) {
294    const int nb32 = k / 32;
295    const int nb = (k + 255) / 256;
296    dequantize_block_q4_1<<<nb, 32, 0, stream>>>(vx, y, nb32);
297}
298
299template<typename dst_t>
300static void dequantize_row_q4_K_cuda(const void * vx, dst_t * y, const int64_t k, cudaStream_t stream) {
301    const int nb = k / QK_K;
302    dequantize_block_q4_K<<<nb, 32, 0, stream>>>(vx, y);
303}
304
305template<typename dst_t>
306static void dequantize_row_q5_K_cuda(const void * vx, dst_t * y, const int64_t k, cudaStream_t stream) {
307    const int nb = k / QK_K;
308    dequantize_block_q5_K<<<nb, 64, 0, stream>>>(vx, y);
309}
310
311template<typename dst_t>
312static void dequantize_row_q6_K_cuda(const void * vx, dst_t * y, const int64_t k, cudaStream_t stream) {
313    const int nb = k / QK_K;
314    dequantize_block_q6_K<<<nb, 64, 0, stream>>>(vx, y);
315}
316
317template<typename dst_t>
318static void dequantize_row_iq2_xxs_cuda(const void * vx, dst_t * y, const int64_t k, cudaStream_t stream) {
319    const int nb = k / QK_K;
320    dequantize_block_iq2_xxs<<<nb, 32, 0, stream>>>(vx, y);
321}
322
323template<typename dst_t>
324static void dequantize_row_iq2_xs_cuda(const void * vx, dst_t * y, const int64_t k, cudaStream_t stream) {
325    const int nb = k / QK_K;
326    dequantize_block_iq2_xs<<<nb, 32, 0, stream>>>(vx, y);
327}
328
329template<typename dst_t>
330static void dequantize_row_iq2_s_cuda(const void * vx, dst_t * y, const int64_t k, cudaStream_t stream) {
331    const int nb = k / QK_K;
332    dequantize_block_iq2_s<<<nb, 32, 0, stream>>>(vx, y);
333}
334
335template<typename dst_t>
336static void dequantize_row_iq3_xxs_cuda(const void * vx, dst_t * y, const int64_t k, cudaStream_t stream) {
337    const int nb = k / QK_K;
338    dequantize_block_iq3_xxs<<<nb, 32, 0, stream>>>(vx, y);
339}
340
341template<typename dst_t>
342static void dequantize_row_iq3_s_cuda(const void * vx, dst_t * y, const int64_t k, cudaStream_t stream) {
343    const int nb = k / QK_K;
344    dequantize_block_iq3_s<<<nb, 32, 0, stream>>>(vx, y);
345}
346
347template<typename dst_t>
348static void dequantize_row_iq1_s_cuda(const void * vx, dst_t * y, const int64_t k, cudaStream_t stream) {
349    const int nb = k / QK_K;
350    dequantize_block_iq1_s<<<nb, 32, 0, stream>>>(vx, y);
351}
352
353template<typename dst_t>
354static void dequantize_row_iq4_nl_cuda(const void * vx, dst_t * y, const int64_t k, cudaStream_t stream) {
355    const int nb = (k + QK_K - 1) / QK_K;
356    dequantize_block_iq4_nl<<<nb, 32, 0, stream>>>(vx, y);
357}
358
359template<typename dst_t>
360static void dequantize_row_iq1_m_cuda(const void * vx, dst_t * y, const int64_t k, cudaStream_t stream) {
361    const int nb = k / QK_K;
362    dequantize_block_iq1_m<<<nb, 32, 0, stream>>>(vx, y);
363}
364
365template<typename dst_t>
366static void dequantize_row_iq4_xs_cuda(const void * vx, dst_t * y, const int64_t k, cudaStream_t stream) {
367    const int nb = (k + QK_K - 1) / QK_K;
368    dequantize_block_iq4_xs<<<nb, 32, 0, stream>>>(vx, y);
369}
370
371template<typename dst_t>
372static void dequantize_row_mxfp4_cuda(const void * vx, dst_t * y, const int64_t k, cudaStream_t stream) {
373    const int nb = (k + QK_K - 1) / QK_K;
374    dequantize_block_mxfp4<<<nb, 32, 0, stream>>>(vx, y);
375}
376
377template <typename dst_t>
378static __global__ void dequantize_block_nvfp4(
379        const void * __restrict__ vx,
380        dst_t * __restrict__ yy,
381        const int64_t ne) {
382    const int64_t i = blockIdx.x;
383    const int     tid = threadIdx.x;
384
385    const int64_t base = i * QK_NVFP4;
386    if (base >= ne) {
387        return;
388    }
389
390    const block_nvfp4 * x = (const block_nvfp4 *) vx;
391    const block_nvfp4 & xb = x[i];
392
393    const int sub = tid / (QK_NVFP4_SUB / 2);
394    const int j = tid % (QK_NVFP4_SUB / 2);
395
396    const float d = ggml_cuda_ue4m3_to_fp32(xb.d[sub]);
397    const uint8_t q = xb.qs[sub * (QK_NVFP4_SUB / 2) + j];
398
399    const int64_t y0 = base + sub * QK_NVFP4_SUB + j;
400    const int64_t y1 = y0 + QK_NVFP4_SUB / 2;
401
402    yy[y0] = ggml_cuda_cast<dst_t>(d * kvalues_mxfp4[q & 0x0F]);
403    yy[y1] = ggml_cuda_cast<dst_t>(d * kvalues_mxfp4[q >> 4]);
404}
405
406template <typename dst_t>
407static void dequantize_row_nvfp4_cuda(
408        const void * vx,
409        dst_t * y,
410        const int64_t k,
411        cudaStream_t stream) {
412    GGML_ASSERT(k % QK_NVFP4 == 0);
413    const int nb = k / QK_NVFP4;
414    dequantize_block_nvfp4<<<nb, 32, 0, stream>>>(vx, y, k);
415}
416template <typename src_t, typename dst_t>
417static __global__ void convert_unary(
418        const void * __restrict__ vx, dst_t * __restrict__ y, const int64_t ne00, const int64_t ne01,
419        const int64_t ne0203, const uint3 ne02,
420        const int64_t s01, const int64_t s02, const int64_t s03) {
421    const int64_t i00 = (int64_t)blockDim.x*blockIdx.x + threadIdx.x;
422
423    if (i00 >= ne00) {
424        return;
425    }
426
427    const src_t * x = (const src_t *) vx;
428
429    for (int64_t i01 = blockIdx.y; i01 < ne01; i01 += gridDim.y) {
430        for (int64_t i0203 = blockIdx.z; i0203 < ne0203; i0203 += gridDim.z) {
431            const uint2 dm = fast_div_modulo((uint32_t)i0203, ne02);
432            const int64_t i02 = dm.y;
433            const int64_t i03 = dm.x;
434
435            const int64_t ix = i03*s03 + i02*s02 + i01*s01 + i00;
436            const int64_t iy = (i0203*ne01 + i01)*ne00 + i00;
437            y[iy] = ggml_cuda_cast<dst_t>(x[ix]);
438        }
439    }
440}
441
442template <typename src_t, typename dst_t>
443static void convert_unary_cuda(const void * vx, dst_t * y,
444        const int64_t ne00, const int64_t ne01, const int64_t ne02, const int64_t ne03,
445        const int64_t s01, const int64_t s02, const int64_t s03, cudaStream_t stream) {
446    const int64_t ne0203 = ne02*ne03;
447    const uint3 ne02_fdv = init_fastdiv_values(ne02);
448    const dim3 num_blocks((ne00 + CUDA_DEQUANTIZE_BLOCK_SIZE - 1) / CUDA_DEQUANTIZE_BLOCK_SIZE, (int)std::min(ne01, (int64_t)65535), (int)std::min(ne0203, (int64_t)65535));
449    convert_unary<src_t><<<num_blocks, CUDA_DEQUANTIZE_BLOCK_SIZE, 0, stream>>>
450        (vx, y, ne00, ne01, ne0203, ne02_fdv, s01, s02, s03);
451}
452
453template <typename src_t, typename dst_t>
454static void convert_unary_cont_cuda(const void * vx, dst_t * y, const int64_t k, cudaStream_t stream) {
455    convert_unary_cuda<src_t>(vx, y, k, 1, 1, 1, k, k, k, stream);
456}
457
458to_bf16_cuda_t ggml_get_to_bf16_cuda(ggml_type type) {
459    switch (type) {
460        case GGML_TYPE_Q1_0:
461            return dequantize_block_cont_cuda<QK1_0, QR1_0, dequantize_q1_0>;
462        case GGML_TYPE_Q2_0:
463            return dequantize_block_cont_cuda<QK2_0, QR2_0, dequantize_q2_0>;
464        case GGML_TYPE_Q4_0:
465            return dequantize_row_q4_0_cuda;
466        case GGML_TYPE_Q4_1:
467            return dequantize_row_q4_1_cuda;
468        case GGML_TYPE_Q5_0:
469            return dequantize_block_cont_cuda<QK5_0, QR5_0, dequantize_q5_0>;
470        case GGML_TYPE_Q5_1:
471            return dequantize_block_cont_cuda<QK5_1, QR5_1, dequantize_q5_1>;
472        case GGML_TYPE_Q8_0:
473            return dequantize_block_cont_cuda<QK8_0, QR8_0, dequantize_q8_0>;
474        case GGML_TYPE_Q2_K:
475            return dequantize_row_q2_K_cuda;
476        case GGML_TYPE_Q3_K:
477            return dequantize_row_q3_K_cuda;
478        case GGML_TYPE_Q4_K:
479            return dequantize_row_q4_K_cuda;
480        case GGML_TYPE_Q5_K:
481            return dequantize_row_q5_K_cuda;
482        case GGML_TYPE_Q6_K:
483            return dequantize_row_q6_K_cuda;
484        case GGML_TYPE_IQ2_XXS:
485            return dequantize_row_iq2_xxs_cuda;
486        case GGML_TYPE_IQ2_XS:
487            return dequantize_row_iq2_xs_cuda;
488        case GGML_TYPE_IQ2_S:
489            return dequantize_row_iq2_s_cuda;
490        case GGML_TYPE_IQ3_XXS:
491            return dequantize_row_iq3_xxs_cuda;
492        case GGML_TYPE_IQ1_S:
493            return dequantize_row_iq1_s_cuda;
494        case GGML_TYPE_IQ1_M:
495            return dequantize_row_iq1_m_cuda;
496        case GGML_TYPE_IQ4_NL:
497            return dequantize_row_iq4_nl_cuda;
498        case GGML_TYPE_IQ4_XS:
499            return dequantize_row_iq4_xs_cuda;
500        case GGML_TYPE_IQ3_S:
501            return dequantize_row_iq3_s_cuda;
502        case GGML_TYPE_MXFP4:
503            return dequantize_row_mxfp4_cuda;
504        case GGML_TYPE_NVFP4:
505            return dequantize_row_nvfp4_cuda;
506        case GGML_TYPE_F32:
507            return convert_unary_cont_cuda<float>;
508        case GGML_TYPE_F16:
509            return convert_unary_cont_cuda<half>;
510        default:
511            return nullptr;
512    }
513}
514
515to_fp16_cuda_t ggml_get_to_fp16_cuda(ggml_type type) {
516    switch (type) {
517        case GGML_TYPE_Q1_0:
518            return dequantize_block_cont_cuda<QK1_0, QR1_0, dequantize_q1_0>;
519        case GGML_TYPE_Q2_0:
520            return dequantize_block_cont_cuda<QK2_0, QR2_0, dequantize_q2_0>;
521        case GGML_TYPE_Q4_0:
522            return dequantize_row_q4_0_cuda;
523        case GGML_TYPE_Q4_1:
524            return dequantize_row_q4_1_cuda;
525        case GGML_TYPE_Q5_0:
526            return dequantize_block_cont_cuda<QK5_0, QR5_0, dequantize_q5_0>;
527        case GGML_TYPE_Q5_1:
528            return dequantize_block_cont_cuda<QK5_1, QR5_1, dequantize_q5_1>;
529        case GGML_TYPE_Q8_0:
530            if (fp16_available(ggml_cuda_info().devices[ggml_cuda_get_device()].cc)) {
531                return dequantize_block_q8_0_f16_cuda;
532            }
533            return dequantize_block_cont_cuda<QK8_0, QR8_0, dequantize_q8_0>;
534        case GGML_TYPE_Q2_K:
535            return dequantize_row_q2_K_cuda;
536        case GGML_TYPE_Q3_K:
537            return dequantize_row_q3_K_cuda;
538        case GGML_TYPE_Q4_K:
539            return dequantize_row_q4_K_cuda;
540        case GGML_TYPE_Q5_K:
541            return dequantize_row_q5_K_cuda;
542        case GGML_TYPE_Q6_K:
543            return dequantize_row_q6_K_cuda;
544        case GGML_TYPE_IQ2_XXS:
545            return dequantize_row_iq2_xxs_cuda;
546        case GGML_TYPE_IQ2_XS:
547            return dequantize_row_iq2_xs_cuda;
548        case GGML_TYPE_IQ2_S:
549            return dequantize_row_iq2_s_cuda;
550        case GGML_TYPE_IQ3_XXS:
551            return dequantize_row_iq3_xxs_cuda;
552        case GGML_TYPE_IQ1_S:
553            return dequantize_row_iq1_s_cuda;
554        case GGML_TYPE_IQ1_M:
555            return dequantize_row_iq1_m_cuda;
556        case GGML_TYPE_IQ4_NL:
557            return dequantize_row_iq4_nl_cuda;
558        case GGML_TYPE_IQ4_XS:
559            return dequantize_row_iq4_xs_cuda;
560        case GGML_TYPE_IQ3_S:
561            return dequantize_row_iq3_s_cuda;
562        case GGML_TYPE_MXFP4:
563            return dequantize_row_mxfp4_cuda;
564        case GGML_TYPE_NVFP4:
565            return dequantize_row_nvfp4_cuda;
566        case GGML_TYPE_F32:
567            return convert_unary_cont_cuda<float>;
568        case GGML_TYPE_BF16:
569            return convert_unary_cont_cuda<nv_bfloat16>;
570        default:
571            return nullptr;
572    }
573}
574
575to_fp32_cuda_t ggml_get_to_fp32_cuda(ggml_type type) {
576    switch (type) {
577        case GGML_TYPE_Q1_0:
578            return dequantize_block_cont_cuda<QK1_0, QR1_0, dequantize_q1_0>;
579        case GGML_TYPE_Q2_0:
580            return dequantize_block_cont_cuda<QK2_0, QR2_0, dequantize_q2_0>;
581        case GGML_TYPE_Q4_0:
582            return dequantize_row_q4_0_cuda;
583        case GGML_TYPE_Q4_1:
584            return dequantize_row_q4_1_cuda;
585        case GGML_TYPE_Q5_0:
586            return dequantize_block_cont_cuda<QK5_0, QR5_0, dequantize_q5_0>;
587        case GGML_TYPE_Q5_1:
588            return dequantize_block_cont_cuda<QK5_1, QR5_1, dequantize_q5_1>;
589        case GGML_TYPE_Q8_0:
590            return dequantize_block_cont_cuda<QK8_0, QR8_0, dequantize_q8_0>;
591        case GGML_TYPE_Q2_K:
592            return dequantize_row_q2_K_cuda;
593        case GGML_TYPE_Q3_K:
594            return dequantize_row_q3_K_cuda;
595        case GGML_TYPE_Q4_K:
596            return dequantize_row_q4_K_cuda;
597        case GGML_TYPE_Q5_K:
598            return dequantize_row_q5_K_cuda;
599        case GGML_TYPE_Q6_K:
600            return dequantize_row_q6_K_cuda;
601        case GGML_TYPE_IQ2_XXS:
602            return dequantize_row_iq2_xxs_cuda;
603        case GGML_TYPE_IQ2_XS:
604            return dequantize_row_iq2_xs_cuda;
605        case GGML_TYPE_IQ2_S:
606            return dequantize_row_iq2_s_cuda;
607        case GGML_TYPE_IQ3_XXS:
608            return dequantize_row_iq3_xxs_cuda;
609        case GGML_TYPE_IQ1_S:
610            return dequantize_row_iq1_s_cuda;
611        case GGML_TYPE_IQ1_M:
612            return dequantize_row_iq1_m_cuda;
613        case GGML_TYPE_IQ4_NL:
614            return dequantize_row_iq4_nl_cuda;
615        case GGML_TYPE_IQ4_XS:
616            return dequantize_row_iq4_xs_cuda;
617        case GGML_TYPE_IQ3_S:
618            return dequantize_row_iq3_s_cuda;
619        case GGML_TYPE_MXFP4:
620            return dequantize_row_mxfp4_cuda;
621        case GGML_TYPE_NVFP4:
622            return dequantize_row_nvfp4_cuda;
623        case GGML_TYPE_F16:
624            return convert_unary_cont_cuda<half>;
625        case GGML_TYPE_BF16:
626            return convert_unary_cont_cuda<nv_bfloat16>;
627        default:
628            return nullptr;
629    }
630}
631
632to_fp16_nc_cuda_t ggml_get_to_fp16_nc_cuda(ggml_type type) {
633    switch (type) {
634        case GGML_TYPE_F32:
635            return convert_unary_cuda<float>;
636        case GGML_TYPE_Q1_0:
637            return dequantize_block_cuda<QK1_0, QR1_0, dequantize_q1_0>;
638        case GGML_TYPE_Q2_0:
639            return dequantize_block_cuda<QK2_0, QR2_0, dequantize_q2_0>;
640        case GGML_TYPE_Q4_0:
641            return dequantize_block_cuda<QK4_0, QR4_0, dequantize_q4_0>;
642        case GGML_TYPE_Q4_1:
643            return dequantize_block_cuda<QK4_1, QR4_1, dequantize_q4_1>;
644        case GGML_TYPE_Q5_0:
645            return dequantize_block_cuda<QK5_0, QR5_0, dequantize_q5_0>;
646        case GGML_TYPE_Q5_1:
647            return dequantize_block_cuda<QK5_1, QR5_1, dequantize_q5_1>;
648        case GGML_TYPE_Q8_0:
649            return dequantize_block_cuda<QK8_0, QR8_0, dequantize_q8_0>;
650        case GGML_TYPE_BF16:
651            return convert_unary_cuda<nv_bfloat16>;
652        default:
653            return nullptr;
654    }
655}
656
657to_bf16_nc_cuda_t ggml_get_to_bf16_nc_cuda(ggml_type type) {
658    switch (type) {
659        case GGML_TYPE_F32:
660            return convert_unary_cuda<float, nv_bfloat16>;
661        case GGML_TYPE_Q1_0:
662            return dequantize_block_cuda<QK1_0, QR1_0, dequantize_q1_0>;
663        case GGML_TYPE_Q2_0:
664            return dequantize_block_cuda<QK2_0, QR2_0, dequantize_q2_0>;
665        case GGML_TYPE_Q4_0:
666            return dequantize_block_cuda<QK4_0, QR4_0, dequantize_q4_0>;
667        case GGML_TYPE_Q4_1:
668            return dequantize_block_cuda<QK4_1, QR4_1, dequantize_q4_1>;
669        case GGML_TYPE_Q5_0:
670            return dequantize_block_cuda<QK5_0, QR5_0, dequantize_q5_0>;
671        case GGML_TYPE_Q5_1:
672            return dequantize_block_cuda<QK5_1, QR5_1, dequantize_q5_1>;
673        case GGML_TYPE_Q8_0:
674            return dequantize_block_cuda<QK8_0, QR8_0, dequantize_q8_0>;
675        case GGML_TYPE_F16:
676            return convert_unary_cuda<half, nv_bfloat16>;
677        default:
678            return nullptr;
679    }
680}
681
682to_fp32_nc_cuda_t ggml_get_to_fp32_nc_cuda(ggml_type type) {
683    switch (type) {
684        case GGML_TYPE_F16:
685            return convert_unary_cuda<half, float>;
686        case GGML_TYPE_Q1_0:
687            return dequantize_block_cuda<QK1_0, QR1_0, dequantize_q1_0>;
688        case GGML_TYPE_Q2_0:
689            return dequantize_block_cuda<QK2_0, QR2_0, dequantize_q2_0>;
690        case GGML_TYPE_Q4_0:
691            return dequantize_block_cuda<QK4_0, QR4_0, dequantize_q4_0>;
692        case GGML_TYPE_Q4_1:
693            return dequantize_block_cuda<QK4_1, QR4_1, dequantize_q4_1>;
694        case GGML_TYPE_Q5_0:
695            return dequantize_block_cuda<QK5_0, QR5_0, dequantize_q5_0>;
696        case GGML_TYPE_Q5_1:
697            return dequantize_block_cuda<QK5_1, QR5_1, dequantize_q5_1>;
698        case GGML_TYPE_Q8_0:
699            return dequantize_block_cuda<QK8_0, QR8_0, dequantize_q8_0>;
700        case GGML_TYPE_BF16:
701            return convert_unary_cuda<nv_bfloat16, float>;
702        default:
703            return nullptr;
704    }
705}