char-fork/llama.cpp

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

shaofeiqiopencl: refactor expm1 and softplus (#19404)e2f19b320

char/aliceai-t5-moe
brancheschar/aliceai-t5-moeall refs →
3.7 KiB116 linesraw
1#pragma OPENCL EXTENSION cl_khr_fp16 : enable
2
3//------------------------------------------------------------------------------
4// softplus
5//------------------------------------------------------------------------------
6
7kernel void kernel_softplus_f32(
8        global const float * src0,
9        ulong                offset0,
10        global       float * dst,
11        ulong                offsetd
12) {
13    src0 = (global float*)((global char*)src0 + offset0);
14    dst  = (global float*)((global char*)dst + offsetd);
15
16    dst[get_global_id(0)] = (src0[get_global_id(0)] > 20.0f) ? src0[get_global_id(0)] : log(1.0f + exp(src0[get_global_id(0)]));
17}
18
19kernel void kernel_softplus_f32_4(
20        global const float4 * src0,
21        ulong                 offset0,
22        global       float4 * dst,
23        ulong                 offsetd
24) {
25    src0 = (global float4*)((global char*)src0 + offset0);
26    dst  = (global float4*)((global char*)dst + offsetd);
27
28    dst[get_global_id(0)] = (src0[get_global_id(0)] > 20.0f) ? src0[get_global_id(0)] : log(1.0f + exp(src0[get_global_id(0)]));
29}
30
31kernel void kernel_softplus_f16(
32        global const half * src0,
33        ulong               offset0,
34        global       half * dst,
35        ulong               offsetd
36) {
37    src0 = (global half*)((global char*)src0 + offset0);
38    dst  = (global half*)((global char*)dst + offsetd);
39
40    const float x = convert_float(src0[get_global_id(0)]);
41    dst[get_global_id(0)] = convert_half_rte((x > 20.0f) ? x : log(1.0f + exp(x)));
42}
43
44kernel void kernel_softplus_f16_4(
45        global const half4 * src0,
46        ulong                offset0,
47        global       half4 * dst,
48        ulong                offsetd
49) {
50    src0 = (global half4*)((global char*)src0 + offset0);
51    dst  = (global half4*)((global char*)dst + offsetd);
52
53    const float4 x = convert_float4(src0[get_global_id(0)]);
54    dst[get_global_id(0)] = convert_half4_rte((x > 20.0f) ? x : log(1.0f + exp(x)));
55}
56
57kernel void kernel_softplus_f32_nc(
58        global const char * src0,
59        ulong               offset0,
60        global       char * dst,
61        ulong               offsetd,
62        int   ne00,
63        ulong nb00,
64        ulong nb01,
65        ulong nb02,
66        ulong nb03,
67        ulong nb0,
68        ulong nb1,
69        ulong nb2,
70        ulong nb3
71) {
72    src0 = src0 + offset0;
73    dst  = dst + offsetd;
74
75    const int i3 = get_group_id(2);
76    const int i2 = get_group_id(1);
77    const int i1 = get_group_id(0);
78
79    for (int i0 = get_local_id(0); i0 < ne00; i0 += get_local_size(0)) {
80        global const float * x = (global const float *)(src0 + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00);
81        global       float * y = (global       float *)(dst  + i3*nb3  + i2*nb2  + i1*nb1  + i0*nb0);
82
83        *y = (*x > 20.0f) ? *x : log(1.0f + exp(*x));
84    }
85}
86
87kernel void kernel_softplus_f16_nc(
88        global const char * src0,
89        ulong               offset0,
90        global       char * dst,
91        ulong               offsetd,
92        int   ne00,
93        ulong nb00,
94        ulong nb01,
95        ulong nb02,
96        ulong nb03,
97        ulong nb0,
98        ulong nb1,
99        ulong nb2,
100        ulong nb3
101) {
102    src0 = src0 + offset0;
103    dst  = dst + offsetd;
104
105    const int i3 = get_group_id(2);
106    const int i2 = get_group_id(1);
107    const int i1 = get_group_id(0);
108
109    for (int i0 = get_local_id(0); i0 < ne00; i0 += get_local_size(0)) {
110        global const half * hx = (global const half *)(src0 + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00);
111        global       half * hy = (global       half *)(dst  + i3*nb3  + i2*nb2  + i1*nb1  + i0*nb0);
112
113        const float x = convert_float(*hx);
114        *hy = convert_half_rte((x > 20.0f) ? x : log(1.0f + exp(x)));
115    }
116}