char-fork/llama.cpp
git clone https://git.t4t.associates/char-fork/llama.cpp
0190529ec
char/aliceai-t5-moe
1#include "common.h" 2 3constant short FC_unary_op [[function_constant(FC_UNARY + 0)]]; 4constant bool FC_unary_cnt[[function_constant(FC_UNARY + 1)]]; 5 6template <typename T0, typename T, typename TC> 7kernel void kernel_unary_impl( 8 constant ggml_metal_kargs_unary & args, 9 device const char * src0, 10 device char * dst, 11 uint3 tgpig[[threadgroup_position_in_grid]], 12 ushort3 tpitg[[thread_position_in_threadgroup]], 13 ushort3 ntg[[threads_per_threadgroup]]) { 14#define FC_OP FC_unary_op 15#define FC_CNT FC_unary_cnt 16 17 device const T0 * src0_ptr; 18 device T * dst_ptr; 19 20 int i0; 21 22 if (FC_CNT) { 23 i0 = tgpig.x; 24 25 src0_ptr = (device const T0 *) (src0); 26 dst_ptr = (device T *) (dst); 27 } else { 28 const int i03 = tgpig.z; 29 const int i02 = tgpig.y; 30 const int k0 = tgpig.x/args.ne01; 31 const int i01 = tgpig.x - k0*args.ne01; 32 33 i0 = k0*ntg.x + tpitg.x; 34 35 src0_ptr = (device const T0 *) (src0 + i03*args.nb03 + i02*args.nb02 + i01*args.nb01); 36 dst_ptr = (device T *) (dst + i03*args.nb3 + i02*args.nb2 + i01*args.nb1 ); 37 } 38 39 { 40 //threadgroup_barrier(mem_flags::mem_none); 41 42 if (!FC_CNT) { 43 if (i0 >= args.ne0) { 44 return; 45 } 46 } 47 48 const TC x = (TC) src0_ptr[i0]; 49 50 if (FC_OP == OP_UNARY_NUM_SCALE) { 51 dst_ptr[i0] = (T) (args.scale * x + args.bias); 52 } 53 54 if (FC_OP == OP_UNARY_NUM_FILL) { 55 dst_ptr[i0] = (T) args.val; 56 } 57 58 if (FC_OP == OP_UNARY_NUM_CLAMP) { 59 dst_ptr[i0] = (T) clamp(x, args.min, args.max); 60 } 61 62 if (FC_OP == OP_UNARY_NUM_SQR) { 63 dst_ptr[i0] = (T) (x * x); 64 } 65 66 if (FC_OP == OP_UNARY_NUM_SQRT) { 67 dst_ptr[i0] = (T) sqrt(x); 68 } 69 70 if (FC_OP == OP_UNARY_NUM_SIN) { 71 dst_ptr[i0] = (T) sin(x); 72 } 73 74 if (FC_OP == OP_UNARY_NUM_COS) { 75 dst_ptr[i0] = (T) cos(x); 76 } 77 78 if (FC_OP == OP_UNARY_NUM_LOG) { 79 dst_ptr[i0] = (T) log(x); 80 } 81 82 if (FC_OP == OP_UNARY_NUM_LEAKY_RELU) { 83 dst_ptr[i0] = (T) (TC(x > 0)*x + TC(x <= 0)*(x * args.slope)); 84 } 85 86 if (FC_OP == OP_UNARY_NUM_TANH) { 87 dst_ptr[i0] = (T) precise::tanh(x); 88 } 89 90 if (FC_OP == OP_UNARY_NUM_RELU) { 91 dst_ptr[i0] = (T) fmax(0, x); 92 } 93 94 if (FC_OP == OP_UNARY_NUM_SIGMOID) { 95 dst_ptr[i0] = (T) (1 / (1 + exp(-x))); 96 } 97 98 if (FC_OP == OP_UNARY_NUM_GELU) { 99 dst_ptr[i0] = (T) (0.5*x*(1 + precise::tanh(SQRT_2_OVER_PI*x*(1 + GELU_COEF_A*x*x)))); 100 } 101 102 if (FC_OP == OP_UNARY_NUM_GELU_ERF) { 103 dst_ptr[i0] = (T) (0.5*x*(1 + erf_approx(SQRT_2_INV*x))); 104 } 105 106 if (FC_OP == OP_UNARY_NUM_GELU_QUICK) { 107 dst_ptr[i0] = (T) (x * (1/(1 + exp(GELU_QUICK_COEF*x)))); 108 } 109 110 if (FC_OP == OP_UNARY_NUM_SILU) { 111 dst_ptr[i0] = (T) (x / (1 + exp(-x))); 112 } 113 114 if (FC_OP == OP_UNARY_NUM_ELU) { 115 dst_ptr[i0] = (T) elu_approx(x); 116 } 117 118 if (FC_OP == OP_UNARY_NUM_NEG) { 119 dst_ptr[i0] = (T) -x; 120 } 121 122 if (FC_OP == OP_UNARY_NUM_ABS) { 123 dst_ptr[i0] = (T) fabs(x); 124 } 125 126 if (FC_OP == OP_UNARY_NUM_SGN) { 127 dst_ptr[i0] = T(x > 0) - T(x < 0); 128 } 129 130 if (FC_OP == OP_UNARY_NUM_STEP) { 131 dst_ptr[i0] = T(x > 0); 132 } 133 134 if (FC_OP == OP_UNARY_NUM_HARDSWISH) { 135 dst_ptr[i0] = (T) (x * fmax(0, fmin(1, x/6 + 0.5))); 136 } 137 138 if (FC_OP == OP_UNARY_NUM_HARDSIGMOID) { 139 dst_ptr[i0] = (T) fmax(0, fmin(1, x/6 + 0.5)); 140 } 141 142 if (FC_OP == OP_UNARY_NUM_EXP) { 143 dst_ptr[i0] = (T) exp(x); 144 } 145 146 if (FC_OP == OP_UNARY_NUM_SOFTPLUS) { 147 dst_ptr[i0] = (T) select(log(1 + exp(x)), x, x > 20); 148 } 149 150 if (FC_OP == OP_UNARY_NUM_EXPM1) { 151 // TODO: precise implementation 152 dst_ptr[i0] = (T) (exp(x) - 1); 153 } 154 155 if (FC_OP == OP_UNARY_NUM_FLOOR) { 156 dst_ptr[i0] = (T) floor(x); 157 } 158 159 if (FC_OP == OP_UNARY_NUM_CEIL) { 160 dst_ptr[i0] = (T) ceil(x); 161 } 162 163 if (FC_OP == OP_UNARY_NUM_ROUND) { 164 dst_ptr[i0] = (T) round(x); 165 } 166 167 if (FC_OP == OP_UNARY_NUM_TRUNC) { 168 dst_ptr[i0] = (T) trunc(x); 169 } 170 171 if (FC_OP == OP_UNARY_NUM_XIELU) { 172 const TC xi = x; 173 const TC gate = TC(xi > TC(0.0f)); 174 const TC clamped = fmin(xi, TC(args.val)); 175 const TC y_pos = TC(args.scale) * xi * xi + TC(args.bias) * xi; 176 const TC y_neg = (exp(clamped) - TC(1.0f) - xi) * TC(args.slope) + TC(args.bias) * xi; 177 dst_ptr[i0] = (T) (gate * y_pos + (TC(1.0f) - gate) * y_neg); 178 } 179 } 180 181#undef FC_OP 182#undef FC_CNT 183} 184 185typedef decltype(kernel_unary_impl<float, float, float>) kernel_unary_t; 186 187template [[host_name("kernel_unary_f32_f32")]] kernel kernel_unary_t kernel_unary_impl<float, float, float>; 188template [[host_name("kernel_unary_f32_f32_4")]] kernel kernel_unary_t kernel_unary_impl<float4, float4, float4>; 189template [[host_name("kernel_unary_f16_f16")]] kernel kernel_unary_t kernel_unary_impl<half, half, float>; 190template [[host_name("kernel_unary_f16_f16_4")]] kernel kernel_unary_t kernel_unary_impl<half4, half4, float4>; 191 192kernel void kernel_silu_back_f32( 193 constant ggml_metal_kargs_silu_back & args, 194 device const float * dy, 195 device const float * x, 196 device float * dx, 197 uint gid [[thread_position_in_grid]]) { 198 if (gid >= args.ne) { 199 return; 200 } 201 202 const float s = 1.0f / (1.0f + exp(-x[gid])); 203 dx[gid] = dy[gid] * s * (1.0f + x[gid] * (1.0f - s)); 204} 205 206template<typename T> 207kernel void kernel_reglu( 208 constant ggml_metal_kargs_glu & args, 209 device const char * src0, 210 device const char * src1, 211 device char * dst, 212 uint tgpig[[threadgroup_position_in_grid]], 213 uint tpitg[[thread_position_in_threadgroup]], 214 uint ntg[[threads_per_threadgroup]]) { 215 device const T * src0_row = (device const T *) ((device const char *) src0 + tgpig*args.nb01) + args.i00; 216 device const T * src1_row = (device const T *) ((device const char *) src1 + tgpig*args.nb11) + args.i10; 217 device T * dst_row = (device T *) ((device char *) dst + tgpig*args.nb1); 218 219 for (int i0 = tpitg; i0 < args.ne0; i0 += ntg) { 220 const float x0 = src0_row[i0]; 221 const float x1 = src1_row[i0]; 222 223 dst_row[i0] = (T)(x0*x1*(x0 > 0.0f)); 224 } 225} 226 227typedef decltype(kernel_reglu<float>) kernel_reglu_t; 228 229template [[host_name("kernel_reglu_f32")]] kernel kernel_reglu_t kernel_reglu<float>; 230template [[host_name("kernel_reglu_f16")]] kernel kernel_reglu_t kernel_reglu<half>; 231 232template<typename T> 233kernel void kernel_geglu( 234 constant ggml_metal_kargs_glu & args, 235 device const char * src0, 236 device const char * src1, 237 device char * dst, 238 uint tgpig[[threadgroup_position_in_grid]], 239 uint tpitg[[thread_position_in_threadgroup]], 240 uint ntg[[threads_per_threadgroup]]) { 241 device const T * src0_row = (device const T *) ((device const char *) src0 + tgpig*args.nb01) + args.i00; 242 device const T * src1_row = (device const T *) ((device const char *) src1 + tgpig*args.nb11) + args.i10; 243 device T * dst_row = (device T *) ((device char *) dst + tgpig*args.nb1); 244 245 for (int i0 = tpitg; i0 < args.ne0; i0 += ntg) { 246 const float x0 = src0_row[i0]; 247 const float x1 = src1_row[i0]; 248 249 const float gelu = 0.5f*x0*(1.0f + precise::tanh(SQRT_2_OVER_PI*x0*(1.0f + GELU_COEF_A*x0*x0))); 250 251 dst_row[i0] = (T)(gelu*x1); 252 } 253} 254 255typedef decltype(kernel_geglu<float>) kernel_geglu_t; 256 257template [[host_name("kernel_geglu_f32")]] kernel kernel_geglu_t kernel_geglu<float>; 258template [[host_name("kernel_geglu_f16")]] kernel kernel_geglu_t kernel_geglu<half>; 259 260template<typename T> 261kernel void kernel_swiglu( 262 constant ggml_metal_kargs_glu & args, 263 device const char * src0, 264 device const char * src1, 265 device char * dst, 266 uint tgpig[[threadgroup_position_in_grid]], 267 uint tpitg[[thread_position_in_threadgroup]], 268 uint ntg[[threads_per_threadgroup]]) { 269 device const T * src0_row = (device const T *) ((device const char *) src0 + tgpig*args.nb01) + args.i00; 270 device const T * src1_row = (device const T *) ((device const char *) src1 + tgpig*args.nb11) + args.i10; 271 device T * dst_row = (device T *) ((device char *) dst + tgpig*args.nb1); 272 273 for (int i0 = tpitg; i0 < args.ne0; i0 += ntg) { 274 const float x0 = src0_row[i0]; 275 const float x1 = src1_row[i0]; 276 277 const float silu = x0 / (1.0f + exp(-x0)); 278 279 dst_row[i0] = (T)(silu*x1); 280 } 281} 282 283typedef decltype(kernel_swiglu<float>) kernel_swiglu_t; 284 285template [[host_name("kernel_swiglu_f32")]] kernel kernel_swiglu_t kernel_swiglu<float>; 286template [[host_name("kernel_swiglu_f16")]] kernel kernel_swiglu_t kernel_swiglu<half>; 287 288template<typename T> 289kernel void kernel_swiglu_oai( 290 constant ggml_metal_kargs_glu & args, 291 device const char * src0, 292 device const char * src1, 293 device char * dst, 294 uint tgpig[[threadgroup_position_in_grid]], 295 uint tpitg[[thread_position_in_threadgroup]], 296 uint ntg[[threads_per_threadgroup]]) { 297 device const T * src0_row = (device const T *) ((device const char *) src0 + tgpig*args.nb01) + args.i00; 298 device const T * src1_row = (device const T *) ((device const char *) src1 + tgpig*args.nb11) + args.i10; 299 device T * dst_row = (device T *) ((device char *) dst + tgpig*args.nb1); 300 301 for (int i0 = tpitg; i0 < args.ne0; i0 += ntg) { 302 float x0 = src0_row[i0]; 303 float x1 = src1_row[i0]; 304 305 x0 = min(x0, args.limit); 306 x1 = max(min(x1, args.limit), -args.limit); 307 308 float out_glu = x0 / (1.0f + exp(-x0 * args.alpha)); 309 out_glu = out_glu * (1.0f + x1); 310 311 dst_row[i0] = (T)out_glu; 312 } 313} 314 315typedef decltype(kernel_swiglu_oai<float>) kernel_swiglu_oai_t; 316 317template [[host_name("kernel_swiglu_oai_f32")]] kernel kernel_swiglu_oai_t kernel_swiglu_oai<float>; 318template [[host_name("kernel_swiglu_oai_f16")]] kernel kernel_swiglu_oai_t kernel_swiglu_oai<half>; 319 320template<typename T> 321kernel void kernel_swiglu_clamp( 322 constant ggml_metal_kargs_glu & args, 323 device const char * src0, 324 device const char * src1, 325 device char * dst, 326 uint tgpig[[threadgroup_position_in_grid]], 327 uint tpitg[[thread_position_in_threadgroup]], 328 uint ntg[[threads_per_threadgroup]]) { 329 device const T * src0_row = (device const T *) ((device const char *) src0 + tgpig*args.nb01) + args.i00; 330 device const T * src1_row = (device const T *) ((device const char *) src1 + tgpig*args.nb11) + args.i10; 331 device T * dst_row = (device T *) ((device char *) dst + tgpig*args.nb1); 332 333 for (int i0 = tpitg; i0 < args.ne0; i0 += ntg) { 334 const float gate = min((float) src0_row[i0], args.limit); 335 const float up = clamp((float) src1_row[i0], -args.limit, args.limit); 336 337 dst_row[i0] = (T)(gate / (1.0f + exp(-gate)) * up); 338 } 339} 340 341typedef decltype(kernel_swiglu_clamp<float>) kernel_swiglu_clamp_t; 342 343template [[host_name("kernel_swiglu_clamp_f32")]] kernel kernel_swiglu_clamp_t kernel_swiglu_clamp<float>; 344template [[host_name("kernel_swiglu_clamp_f16")]] kernel kernel_swiglu_clamp_t kernel_swiglu_clamp<half>; 345 346template<typename T> 347kernel void kernel_geglu_erf( 348 constant ggml_metal_kargs_glu & args, 349 device const char * src0, 350 device const char * src1, 351 device char * dst, 352 uint tgpig[[threadgroup_position_in_grid]], 353 uint tpitg[[thread_position_in_threadgroup]], 354 uint ntg[[threads_per_threadgroup]]) { 355 device const T * src0_row = (device const T *) ((device const char *) src0 + tgpig*args.nb01) + args.i00; 356 device const T * src1_row = (device const T *) ((device const char *) src1 + tgpig*args.nb11) + args.i10; 357 device T * dst_row = (device T *) ((device char *) dst + tgpig*args.nb1); 358 359 for (int i0 = tpitg; i0 < args.ne0; i0 += ntg) { 360 const float x0 = src0_row[i0]; 361 const float x1 = src1_row[i0]; 362 363 const float gelu_erf = 0.5f*x0*(1.0f+erf_approx<float>(x0*SQRT_2_INV)); 364 365 dst_row[i0] = (T)(gelu_erf*x1); 366 } 367} 368 369typedef decltype(kernel_geglu_erf<float>) kernel_geglu_erf_t; 370 371template [[host_name("kernel_geglu_erf_f32")]] kernel kernel_geglu_erf_t kernel_geglu_erf<float>; 372template [[host_name("kernel_geglu_erf_f16")]] kernel kernel_geglu_erf_t kernel_geglu_erf<half>; 373 374template<typename T> 375kernel void kernel_geglu_quick( 376 constant ggml_metal_kargs_glu & args, 377 device const char * src0, 378 device const char * src1, 379 device char * dst, 380 uint tgpig[[threadgroup_position_in_grid]], 381 uint tpitg[[thread_position_in_threadgroup]], 382 uint ntg[[threads_per_threadgroup]]) { 383 device const T * src0_row = (device const T *) ((device const char *) src0 + tgpig*args.nb01) + args.i00; 384 device const T * src1_row = (device const T *) ((device const char *) src1 + tgpig*args.nb11) + args.i10; 385 device T * dst_row = (device T *) ((device char *) dst + tgpig*args.nb1); 386 387 for (int i0 = tpitg; i0 < args.ne0; i0 += ntg) { 388 const float x0 = src0_row[i0]; 389 const float x1 = src1_row[i0]; 390 391 const float gelu_quick = x0*(1.0f/(1.0f+exp(GELU_QUICK_COEF*x0))); 392 393 dst_row[i0] = (T)(gelu_quick*x1); 394 } 395} 396 397typedef decltype(kernel_geglu_quick<float>) kernel_geglu_quick_t; 398 399template [[host_name("kernel_geglu_quick_f32")]] kernel kernel_geglu_quick_t kernel_geglu_quick<float>; 400template [[host_name("kernel_geglu_quick_f16")]] kernel kernel_geglu_quick_t kernel_geglu_quick<half>;