1#pragma OPENCL EXTENSION cl_khr_fp16 : enable
 2
 3//------------------------------------------------------------------------------
 4// silu
 5//------------------------------------------------------------------------------
 6kernel void kernel_silu(
 7        global float * src0,
 8        ulong offset0,
 9        global float * dst,
10        ulong offsetd
11) {
12    src0 = (global float*)((global char*)src0 + offset0);
13    dst = (global float*)((global char*)dst + offsetd);
14
15    float x = src0[get_global_id(0)];
16    dst[get_global_id(0)] = x / (1.0f + exp(-x));
17}
18
19kernel void kernel_silu_4(
20        global 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    float4 x = src0[get_global_id(0)];
29    dst[get_global_id(0)] = x / (1.0f + exp(-x));
30}