char-fork/llama.cpp

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

lhezopencl: improve get_rows, cpy, concat and q6_k flat gemv (#24160)308f61c31

char/aliceai-t5-moe
brancheschar/aliceai-t5-moeall refs →
5.4 KiB193 linesraw
1#pragma OPENCL EXTENSION cl_khr_fp16 : enable
2
3typedef char int8_t;
4typedef uchar uint8_t;
5typedef short int16_t;
6typedef ushort uint16_t;
7typedef int int32_t;
8typedef uint uint32_t;
9
10#define QK4_0                   32
11
12//------------------------------------------------------------------------------
13// block_q4_0
14//------------------------------------------------------------------------------
15struct block_q4_0
16{
17    half d;
18    uint8_t qs[QK4_0 / 2];
19};
20
21
22//------------------------------------------------------------------------------
23// dequantize_q4_0_f32, dequantize_q4_0_f16
24//------------------------------------------------------------------------------
25void dequantize_q4_0_f32(global struct block_q4_0 * xb, short il, float16 * reg) {
26    global ushort * qs = ((global ushort *)xb + 1);
27    float d1 = il ? (xb->d / 16.h) : xb->d;
28    float d2 = d1 / 256.f;
29    float md = -8.h * xb->d;
30    ushort mask0 = il ? 0x00F0 : 0x000F;
31    ushort mask1 = mask0 << 8;
32
33    reg->s0 = d1 * (qs[0] & mask0) + md;
34    reg->s1 = d2 * (qs[0] & mask1) + md;
35
36    reg->s2 = d1 * (qs[1] & mask0) + md;
37    reg->s3 = d2 * (qs[1] & mask1) + md;
38
39    reg->s4 = d1 * (qs[2] & mask0) + md;
40    reg->s5 = d2 * (qs[2] & mask1) + md;
41
42    reg->s6 = d1 * (qs[3] & mask0) + md;
43    reg->s7 = d2 * (qs[3] & mask1) + md;
44
45    reg->s8 = d1 * (qs[4] & mask0) + md;
46    reg->s9 = d2 * (qs[4] & mask1) + md;
47
48    reg->sa = d1 * (qs[5] & mask0) + md;
49    reg->sb = d2 * (qs[5] & mask1) + md;
50
51    reg->sc = d1 * (qs[6] & mask0) + md;
52    reg->sd = d2 * (qs[6] & mask1) + md;
53
54    reg->se = d1 * (qs[7] & mask0) + md;
55    reg->sf = d2 * (qs[7] & mask1) + md;
56}
57
58
59//------------------------------------------------------------------------------
60// get_rows
61//------------------------------------------------------------------------------
62kernel void kernel_get_rows_f32(
63        global void * src0,
64        ulong offset0,
65        global int * src1,
66        ulong offset1,
67        global float * dst,
68        ulong offsetd,
69        int ne00,
70        ulong nb01,
71        ulong nb02,
72        ulong nb03,
73        int ne10,
74        ulong nb10,
75        ulong nb11,
76        ulong nb12,
77        ulong nb1,
78        ulong nb2,
79        ulong nb3
80) {
81    src0 = (global void*)((global char*)src0 + offset0);
82    src1 = (global int*)((global char*)src1 + offset1);
83    dst = (global float*)((global char*)dst + offsetd);
84
85    int nchunks = get_num_groups(0) / ne10;
86    int g       = get_group_id(0);
87    int i10     = g / nchunks;
88    int chunk   = g - i10 * nchunks;
89    int i11     = get_group_id(1);
90    int i12     = get_group_id(2);
91
92    int r = ((global int *) ((global char *) src1 + i12*nb12 + i11*nb11 + i10*nb10))[0];
93
94    int i02 = i11;
95    int i03 = i12;
96
97    global float * dst_row = (global float *) ((global char *) dst  + i12*nb3 + i11*nb2 + i10*nb1);
98    global float * src_row = (global float *) ((global char *) src0 + r*nb01 + i02*nb02 + i03*nb03);
99
100    int span  = (ne00 + nchunks - 1) / nchunks;
101    int start = chunk * span;
102    int end   = min(start + span, ne00);
103
104    for (int ind = start + get_local_id(0); ind < end; ind += get_local_size(0)) {
105        dst_row[ind] = src_row[ind];
106    }
107}
108
109kernel void kernel_get_rows_f16(
110        global void * src0,
111        ulong offset0,
112        global int * src1,
113        ulong offset1,
114        global float * dst,
115        ulong offsetd,
116        int ne00,
117        ulong nb01,
118        ulong nb02,
119        ulong nb03,
120        int ne10,
121        ulong nb10,
122        ulong nb11,
123        ulong nb12,
124        ulong nb1,
125        ulong nb2,
126        ulong nb3
127) {
128    src0 = (global void*)((global char*)src0 + offset0);
129    src1 = (global int*)((global char*)src1 + offset1);
130    dst = (global float*)((global char*)dst + offsetd);
131
132    int i10 = get_group_id(0);
133    int i11 = get_group_id(1);
134    int i12 = get_group_id(2);
135
136    int r = ((global int32_t *) ((global char *) src1 + i12*nb12 + i11*nb11 + i10*nb10))[0];
137
138    int i02 = i11;
139    int i03 = i12;
140
141    for (int ind = get_local_id(0); ind < ne00; ind += get_local_size(0)) {
142        if (ind >= ne00) {
143            return;
144        }
145        ((global float *) ((global char *) dst + i12*nb3 + i11*nb2 + i10*nb1))[ind] =
146            ((global half *) ((global char *) src0 + r*nb01 + i02*nb02 + i03*nb03))[ind];
147    }
148}
149
150kernel void kernel_get_rows_q4_0(
151        global void * src0,
152        ulong offset0,
153        global int * src1,
154        ulong offset1,
155        global float * dst,
156        ulong offsetd,
157        int ne00,
158        ulong nb01,
159        ulong nb02,
160        ulong nb03,
161        int ne10,
162        ulong nb10,
163        ulong nb11,
164        ulong nb12,
165        ulong nb1,
166        ulong nb2,
167        ulong nb3
168) {
169    src0 = (global void*)((global char*)src0 + offset0);
170    src1 = (global int*)((global char*)src1 + offset1);
171    dst = (global float*)((global char*)dst + offsetd);
172
173    const int NL = 2;
174
175    int i10 = get_group_id(0);
176    int i11 = get_group_id(1);
177    int i12 = get_group_id(2);
178
179    int r = ((global int32_t *) ((global char *) src1 + i12*nb12 + i11*nb11 + i10*nb10))[0];
180
181    int i02 = i11;
182    int i03 = i12;
183
184    for (int ind = get_local_id(0); ind < ne00/16; ind += get_local_size(0)) {
185        float16 temp;
186        if (ind >= ne00) {
187            return;
188        }
189        dequantize_q4_0_f32(
190            ((global struct block_q4_0 *) ((global char *) src0 + r*nb01 + i02*nb02 + i03*nb03)) + ind/NL, ind%NL, &temp);
191        *(((global float16 *) ((global char *) dst + i12*nb3 + i11*nb2 + i10*nb1)) + ind) = temp;
192    }
193}