char-fork/llama.cpp

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

Hongqiang Wangopencl: extend the elementwise and data‐movement op coverage (#27633)1548a240e

char/aliceai-t5-moe
brancheschar/aliceai-t5-moeall refs →
8.6 KiB313 linesraw
1#pragma OPENCL EXTENSION cl_khr_fp16 : enable
2
3//------------------------------------------------------------------------------
4// cpy
5//------------------------------------------------------------------------------
6
7kernel void kernel_cpy_f16_f16(
8        global half * src0,
9        ulong offset0,
10        global half * dst,
11        ulong offsetd,
12        int ne00,
13        int ne01,
14        int ne02,
15        int ne03,
16        ulong nb00,
17        ulong nb01,
18        ulong nb02,
19        ulong nb03,
20        int ne0,
21        int ne1,
22        int ne2,
23        int ne3,
24        ulong nb0,
25        ulong nb1,
26        ulong nb2,
27        ulong nb3
28) {
29    src0 = (global half*)((global char*)src0 + offset0);
30    dst = (global half*)((global char*)dst + offsetd);
31
32    int i03 = get_group_id(2);
33    int i02 = get_group_id(1);
34    int i01 = get_group_id(0);
35
36    int n = i03*ne02*ne01*ne00 + i02*ne01*ne00 + i01*ne00;
37
38    int i3 = n / (ne2*ne1*ne0);
39    int i2 = (n - i3*ne2*ne1*ne0) / (ne1*ne0);
40    int i1 = (n - i3*ne2*ne1*ne0 - i2*ne1*ne0) / ne0;
41    int i0 = (n - i3*ne2*ne1*ne0 - i2*ne1*ne0 - i1*ne0);
42
43    global half * dst_data = (global half *) ((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0);
44
45    for (int i00 = get_local_id(0); i00 < ne00; i00 += get_local_size(0)) {
46        global const half * src = (global half *)((global char *) src0 + i03*nb03 + i02*nb02 + i01*nb01 + i00*nb00);
47        dst_data[i00] = src[0];
48    }
49}
50
51kernel void kernel_cpy_f16_f32(
52        global half * src0,
53        ulong offset0,
54        global float * dst,
55        ulong offsetd,
56        int ne00,
57        int ne01,
58        int ne02,
59        int ne03,
60        ulong nb00,
61        ulong nb01,
62        ulong nb02,
63        ulong nb03,
64        int ne0,
65        int ne1,
66        int ne2,
67        int ne3,
68        ulong nb0,
69        ulong nb1,
70        ulong nb2,
71        ulong nb3
72) {
73
74    src0 = (global half*)((global char*)src0 + offset0);
75    dst = (global float*)((global char*)dst + offsetd);
76
77    int i03 = get_group_id(2);
78    int i02 = get_group_id(1);
79    int i01 = get_group_id(0);
80
81    int n = i03*ne02*ne01*ne00 + i02*ne01*ne00 + i01*ne00;
82
83    int i3 = n / (ne2*ne1*ne0);
84    int i2 = (n - i3*ne2*ne1*ne0) / (ne1*ne0);
85    int i1 = (n - i3*ne2*ne1*ne0 - i2*ne1*ne0) / ne0;
86    int i0 = (n - i3*ne2*ne1*ne0 - i2*ne1*ne0 - i1*ne0);
87
88    global float * dst_data = (global float *) ((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0);
89
90    for (int i00 = get_local_id(0); i00 < ne00; i00 += get_local_size(0)) {
91        global half * src = (global half *)((global char *) src0 + i03*nb03 + i02*nb02 + i01*nb01 + i00*nb00);
92        dst_data[i00] = src[0];
93    }
94}
95
96kernel void kernel_cpy_f32_f16(
97        global float * src0,
98        ulong offset0,
99        global half * dst,
100        ulong offsetd,
101        int ne00,
102        int ne01,
103        int ne02,
104        int ne03,
105        ulong nb00,
106        ulong nb01,
107        ulong nb02,
108        ulong nb03,
109        int ne0,
110        int ne1,
111        int ne2,
112        int ne3,
113        ulong nb0,
114        ulong nb1,
115        ulong nb2,
116        ulong nb3
117) {
118    src0 = (global float*)((global char*)src0 + offset0);
119    dst = (global half*)((global char*)dst + offsetd);
120
121    int i03 = get_group_id(2);
122    int i02 = get_group_id(1);
123    int i01 = get_group_id(0);
124
125    int n = i03*ne02*ne01*ne00 + i02*ne01*ne00 + i01*ne00;
126
127    int i3 = n / (ne2*ne1*ne0);
128    int i2 = (n - i3*ne2*ne1*ne0) / (ne1*ne0);
129    int i1 = (n - i3*ne2*ne1*ne0 - i2*ne1*ne0) / ne0;
130    int i0 = (n - i3*ne2*ne1*ne0 - i2*ne1*ne0 - i1*ne0);
131
132    global half * dst_data = (global half *) ((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0);
133
134    for (int i00 = get_local_id(0); i00 < ne00; i00 += get_local_size(0)) {
135        global const float * src = (global float *)((global char *) src0 + i03*nb03 + i02*nb02 + i01*nb01 + i00*nb00);
136
137        dst_data[i00] = src[0];
138    }
139}
140
141kernel void kernel_cpy_f32_f32(
142        global float * src0,
143        ulong offset0,
144        global float * dst,
145        ulong offsetd,
146        int ne00,
147        int ne01,
148        int ne02,
149        int ne03,
150        ulong nb00,
151        ulong nb01,
152        ulong nb02,
153        ulong nb03,
154        int ne0,
155        int ne1,
156        int ne2,
157        int ne3,
158        ulong nb0,
159        ulong nb1,
160        ulong nb2,
161        ulong nb3
162) {
163    src0 = (global float*)((global char*)src0 + offset0);
164    dst = (global float*)((global char*)dst + offsetd);
165
166    int i03 = get_group_id(2);
167    int i02 = get_group_id(1);
168    int i01 = get_group_id(0);
169
170    int n = i03*ne02*ne01*ne00 + i02*ne01*ne00 + i01*ne00;
171
172    int i3 = n / (ne2*ne1*ne0);
173    int i2 = (n - i3*ne2*ne1*ne0) / (ne1*ne0);
174    int i1 = (n - i3*ne2*ne1*ne0 - i2*ne1*ne0) / ne0;
175    int i0 = (n - i3*ne2*ne1*ne0 - i2*ne1*ne0 - i1*ne0);
176
177    global float * dst_data = (global float *) ((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0);
178
179    for (int i00 = get_local_id(0); i00 < ne00; i00 += get_local_size(0)) {
180        global const float * src = (global float *)((global char *) src0 + i03*nb03 + i02*nb02 + i01*nb01 + i00*nb00);
181
182        dst_data[i00] = src[0];
183    }
184}
185
186kernel void kernel_cpy_f32_f32_pack(
187        global float * src0,
188        ulong offset0,
189        global float * dst,
190        ulong offsetd,
191        int ne00,
192        int ne01,
193        int ne02,
194        int ne03,
195        ulong nb00,
196        ulong nb01,
197        ulong nb02,
198        ulong nb03,
199        int ne0,
200        int ne1,
201        int ne2,
202        int ne3,
203        ulong nb0,
204        ulong nb1,
205        ulong nb2,
206        ulong nb3
207) {
208    src0 = (global float*)((global char*)src0 + offset0);
209    dst = (global float*)((global char*)dst + offsetd);
210
211    int lsz = get_local_size(0);
212    int tpr = min(ne00, lsz);          // threads per row
213    int rpw = lsz / tpr;               // rows per workgroup
214    int lid = get_local_id(0);
215    int row = get_group_id(0)*rpw + lid / tpr;
216    int lane = lid - (lid / tpr) * tpr;
217
218    int nrows = ne01*ne02*ne03;
219    if (row >= nrows) {
220        return;
221    }
222
223    int i01 = row % ne01;
224    int t   = row / ne01;
225    int i02 = t % ne02;
226    int i03 = t / ne02;
227
228    // linear index of the first element of this row, unflattened over dst dims
229    long n  = (long)row * ne00;
230    int i3  = (int)(n / ((long)ne2*ne1*ne0));
231    long rm = n - (long)i3*ne2*ne1*ne0;
232    int i2  = (int)(rm / ((long)ne1*ne0));
233    rm     -= (long)i2*ne1*ne0;
234    int i1  = (int)(rm / ne0);
235    int i0  = (int)(rm - (long)i1*ne0);
236
237    global float * dst_data = (global float *) ((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0);
238
239    for (int i00 = lane; i00 < ne00; i00 += tpr) {
240        global const float * src = (global float *)((global char *) src0 + i03*nb03 + i02*nb02 + i01*nb01 + i00*nb00);
241        dst_data[i00] = src[0];
242    }
243}
244
245kernel void kernel_cpy_i32_i32(
246        global int * src0,
247        ulong offset0,
248        global int * dst,
249        ulong offsetd,
250        int ne00,
251        int ne01,
252        int ne02,
253        int ne03,
254        ulong nb00,
255        ulong nb01,
256        ulong nb02,
257        ulong nb03,
258        int ne0,
259        int ne1,
260        int ne2,
261        int ne3,
262        ulong nb0,
263        ulong nb1,
264        ulong nb2,
265        ulong nb3
266) {
267    src0 = (global int*)((global char*)src0 + offset0);
268    dst = (global int*)((global char*)dst + offsetd);
269
270    int i03 = get_group_id(2);
271    int i02 = get_group_id(1);
272    int i01 = get_group_id(0);
273
274    int n = i03*ne02*ne01*ne00 + i02*ne01*ne00 + i01*ne00;
275
276    int i3 = n / (ne2*ne1*ne0);
277    int i2 = (n - i3*ne2*ne1*ne0) / (ne1*ne0);
278    int i1 = (n - i3*ne2*ne1*ne0 - i2*ne1*ne0) / ne0;
279    int i0 = (n - i3*ne2*ne1*ne0 - i2*ne1*ne0 - i1*ne0);
280
281    global int * dst_data = (global int *) ((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0);
282
283    for (int i00 = get_local_id(0); i00 < ne00; i00 += get_local_size(0)) {
284        global const int * src = (global int *)((global char *) src0 + i03*nb03 + i02*nb02 + i01*nb01 + i00*nb00);
285
286        dst_data[i00] = src[0];
287    }
288}
289
290// Contiguous f32 copy, one work item per float4 over the whole tensor. The kernels above map
291// one workgroup to each row, which leaves a tensor with few long rows on a single compute unit.
292// vload4/vstore4 rather than a float4 cast: these buffers carry an arbitrary 4-byte view offset.
293kernel void kernel_cpy_f32_f32_flat(
294        global float * src0,
295        ulong offset0,
296        global float * dst,
297        ulong offsetd,
298        ulong ne,
299        ulong n4
300) {
301    src0 = (global float*)((global char*)src0 + offset0);
302    dst  = (global float*)((global char*)dst  + offsetd);
303
304    const ulong i = get_global_id(0);
305
306    if (i < n4) {
307        vstore4(vload4(i, src0), i, dst);
308    } else if (i == n4) {
309        for (ulong t = n4 * 4; t < ne; ++t) {
310            dst[t] = src0[t];
311        }
312    }
313}