char-fork/llama.cpp

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

lhezopencl: add neg, exp and diag (#20127)6c97bffd6

char/aliceai-t5-moe
brancheschar/aliceai-t5-moeall refs →
3.3 KiB125 linesraw
1#pragma OPENCL EXTENSION cl_khr_fp16 : enable
2
3kernel void kernel_neg_f32(
4        global const float * src0,
5        ulong                offset0,
6        global       float * dst,
7        ulong                offsetd,
8        int                  n
9) {
10    if (get_global_id(0) >= n) {
11        return;
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)];
17}
18
19kernel void kernel_neg_f32_4(
20        global const float4 * src0,
21        ulong                 offset0,
22        global       float4 * dst,
23        ulong                 offsetd,
24        int                   n
25) {
26    if (get_global_id(0) >= n) {
27        return;
28    }
29    src0 = (global float4*)((global char*)src0 + offset0);
30    dst  = (global float4*)((global char*)dst + offsetd);
31
32    dst[get_global_id(0)] = -src0[get_global_id(0)];
33}
34
35kernel void kernel_neg_f16(
36        global const half * src0,
37        ulong               offset0,
38        global       half * dst,
39        ulong               offsetd,
40        int                 n
41) {
42    if (get_global_id(0) >= n) {
43        return;
44    }
45    src0 = (global half*)((global char*)src0 + offset0);
46    dst  = (global half*)((global char*)dst + offsetd);
47
48    dst[get_global_id(0)] = -src0[get_global_id(0)];
49}
50
51kernel void kernel_neg_f16_4(
52        global const half4 * src0,
53        ulong                offset0,
54        global       half4 * dst,
55        ulong                offsetd,
56        int                  n
57) {
58    if (get_global_id(0) >= n) {
59        return;
60    }
61    src0 = (global half4*)((global char*)src0 + offset0);
62    dst  = (global half4*)((global char*)dst + offsetd);
63
64    dst[get_global_id(0)] = -src0[get_global_id(0)];
65}
66
67kernel void kernel_neg_f32_nc(
68        global const char * src0,
69        ulong               offset0,
70        global       char * dst,
71        ulong               offsetd,
72        int   ne00,
73        ulong nb00,
74        ulong nb01,
75        ulong nb02,
76        ulong nb03,
77        ulong nb0,
78        ulong nb1,
79        ulong nb2,
80        ulong nb3
81) {
82    src0 = src0 + offset0;
83    dst  = dst + offsetd;
84
85    const int i3 = get_group_id(2);
86    const int i2 = get_group_id(1);
87    const int i1 = get_group_id(0);
88
89    for (int i0 = get_local_id(0); i0 < ne00; i0 += get_local_size(0)) {
90        global const float * x = (global const float *)(src0 + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00);
91        global       float * y = (global       float *)(dst  + i3*nb3  + i2*nb2  + i1*nb1  + i0*nb0);
92
93        *y = -*x;
94    }
95}
96
97kernel void kernel_neg_f16_nc(
98        global const char * src0,
99        ulong               offset0,
100        global       char * dst,
101        ulong               offsetd,
102        int   ne00,
103        ulong nb00,
104        ulong nb01,
105        ulong nb02,
106        ulong nb03,
107        ulong nb0,
108        ulong nb1,
109        ulong nb2,
110        ulong nb3
111) {
112    src0 = src0 + offset0;
113    dst  = dst + offsetd;
114
115    const int i3 = get_group_id(2);
116    const int i2 = get_group_id(1);
117    const int i1 = get_group_id(0);
118
119    for (int i0 = get_local_id(0); i0 < ne00; i0 += get_local_size(0)) {
120        global const half * x = (global const half *)(src0 + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00);
121        global       half * y = (global       half *)(dst  + i3*nb3  + i2*nb2  + i1*nb1  + i0*nb0);
122
123        *y = -*x;
124    }
125}