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 →
6.6 KiB128 linesraw
1// concat is a pure copy, so the kernels are keyed by element byte size
2// (1/2/4/8) rather than logical type, matching the CUDA backend.
3
4#define KERNEL_CONCAT(SUFFIX, T)                                                     \
5kernel void kernel_concat_##SUFFIX(                                                  \
6    global  const char * src0,                                                       \
7    ulong                offset0,                                                    \
8    global  const char * src1,                                                       \
9    ulong                offset1,                                                    \
10    global        char * 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    ulong           nb10,                                                            \
21    ulong           nb11,                                                            \
22    ulong           nb12,                                                            \
23    ulong           nb13,                                                            \
24    int             ne0,                                                             \
25    ulong           nb0,                                                             \
26    ulong           nb1,                                                             \
27    ulong           nb2,                                                             \
28    ulong           nb3,                                                             \
29    int             dim                                                              \
30) {                                                                                  \
31    src0 = src0 + offset0;                                                           \
32    src1 = src1 + offset1;                                                           \
33    dst  = dst  + offsetd;                                                           \
34                                                                                     \
35    const int i3 = get_group_id(2);                                                  \
36    const int i2 = get_group_id(1);                                                  \
37    const int i1 = get_group_id(0);                                                  \
38                                                                                     \
39    int o[4] = {0, 0, 0, 0};                                                         \
40    o[dim] = dim == 0 ? ne00 : (dim == 1 ? ne01 : (dim == 2 ? ne02 : ne03));         \
41                                                                                     \
42    global const T * x;                                                              \
43                                                                                     \
44    for (int i0 = get_local_id(0); i0 < ne0; i0 += get_local_size(0)) {              \
45        if (i0 < ne00 && i1 < ne01 && i2 < ne02 && i3 < ne03) {                      \
46            x = (global const T *)(src0 + (i3       )*nb03 + (i2       )*nb02 + (i1       )*nb01 + (i0       )*nb00); \
47        } else {                                                                     \
48            x = (global const T *)(src1 + (i3 - o[3])*nb13 + (i2 - o[2])*nb12 + (i1 - o[1])*nb11 + (i0 - o[0])*nb10); \
49        }                                                                            \
50                                                                                     \
51        global T * y = (global T *)(dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0);        \
52                                                                                     \
53        *y = *x;                                                                     \
54    }                                                                                \
55}
56
57KERNEL_CONCAT(b1, char)
58KERNEL_CONCAT(b2, short)
59KERNEL_CONCAT(b4, int)
60KERNEL_CONCAT(b8, long)
61
62// packed variant for the common dim==0, small-ne0 case (4-byte elements only).
63kernel void kernel_concat_b4_pack(
64    global  const char * src0,
65    ulong                offset0,
66    global  const char * src1,
67    ulong                offset1,
68    global        char * dst,
69    ulong                offsetd,
70    int             ne00,
71    int             ne01,
72    int             ne02,
73    int             ne03,
74    ulong           nb00,
75    ulong           nb01,
76    ulong           nb02,
77    ulong           nb03,
78    ulong           nb10,
79    ulong           nb11,
80    ulong           nb12,
81    ulong           nb13,
82    int             ne0,
83    ulong           nb0,
84    ulong           nb1,
85    ulong           nb2,
86    ulong           nb3,
87    int             dim,
88    int             ne1,
89    int             ne2,
90    int             ne3
91) {
92    src0 = src0 + offset0;
93    src1 = src1 + offset1;
94    dst  = dst  + offsetd;
95
96    int lsz = get_local_size(0);
97    int tpr = min(ne0, lsz);          // threads per row
98    int rpw = lsz / tpr;              // rows per workgroup
99    int lid = get_local_id(0);
100    int row = get_group_id(0)*rpw + lid / tpr;
101    int lane = lid - (lid / tpr) * tpr;
102
103    int nrows = ne1*ne2*ne3;
104    if (row >= nrows) {
105        return;
106    }
107
108    int i1 = row % ne1;
109    int t  = row / ne1;
110    int i2 = t % ne2;
111    int i3 = t / ne2;
112
113    int o[4] = {0, 0, 0, 0};
114    o[dim] = dim == 0 ? ne00 : (dim == 1 ? ne01 : (dim == 2 ? ne02 : ne03));
115
116    for (int i0 = lane; i0 < ne0; i0 += tpr) {
117        global const int * x;
118        if (i0 < ne00 && i1 < ne01 && i2 < ne02 && i3 < ne03) {
119            x = (global const int *)(src0 + (i3       )*nb03 + (i2       )*nb02 + (i1       )*nb01 + (i0       )*nb00);
120        } else {
121            x = (global const int *)(src1 + (i3 - o[3])*nb13 + (i2 - o[2])*nb12 + (i1 - o[1])*nb11 + (i0 - o[0])*nb10);
122        }
123
124        global int * y = (global int *)(dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0);
125
126        *y = *x;
127    }
128}