#include "opencl_source_map.hpp" namespace MNN { const char* unary = "#ifdef MNN_SUPPORT_FP16\n" "#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n" "#endif\n" "#define GLOBAL_SIZE_3_DIMS "" __private const int global_size_dim0,__private const int global_size_dim1,__private const int global_size_dim2,\n" "#define DEAL_NON_UNIFORM_DIM3(input1, input2, input3) "" if (input1 >= global_size_dim0 || input2 >= global_size_dim1 || input3 >= global_size_dim2) { "" return; "" }\n" "inline float4 gelu(float4 in){\n" " float4 value=0.79788458f*(0.044715f*in*in*in+in);\n" " float4 x2=value*value;\n" " float4 dst=value>(float4)5.0f ? (float4)1.0f : (value <= -(float4)5.0f ? -(float4)1.0f :\n" " (value*(135135.0f+x2*(17325.0f+x2*(378.0f+x2))))/(135135.0f+x2*(62370.0f+x2*(3150.0f+x2*28.0f))));\n" " return (1.0f+dst)*in*0.5f;\n" "}\n" "// MNN: erfinv via TensorFlow's two-branch polynomial (kept in lockstep with\n" "// CPU's UnaryErfinv in source/backend/cpu/UnaryUtils.hpp). Avoids the OpenCL\n" "// CPU-fallback path which is broken when the runtime is in IMAGE memtype.\n" "inline float4 erfinv4(float4 x){\n" " float4 w=-log((float4)1.0f-x*x+(float4)1e-12f);\n" " int4 lt=isless(w,(float4)5.0f);\n" " float4 w_sel=select(sqrt(w)-(float4)3.0f,w-(float4)2.5f,lt);\n" " float4 p_lt=(float4)2.81022636e-08f;\n" " p_lt=(float4)3.43273939e-07f+p_lt*w_sel;\n" " p_lt=(float4)(-3.5233877e-06f)+p_lt*w_sel;\n" " p_lt=(float4)(-4.39150654e-06f)+p_lt*w_sel;\n" " p_lt=(float4)0.00021858087f+p_lt*w_sel;\n" " p_lt=(float4)(-0.00125372503f)+p_lt*w_sel;\n" " p_lt=(float4)(-0.00417768164f)+p_lt*w_sel;\n" " p_lt=(float4)0.246640727f+p_lt*w_sel;\n" " p_lt=(float4)1.50140941f+p_lt*w_sel;\n" " float4 p_ge=(float4)(-0.000200214257f);\n" " p_ge=(float4)0.000100950558f+p_ge*w_sel;\n" " p_ge=(float4)0.00134934322f+p_ge*w_sel;\n" " p_ge=(float4)(-0.00367342844f)+p_ge*w_sel;\n" " p_ge=(float4)0.00573950773f+p_ge*w_sel;\n" " p_ge=(float4)(-0.0076224613f)+p_ge*w_sel;\n" " p_ge=(float4)0.00943887047f+p_ge*w_sel;\n" " p_ge=(float4)1.00167406f+p_ge*w_sel;\n" " p_ge=(float4)2.83297682f+p_ge*w_sel;\n" " return select(p_ge,p_lt,lt)*x;\n" "}\n" "__constant sampler_t SAMPLER=CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;\n" "__kernel void unary(GLOBAL_SIZE_3_DIMS __read_only image2d_t input,__write_only image2d_t output) {\n" " const int channel_block_idx=get_global_id(0);\n" " const int w=get_global_id(1);\n" " const int hb=get_global_id(2);\n" " DEAL_NON_UNIFORM_DIM3(channel_block_idx,w,hb);\n" " const int width=global_size_dim1;\n" " const int pos=mad24(channel_block_idx,width,w);\n" " float4 in=convert_float4(RI_DATA(input,SAMPLER,(int2)(pos,hb)));\n" " OUTPUT_TYPE_I4 out=CONVERT_OUTPUT_I4(OPERATOR);\n" " \n" " WI_DATA(output,(int2)(pos,hb),out);\n" "}\n" ; }