#include "opencl_source_map.hpp" 
namespace MNN { 
const char* select = 
"#ifdef MNN_SUPPORT_FP16\n"
"#pragma OPENCL EXTENSION cl_khr_fp16 : enable\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"
"__constant sampler_t SAMPLER=CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;\n"
"__kernel void select_img(GLOBAL_SIZE_2_DIMS\n"
" __read_only image2d_t input,\n"
" __read_only image2d_t input0,\n"
" __read_only image2d_t input1,\n"
" __write_only image2d_t output\n"
" ) {\n"
" const int idx=get_global_id(0);\n"
" const int idy=get_global_id(1);\n"
" DEAL_NON_UNIFORM_DIM2(idx,idy);\n"
" int4 select_vec=read_imagei(input,SAMPLER,(int2)(idx,idy));\n"
"#ifdef INSIZE1_EUQAL_1\n"
" FLOAT4 in0=RI_F(input0,SAMPLER,(int2)(0,0));\n"
" in0=(FLOAT4)(in0.x);\n"
"#else\n"
" FLOAT4 in0=RI_F(input0,SAMPLER,(int2)(idx,idy));\n"
"#endif\n"
" \n"
"#ifdef INSIZE2_EUQAL_1\n"
" FLOAT4 in1=RI_F(input1,SAMPLER,(int2)(0,0));\n"
" in1=(FLOAT4)(in1.x);\n"
"#else\n"
" FLOAT4 in1=RI_F(input1,SAMPLER,(int2)(idx,idy));\n"
"#endif\n"
" FLOAT4 out=select(in1,in0,CONVERT_FLOAT4(select_vec) == (FLOAT4)(1));\n"
" WI_F(output,(int2)(idx,idy),out);\n"
"}\n"
;
}