char-fork/llama.cpp

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

shaofeiqiopencl: add optimized q4_1 mm kernel for adreno (#19840)24350fdf9

char/aliceai-t5-moe
brancheschar/aliceai-t5-moeall refs →
4.9 KiB143 linesraw
1#pragma OPENCL EXTENSION cl_khr_fp16 : enable
2
3// 16-bit transpose, loading/storing a 4x4 tile of elements
4kernel void kernel_transpose_16(
5    __read_only image1d_buffer_t input,
6    __write_only image1d_buffer_t output,
7    const uint rows,
8    const uint cols
9) {
10
11    const int i = get_global_id(0);
12    const int j = get_global_id(1);
13    const int i_2 = i<<2;
14    const int j_2 = j<<2;
15
16    half4 temp0 = read_imageh(input, (j_2+0)*cols+i);
17    half4 temp1 = read_imageh(input, (j_2+1)*cols+i);
18    half4 temp2 = read_imageh(input, (j_2+2)*cols+i);
19    half4 temp3 = read_imageh(input, (j_2+3)*cols+i);
20
21    write_imageh(output, (i_2+0)*rows+j, (half4)(temp0.s0, temp1.s0, temp2.s0, temp3.s0));
22    write_imageh(output, (i_2+1)*rows+j, (half4)(temp0.s1, temp1.s1, temp2.s1, temp3.s1));
23    write_imageh(output, (i_2+2)*rows+j, (half4)(temp0.s2, temp1.s2, temp2.s2, temp3.s2));
24    write_imageh(output, (i_2+3)*rows+j, (half4)(temp0.s3, temp1.s3, temp2.s3, temp3.s3));
25}
26
27// Padded kernel for irregular shape
28kernel void kernel_transpose_16_4x1(
29    __read_only image1d_buffer_t input,
30    __write_only image1d_buffer_t output,
31    const uint rows,
32    const uint cols
33) {
34
35    const int i = get_global_id(0);
36    const int j = get_global_id(1);
37    const int j_2 = j << 2;
38
39    half temp0 = read_imageh(input, (j_2 + 0) * cols + i).x;
40    half temp1 = read_imageh(input, (j_2 + 1) * cols + i).x;
41    half temp2 = read_imageh(input, (j_2 + 2) * cols + i).x;
42    half temp3 = read_imageh(input, (j_2 + 3) * cols + i).x;
43
44    write_imageh(output, i * rows + j, (half4)(temp0, temp1, temp2, temp3));
45}
46
47// Transpose treating each element as 8-bit using buffer
48kernel void kernel_transpose_8_buf(
49    global const uchar * input,
50    global uchar * output,
51    const int ldi,
52    const int ldo
53) {
54    const int x = get_global_id(0);
55    const int y = get_global_id(1);
56
57    output[x*ldo + y] = input[y*ldi + x];
58}
59
60// Transpose treating each element as 16-bit using buffer
61kernel void kernel_transpose_16_buf(
62    global const ushort * input,
63    global ushort * output,
64    const int ldi,
65    const int ldo
66) {
67    const int x = get_global_id(0);
68    const int y = get_global_id(1);
69
70    output[x*ldo + y] = input[y*ldi + x];
71}
72
73// Transpose treating each element as 32-bit using buffer
74kernel void kernel_transpose_32_buf(
75    global const uint * input,
76    global uint * output,
77    const int ldi,
78    const int ldo
79) {
80    const int x = get_global_id(0);
81    const int y = get_global_id(1);
82
83    output[x*ldo + y] = input[y*ldi + x];
84}
85
86// 32-bit transpose, loading/storing a 4x4 tile of elements
87kernel void kernel_transpose_32(
88    __read_only image1d_buffer_t input,
89    __write_only image1d_buffer_t output,
90    const uint rows,
91    const uint cols
92) {
93
94    const int i = get_global_id(0);
95    const int j = get_global_id(1);
96    const int i_2 = i<<2;
97    const int j_2 = j<<2;
98
99    float4 temp0 = read_imagef(input, (j_2+0)*cols+i);
100    float4 temp1 = read_imagef(input, (j_2+1)*cols+i);
101    float4 temp2 = read_imagef(input, (j_2+2)*cols+i);
102    float4 temp3 = read_imagef(input, (j_2+3)*cols+i);
103
104    write_imagef(output, (i_2+0)*rows+j, (float4)(temp0.s0, temp1.s0, temp2.s0, temp3.s0));
105    write_imagef(output, (i_2+1)*rows+j, (float4)(temp0.s1, temp1.s1, temp2.s1, temp3.s1));
106    write_imagef(output, (i_2+2)*rows+j, (float4)(temp0.s2, temp1.s2, temp2.s2, temp3.s2));
107    write_imagef(output, (i_2+3)*rows+j, (float4)(temp0.s3, temp1.s3, temp2.s3, temp3.s3));
108
109}
110
111// 32-bit transpose, loading/storing a 4x4 tile of elements
112// Only used for activations
113// converts to FP16
114// also adds zero padding for non multiple of 8 prompt lengths
115kernel void kernel_transpose_32_16(__read_only image1d_buffer_t input, __write_only image1d_buffer_t output, const uint rows, const uint cols, const uint padded_rows) {
116
117    const int i = get_global_id(0);
118    const int j = get_global_id(1);
119    const int i_2 = i<<2;
120    const int j_2 = j<<2;
121    half4 temp0 = {0,0,0,0}; // initialize outputs to 0
122    half4 temp1 = {0,0,0,0};
123    half4 temp2 = {0,0,0,0};
124    half4 temp3 = {0,0,0,0};
125
126    if((j_2+0)*cols+i*4+3 < rows*cols*16){ // only load from a valid location. Otherwise keep register data as 0
127        temp0 = read_imageh(input, (j_2+0)*cols+i);
128    }
129    if((j_2+1)*cols+i*4+3 < rows*cols*16){
130        temp1 = read_imageh(input, (j_2+1)*cols+i);
131    }
132    if((j_2+2)*cols+i*4+3 < rows*cols*16){
133        temp2 = read_imageh(input, (j_2+2)*cols+i);
134    }
135    if((j_2+3)*cols+i*4+3 < rows*cols*16){
136        temp3 = read_imageh(input, (j_2+3)*cols+i);
137    }
138
139    write_imageh(output, (i_2+0)*padded_rows+j, (half4)(temp0.s0, temp1.s0, temp2.s0, temp3.s0)); // no conditionals for output, includes zero padding
140    write_imageh(output, (i_2+1)*padded_rows+j, (half4)(temp0.s1, temp1.s1, temp2.s1, temp3.s1));
141    write_imageh(output, (i_2+2)*padded_rows+j, (half4)(temp0.s2, temp1.s2, temp2.s2, temp3.s2));
142    write_imageh(output, (i_2+3)*padded_rows+j, (half4)(temp0.s3, temp1.s3, temp2.s3, temp3.s3));
143}