#include "opencl_source_map.hpp" 
namespace MNN { 
#ifndef MNN_OPENCL_BUFFER_CLOSED
const char* cast_buf = 
"#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"
"__kernel void cast_buf(GLOBAL_SIZE_2_DIMS\n"
" __global INPUT_TYPE* input,\n"
" __global OUTPUT_TYPE* output,\n"
" __private const int size\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"
" const int inp_offset=idx*4;\n"
"#ifdef PACK_LEAVE\n"
" if(inp_offset+3 >= size){\n"
" int remain=size-inp_offset;\n"
" for(int i=0; i<remain; ++i){\n"
" #ifdef TO_BOOL\n"
" int value=(int)input[inp_offset+i];\n"
" value=value == 0 ? 0 : 1;\n"
" output[inp_offset+i]=(OUTPUT_TYPE)value;\n"
" #else\n"
" output[inp_offset+i]=(OUTPUT_TYPE)input[inp_offset+i];\n"
" #endif\n"
" }\n"
" }else {\n"
"#endif\n"
" #ifdef TO_BOOL\n"
" int4 value=convert_int4(vload4(0,input+inp_offset));\n"
" value=value == (int4)0 ? (int4)0 : (int4)1;\n"
" vstore4(CONVERT_OUTPUT4(value),0,output+inp_offset);\n"
" #else\n"
" vstore4(CONVERT_OUTPUT4(vload4(0,input+inp_offset)),0,output+inp_offset);\n"
" #endif\n"
"#ifdef PACK_LEAVE\n"
" }\n"
"#endif\n"
"}\n"
;
#endif
}