char-fork/llama.cpp

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

lhezopencl: fix l2_norm (#20480)3b439504b

char/aliceai-t5-moe
brancheschar/aliceai-t5-moeall refs →
2.1 KiB71 linesraw
1#ifdef cl_intel_required_subgroup_size
2#pragma OPENCL EXTENSION cl_intel_required_subgroup_size : enable
3#define INTEL_GPU 1
4#define REQD_SUBGROUP_SIZE_16 __attribute__((intel_reqd_sub_group_size(16)))
5#define REQD_SUBGROUP_SIZE_32 __attribute__((intel_reqd_sub_group_size(32)))
6#elif defined(cl_qcom_reqd_sub_group_size)
7#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
8#define ADRENO_GPU 1
9#define REQD_SUBGROUP_SIZE_64  __attribute__((qcom_reqd_sub_group_size("half")))
10#define REQD_SUBGROUP_SIZE_128 __attribute__((qcom_reqd_sub_group_size("full")))
11#endif
12
13#ifdef INTEL_GPU
14REQD_SUBGROUP_SIZE_32
15#elif defined (ADRENO_GPU)
16REQD_SUBGROUP_SIZE_64
17#endif
18kernel void kernel_l2_norm_f32(
19        global void * src0,
20        ulong offset0,
21        global float * dst,
22        ulong offsetd,
23        int ne00,
24        int ne01,
25        int ne02,
26        int ne03,
27        ulong nb01,
28        ulong nb02,
29        ulong nb03,
30        float eps,
31        local float * sum
32) {
33    src0 = (global void*)((global char*)src0 + offset0);
34    dst = (global float*)((global char*)dst + offsetd);
35
36    int i03 = get_group_id(2);
37    int i02 = get_group_id(1);
38    int i01 = get_group_id(0);
39
40    global float * x = (global float *) ((global char *) src0 + i03*nb03 + i02*nb02 + i01*nb01);
41    global float * y = (global float *) (dst + i03*ne02*ne01*ne00 + i02*ne01*ne00 + i01*ne00);
42
43    float sumf = 0;
44
45    // parallel sum
46    for (int i00 = get_local_id(0); i00 < ne00; i00 += get_local_size(0)) {
47        sumf += x[i00] * x[i00];
48    }
49    sumf = sub_group_reduce_add(sumf);
50
51    if (get_sub_group_local_id() == 0) {
52        sum[get_sub_group_id()] = sumf;
53    }
54
55    barrier(CLK_LOCAL_MEM_FENCE);
56
57    // broadcast
58    for (uint i = get_local_size(0) / get_max_sub_group_size() / 2; i > 0; i /= 2) {
59       if (get_local_id(0) < i) {
60           sum[get_local_id(0)] += sum[get_local_id(0) + i];
61       }
62    }
63
64    barrier(CLK_LOCAL_MEM_FENCE);
65
66    const float scale = 1.0f/max(sqrt(sum[0]), eps);
67
68    for (int i00 = get_local_id(0); i00 < ne00; i00 += get_local_size(0)) {
69        y[i00] = x[i00] * scale;
70    }
71}