char-fork/llama.cpp

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

fairydreamingcuda : relax tensor contiguity requirements for quantized concat (#25678)a3e5b96ac

char/aliceai-t5-moe
brancheschar/aliceai-t5-moeall refs →
9.5 KiB242 linesraw
1#include "concat.cuh"
2
3#include <stdint.h>
4
5// contiguous kernels
6template <typename T, int dim>
7static __global__ void __launch_bounds__(CUDA_CONCAT_BLOCK_SIZE) concat_cont(const T * x,
8                                                                             const T * y,
9                                                                             T *       dst,
10                                                                             int64_t   ne00,
11                                                                             int64_t   ne01,
12                                                                             int64_t   ne02,
13                                                                             int64_t   ne0,
14                                                                             int64_t   ne1,
15                                                                             int64_t   ne2) {
16    static_assert(dim >= 0 && dim <= 2, "dim must be in [0, 2]");
17
18    const int64_t n = ne0 * ne1 * ne2;
19
20    ggml_cuda_pdl_sync();
21    for (int64_t i = (int64_t) blockIdx.x * blockDim.x + threadIdx.x; i < n; i += (int64_t) blockDim.x * gridDim.x) {
22        if constexpr (dim == 0) {
23            const int64_t row = i / ne0;
24            const int64_t i0  = i - row * ne0;
25
26            if (i0 < ne00) {
27                dst[i] = x[row * ne00 + i0];
28            } else {
29                dst[i] = y[row * (ne0 - ne00) + (i0 - ne00)];
30            }
31        } else if constexpr (dim == 1) {
32            const int64_t dst_plane  = ne0 * ne1;
33            const int64_t src0_plane = ne0 * ne01;
34            const int64_t src1_plane = dst_plane - src0_plane;
35            const int64_t i2         = i / dst_plane;
36            const int64_t i01        = i - i2 * dst_plane;
37
38            if (i01 < src0_plane) {
39                dst[i] = x[i2 * src0_plane + i01];
40            } else {
41                dst[i] = y[i2 * src1_plane + (i01 - src0_plane)];
42            }
43        } else {
44            const int64_t src0_size = ne0 * ne1 * ne02;
45
46            if (i < src0_size) {
47                dst[i] = x[i];
48            } else {
49                dst[i] = y[i - src0_size];
50            }
51        }
52    }
53}
54
55template <typename T>
56static void concat_cont_cuda(const T * x,
57                             const T * y,
58                             T *       dst,
59                             int64_t   ne00,
60                             int64_t   ne01,
61                             int64_t   ne02,
62                             int64_t   ne0,
63                             int64_t   ne1,
64                             int64_t   ne2,
65                             int       dim,
66                             cudaStream_t stream) {
67    const int64_t n          = ne0 * ne1 * ne2;
68    const int     num_blocks = (n + CUDA_CONCAT_BLOCK_SIZE - 1) / CUDA_CONCAT_BLOCK_SIZE;
69
70    if (dim == 0) {
71        const ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params(num_blocks, CUDA_CONCAT_BLOCK_SIZE, 0, stream);
72        ggml_cuda_kernel_launch(concat_cont<T, 0>, launch_params, x, y, dst, ne00, ne01, ne02, ne0, ne1, ne2);
73        return;
74    }
75    if (dim == 1) {
76        concat_cont<T, 1><<<num_blocks, CUDA_CONCAT_BLOCK_SIZE, 0, stream>>>(x, y, dst, ne00, ne01, ne02, ne0, ne1, ne2);
77        return;
78    }
79    concat_cont<T, 2><<<num_blocks, CUDA_CONCAT_BLOCK_SIZE, 0, stream>>>(x, y, dst, ne00, ne01, ne02, ne0, ne1, ne2);
80}
81
82// non-contiguous kernel (slow)
83template <typename T, int dim>
84static __global__ void __launch_bounds__(CUDA_CONCAT_BLOCK_SIZE)
85    concat_non_cont(
86        const char * src0,
87        const char * src1,
88              char * dst,
89           int64_t   ne00,
90           int64_t   ne01,
91           int64_t   ne02,
92           int64_t   ne03,
93          uint64_t   nb00,
94          uint64_t   nb01,
95          uint64_t   nb02,
96          uint64_t   nb03,
97           int64_t /*ne10*/,
98           int64_t /*ne11*/,
99           int64_t /*ne12*/,
100           int64_t /*ne13*/,
101          uint64_t   nb10,
102          uint64_t   nb11,
103          uint64_t   nb12,
104          uint64_t   nb13,
105           int64_t   ne0,
106           int64_t /*ne1*/,
107           int64_t /*ne2*/,
108           int64_t /*ne3*/,
109          uint64_t   nb0,
110          uint64_t   nb1,
111          uint64_t   nb2,
112          uint64_t   nb3) {
113    static_assert(dim >= 0 && dim <= 3, "dim must be in [0, 3]");
114
115    const int64_t i3 = blockIdx.z;
116    const int64_t i2 = blockIdx.y;
117    const int64_t i1 = blockIdx.x;
118
119    const T * x;
120
121    for (int64_t i0 = threadIdx.x; i0 < ne0; i0 += blockDim.x) {
122        if (i0 < ne00 && i1 < ne01 && i2 < ne02 && i3 < ne03) {
123            x = (const T *)(src0 + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00);
124        } else {
125            if constexpr (dim == 0) {
126                x = (const T *)(src1 + i3*nb13 + i2*nb12 + i1*nb11 + (i0 - ne00)*nb10);
127            } else if constexpr (dim == 1) {
128                x = (const T *)(src1 + i3*nb13 + i2*nb12 + (i1 - ne01)*nb11 + i0*nb10);
129            } else if constexpr (dim == 2) {
130                x = (const T *)(src1 + i3*nb13 + (i2 - ne02)*nb12 + i1*nb11 + i0*nb10);
131            } else if constexpr (dim == 3) {
132                x = (const T *)(src1 + (i3 - ne03)*nb13 + i2*nb12 + i1*nb11 + i0*nb10);
133            }
134        }
135
136        T * y = (T *)(dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0);
137
138        *y = *x;
139    }
140}
141
142template <typename T>
143static void concat_cuda(const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst, int dim, cudaStream_t stream) {
144    if (dim != 3 && ggml_is_contiguous_to_3(src0) && ggml_is_contiguous_to_3(src1)) {
145        const T * src0_d = (const T *) src0->data;
146        const T * src1_d = (const T *) src1->data;
147        T *       dst_d  = (T *) dst->data;
148
149        for (int64_t i3 = 0; i3 < dst->ne[3]; i3++) {
150            concat_cont_cuda(
151                    src0_d + i3*(src0->nb[3] / sizeof(T)),
152                    src1_d + i3*(src1->nb[3] / sizeof(T)),
153                    dst_d  + i3*( dst->nb[3] / sizeof(T)),
154                    ggml_row_size(src0->type, src0->ne[0])/sizeof(T), src0->ne[1], src0->ne[2],
155                    ggml_row_size(dst->type, dst->ne[0])/sizeof(T),  dst->ne[1],  dst->ne[2], dim, stream);
156        }
157    } else if (dim == 3 && ggml_is_contiguous(src0) && ggml_is_contiguous(src1)) {
158        const size_t size0 = ggml_nbytes(src0);
159        const size_t size1 = ggml_nbytes(src1);
160
161        CUDA_CHECK(cudaMemcpyAsync((char *) dst->data,         src0->data, size0, cudaMemcpyDeviceToDevice, stream));
162        CUDA_CHECK(cudaMemcpyAsync((char *) dst->data + size0, src1->data, size1, cudaMemcpyDeviceToDevice, stream));
163    } else {
164        GGML_ASSERT(!ggml_is_quantized(src0->type));
165
166        dim3 grid_dim(dst->ne[1], dst->ne[2], dst->ne[3]);
167        auto launch_kernel = [&](auto dim) {
168            concat_non_cont<T, dim><<<grid_dim, CUDA_CONCAT_BLOCK_SIZE, 0, stream>>>(
169                (const char *) src0->data, (const char *) src1->data, (char *) dst->data,
170                src0->ne[0], src0->ne[1], src0->ne[2], src0->ne[3],
171                src0->nb[0], src0->nb[1], src0->nb[2], src0->nb[3],
172                src1->ne[0], src1->ne[1], src1->ne[2], src1->ne[3],
173                src1->nb[0], src1->nb[1], src1->nb[2], src1->nb[3],
174                dst->ne[0], dst->ne[1], dst->ne[2], dst->ne[3],
175                dst->nb[0], dst->nb[1], dst->nb[2], dst->nb[3]);
176        };
177        switch (dim) {
178            case 0:
179                launch_kernel(std::integral_constant<int, 0>{});
180                break;
181            case 1:
182                launch_kernel(std::integral_constant<int, 1>{});
183                break;
184            case 2:
185                launch_kernel(std::integral_constant<int, 2>{});
186                break;
187            case 3:
188                launch_kernel(std::integral_constant<int, 3>{});
189                break;
190            default:
191                GGML_ABORT("Invalid dim: %d", dim);
192                break;
193        }
194    }
195}
196
197void ggml_cuda_op_concat(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
198    const ggml_tensor * src0 = dst->src[0];
199    const ggml_tensor * src1 = dst->src[1];
200
201    cudaStream_t stream = ctx.stream();
202
203    const int32_t dim = ((int32_t *) dst->op_params)[0];
204
205    GGML_ASSERT(src0->type == src1->type);
206    GGML_ASSERT(dst->type  == src0->type);
207
208    if (ggml_is_quantized(src0->type)) {
209        if (dim == 3) {
210            GGML_ASSERT(ggml_is_contiguous(src0));
211            GGML_ASSERT(ggml_is_contiguous(src1));
212        } else {
213            GGML_ASSERT(ggml_is_contiguous_to_3(src0));
214            GGML_ASSERT(ggml_is_contiguous_to_3(src1));
215        }
216        GGML_ASSERT(src0->ne[0] % ggml_blck_size(src0->type) == 0);
217        GGML_ASSERT(src1->ne[0] % ggml_blck_size(src1->type) == 0);
218
219        // if first 3 dimensions are contiguous and ne[0] is multiple of the block size we can concat both tensors as byte tensors
220        concat_cuda<uint8_t>(src0, src1, dst, dim, stream);
221    } else {
222        GGML_ASSERT(ggml_blck_size(src0->type) == 1);
223
224        switch (ggml_type_size(src0->type)) {
225            case 1:
226                concat_cuda<uint8_t>(src0, src1, dst, dim, stream);
227                break;
228            case 2:
229                concat_cuda<uint16_t>(src0, src1, dst, dim, stream);
230                break;
231            case 4:
232                concat_cuda<uint32_t>(src0, src1, dst, dim, stream);
233                break;
234            case 8:
235                concat_cuda<uint64_t>(src0, src1, dst, dim, stream);
236                break;
237            default:
238                GGML_ABORT("Unsupported type size: %zu", ggml_type_size(src0->type));
239                break;
240        }
241    }
242}