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.2 KiB113 linesraw
1#pragma OPENCL EXTENSION cl_khr_fp16 : enable
2
3//------------------------------------------------------------------------------
4// expm1
5//------------------------------------------------------------------------------
6
7kernel void kernel_expm1_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)] = exp(src0[get_global_id(0)]) - 1.0f;
17}
18
19kernel void kernel_expm1_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)] = exp(src0[get_global_id(0)]) - 1.0f;
29}
30
31kernel void kernel_expm1_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    dst[get_global_id(0)] = exp(src0[get_global_id(0)]) - 1.0h;
41}
42
43kernel void kernel_expm1_f16_4(
44        global const half4 * src0,
45        ulong                offset0,
46        global       half4 * dst,
47        ulong                offsetd
48) {
49    src0 = (global half4*)((global char*)src0 + offset0);
50    dst  = (global half4*)((global char*)dst + offsetd);
51
52    dst[get_global_id(0)] = exp(src0[get_global_id(0)]) - 1.0h;
53}
54
55kernel void kernel_expm1_f32_nc(
56        global const char * src0,
57        ulong               offset0,
58        global       char * dst,
59        ulong               offsetd,
60        int   ne00,
61        ulong nb00,
62        ulong nb01,
63        ulong nb02,
64        ulong nb03,
65        ulong nb0,
66        ulong nb1,
67        ulong nb2,
68        ulong nb3
69) {
70    src0 = src0 + offset0;
71    dst  = dst + offsetd;
72
73    const int i3 = get_group_id(2);
74    const int i2 = get_group_id(1);
75    const int i1 = get_group_id(0);
76
77    for (int i0 = get_local_id(0); i0 < ne00; i0 += get_local_size(0)) {
78        global const float * x = (global const float *)(src0 + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00);
79        global       float * y = (global       float *)(dst  + i3*nb3  + i2*nb2  + i1*nb1  + i0*nb0);
80
81        *y = exp(*x) - 1.0f;
82    }
83}
84
85kernel void kernel_expm1_f16_nc(
86        global const char * src0,
87        ulong               offset0,
88        global       char * dst,
89        ulong               offsetd,
90        int   ne00,
91        ulong nb00,
92        ulong nb01,
93        ulong nb02,
94        ulong nb03,
95        ulong nb0,
96        ulong nb1,
97        ulong nb2,
98        ulong nb3
99) {
100    src0 = src0 + offset0;
101    dst  = dst + offsetd;
102
103    const int i3 = get_group_id(2);
104    const int i2 = get_group_id(1);
105    const int i1 = get_group_id(0);
106
107    for (int i0 = get_local_id(0); i0 < ne00; i0 += get_local_size(0)) {
108        global const half * x = (global const half *)(src0 + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00);
109        global       half * y = (global       half *)(dst  + i3*nb3  + i2*nb2  + i1*nb1  + i0*nb0);
110
111        *y = exp(*x) - 1.0f;
112    }
113}