char-fork/llama.cpp
git clone https://git.t4t.associates/char-fork/llama.cpp
1548a240e
char/aliceai-t5-moe
1#pragma OPENCL EXTENSION cl_khr_fp16 : enable 2 3//------------------------------------------------------------------------------ 4// Extended elementwise unary ops, same variant shape as abs.cl: 5// f32, f32_4 (vec4), f16, f16_4 (vec4), f32_nc, f16_nc (stride-addressed). 6// 7// sgn, step, elu, hardswish, hardsigmoid, floor, ceil, round, trunc. 8// 9// Semantics match the ggml CPU reference (ggml.c). Values are computed in float 10// (the f16 variants read/write half and convert), so the conditional ops match 11// the CPU bit-for-bit within tolerance. SEXPR is the scalar form, VEXPR the 12// float4 form (vector ternaries need select()). 13//------------------------------------------------------------------------------ 14 15#define UNARY_EXT(NAME, SEXPR, VEXPR)\ 16kernel void kernel_##NAME##_f32(\ 17 global const float * src0, ulong offset0,\ 18 global float * dst, ulong offsetd) {\ 19 src0 = (global float*)((global char*)src0 + offset0);\ 20 dst = (global float*)((global char*)dst + offsetd);\ 21 float x = src0[get_global_id(0)];\ 22 dst[get_global_id(0)] = (SEXPR);\ 23}\ 24kernel void kernel_##NAME##_f32_4(\ 25 global const float4 * src0, ulong offset0,\ 26 global float4 * dst, ulong offsetd) {\ 27 src0 = (global float4*)((global char*)src0 + offset0);\ 28 dst = (global float4*)((global char*)dst + offsetd);\ 29 float4 x = src0[get_global_id(0)];\ 30 dst[get_global_id(0)] = (VEXPR);\ 31}\ 32kernel void kernel_##NAME##_f16(\ 33 global const half * src0, ulong offset0,\ 34 global half * dst, ulong offsetd) {\ 35 src0 = (global half*)((global char*)src0 + offset0);\ 36 dst = (global half*)((global char*)dst + offsetd);\ 37 float x = src0[get_global_id(0)];\ 38 dst[get_global_id(0)] = (SEXPR);\ 39}\ 40kernel void kernel_##NAME##_f16_4(\ 41 global const half4 * src0, ulong offset0,\ 42 global half4 * dst, ulong offsetd) {\ 43 src0 = (global half4*)((global char*)src0 + offset0);\ 44 dst = (global half4*)((global char*)dst + offsetd);\ 45 float4 x = convert_float4(src0[get_global_id(0)]);\ 46 dst[get_global_id(0)] = convert_half4(VEXPR);\ 47}\ 48kernel void kernel_##NAME##_f32_nc(\ 49 global const char * src0, ulong offset0,\ 50 global char * dst, ulong offsetd,\ 51 int ne00, ulong nb00, ulong nb01, ulong nb02, ulong nb03,\ 52 ulong nb0, ulong nb1, ulong nb2, ulong nb3) {\ 53 src0 = src0 + offset0; dst = dst + offsetd;\ 54 const int i3 = get_group_id(2);\ 55 const int i2 = get_group_id(1);\ 56 const int i1 = get_group_id(0);\ 57 for (int i0 = get_local_id(0); i0 < ne00; i0 += get_local_size(0)) {\ 58 float x = *(global const float *)(src0 + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00);\ 59 *(global float *)(dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0) = (SEXPR);\ 60 }\ 61}\ 62kernel void kernel_##NAME##_f16_nc(\ 63 global const char * src0, ulong offset0,\ 64 global char * dst, ulong offsetd,\ 65 int ne00, ulong nb00, ulong nb01, ulong nb02, ulong nb03,\ 66 ulong nb0, ulong nb1, ulong nb2, ulong nb3) {\ 67 src0 = src0 + offset0; dst = dst + offsetd;\ 68 const int i3 = get_group_id(2);\ 69 const int i2 = get_group_id(1);\ 70 const int i1 = get_group_id(0);\ 71 for (int i0 = get_local_id(0); i0 < ne00; i0 += get_local_size(0)) {\ 72 float x = *(global const half *)(src0 + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00);\ 73 *(global half *)(dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0) = (SEXPR);\ 74 }\ 75} 76 77UNARY_EXT(sgn, sign(x), sign(x)) 78UNARY_EXT(step, x > 0.0f ? 1.0f : 0.0f, select((float4)0.0f, (float4)1.0f, x > 0.0f)) 79UNARY_EXT(elu, x > 0.0f ? x : expm1(x), select(expm1(x), x, x > 0.0f)) 80UNARY_EXT(hardswish, x * fmin(1.0f, fmax(0.0f, (x + 3.0f) / 6.0f)), x * fmin((float4)1.0f, fmax((float4)0.0f, (x + 3.0f) / 6.0f))) 81UNARY_EXT(hardsigmoid, fmin(1.0f, fmax(0.0f, (x + 3.0f) / 6.0f)), fmin((float4)1.0f, fmax((float4)0.0f, (x + 3.0f) / 6.0f))) 82UNARY_EXT(floor, floor(x), floor(x)) 83UNARY_EXT(ceil, ceil(x), ceil(x)) 84UNARY_EXT(round, round(x), round(x)) 85UNARY_EXT(trunc, trunc(x), trunc(x))