wehub-resource-sync
80 lines
3.0 KiB
C++
80 lines
3.0 KiB
C++
#include "opencl_source_map.hpp"
|
|
namespace MNN {
|
|
#ifndef MNN_OPENCL_BUFFER_CLOSED
|
|
const char* unary_buf =
|
|
"#ifdef MNN_SUPPORT_FP16\n"
|
|
"#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n"
|
|
"#define CLAMP(a) "" clamp(a,(float4)(-65504.0f),(float4)(65504.0f));\n"
|
|
"#else\n"
|
|
"#define CLAMP(a) a\n"
|
|
"#endif\n"
|
|
"#define GLOBAL_SIZE_2_DIMS "" __private const int global_size_dim0,__private const int global_size_dim1,\n"
|
|
"#define DEAL_NON_UNIFORM_DIM2(input1, input2) "" if (input1 >= global_size_dim0 || input2 >= global_size_dim1) { "" 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"
|
|
"__kernel void unary_buf(GLOBAL_SIZE_2_DIMS\n"
|
|
" __global const INPUT_TYPE *input,\n"
|
|
" __global OUTPUT_TYPE *output,\n"
|
|
" __private const int size) {\n"
|
|
" const int x=get_global_id(0);\n"
|
|
" const int y=get_global_id(1);\n"
|
|
" DEAL_NON_UNIFORM_DIM2(x,y);\n"
|
|
" const int offset=x << 2;\n"
|
|
"#ifdef PACK_LEAVE\n"
|
|
" if(offset+3 >= size){\n"
|
|
" int remain=size-offset;\n"
|
|
" float4 in;\n"
|
|
" float* in_ptr=(float*)∈\n"
|
|
" for(int i=0; i<remain; ++i){\n"
|
|
" in_ptr[i]=(float)input[offset+i];\n"
|
|
" }\n"
|
|
" float4 out=OPERATOR;\n"
|
|
" float* out_ptr=(float*)&out;\n"
|
|
" for(int i=0; i<remain; ++i){\n"
|
|
" output[offset+i]=(OUTPUT_TYPE)out_ptr[i];\n"
|
|
" }\n"
|
|
" }else {\n"
|
|
"#endif\n"
|
|
" float4 in=convert_float4(vload4(0,input+offset));\n"
|
|
" float4 out=CLAMP(OPERATOR);\n"
|
|
" vstore4(CONVERT_OUTPUT4(out),0,output+offset);\n"
|
|
"#ifdef PACK_LEAVE\n"
|
|
" }\n"
|
|
"#endif\n"
|
|
"}\n"
|
|
;
|
|
#endif
|
|
}
|