wehub-resource-sync
36 lines
1.7 KiB
C++
36 lines
1.7 KiB
C++
#include "opencl_source_map.hpp"
|
|
namespace MNN {
|
|
const char* cast =
|
|
"#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"
|
|
"__constant sampler_t SAMPLER=CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;\n"
|
|
"__kernel void cast(GLOBAL_SIZE_3_DIMS\n"
|
|
" __read_only image2d_t input,\n"
|
|
" __write_only image2d_t output,\n"
|
|
" __private const int width,\n"
|
|
" __private const int height,\n"
|
|
" __private const int channelBlock\n"
|
|
" ) {\n"
|
|
" const int width_idx=get_global_id(0);\n"
|
|
" const int height_idx=get_global_id(1);\n"
|
|
" const int batch_channel_idx=get_global_id(2);\n"
|
|
" DEAL_NON_UNIFORM_DIM3(width_idx,height_idx,batch_channel_idx);\n"
|
|
" \n"
|
|
" const int batch_idx=batch_channel_idx/channelBlock;\n"
|
|
" const int channel_idx=batch_channel_idx % channelBlock;\n"
|
|
" \n"
|
|
"#ifdef TO_BOOL\n"
|
|
" int4 value=convert_int4(RI_DATA(input,SAMPLER,(int2)(channel_idx*width+width_idx,batch_idx*height+height_idx)));\n"
|
|
" value=value == (int4)0 ? (int4)0 : (int4)1;\n"
|
|
" WI_DATA(output,(int2)(channel_idx*width+width_idx,batch_idx*height+height_idx),CONVERT_OUTPUT_I4(value));\n"
|
|
"#else\n"
|
|
" INPUT_TYPE_I4 value=RI_DATA(input,SAMPLER,(int2)(channel_idx*width+width_idx,batch_idx*height+height_idx));\n"
|
|
" WI_DATA(output,(int2)(channel_idx*width+width_idx,batch_idx*height+height_idx),CONVERT_OUTPUT_I4(value));\n"
|
|
"#endif\n"
|
|
"}\n"
|
|
;
|
|
}
|