wehub-resource-sync
47 lines
1.4 KiB
C++
47 lines
1.4 KiB
C++
#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
|
|
}
|