wehub-resource-sync
97 lines
4.3 KiB
C++
97 lines
4.3 KiB
C++
#include "opencl_source_map.hpp"
|
|
namespace MNN {
|
|
const char* roi_pooling =
|
|
"#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"
|
|
"#define MIN_VALUE -FLT_MAX\n"
|
|
"// Supported data type: half/float\n"
|
|
"__kernel void roi_pooling(GLOBAL_SIZE_3_DIMS __read_only image2d_t input,__read_only image2d_t roi,\n"
|
|
" __private const int in_height,__private const int in_width,__private const int in_batch,\n"
|
|
" __private const int out_height,__private const int out_width,__private const float spatial_scale,\n"
|
|
" __write_only image2d_t output) {\n"
|
|
" const int out_channel_idx=get_global_id(0);\n"
|
|
" const int out_width_idx=get_global_id(1);\n"
|
|
" const int out_hb_idx=get_global_id(2);\n"
|
|
" DEAL_NON_UNIFORM_DIM3(out_channel_idx,out_width_idx,out_hb_idx);\n"
|
|
" const int roi_batch_idx=out_hb_idx/out_height;\n"
|
|
" const int out_height_idx=out_hb_idx % out_height;\n"
|
|
"#if defined ROI_C1H1W5\n"
|
|
" FLOAT4 roi_0=RI_F(roi,SAMPLER,(int2)(0,roi_batch_idx));\n"
|
|
" int input_batch=roi_0.x;\n"
|
|
" if(input_batch >= in_batch){\n"
|
|
" return;\n"
|
|
" }\n"
|
|
" FLOAT4 roi_1=RI_F(roi,SAMPLER,(int2)(1,roi_batch_idx));\n"
|
|
" FLOAT4 roi_2=RI_F(roi,SAMPLER,(int2)(2,roi_batch_idx));\n"
|
|
" FLOAT4 roi_3=RI_F(roi,SAMPLER,(int2)(3,roi_batch_idx));\n"
|
|
" FLOAT4 roi_4=RI_F(roi,SAMPLER,(int2)(4,roi_batch_idx));\n"
|
|
" int x1=round(roi_1.x*spatial_scale);\n"
|
|
" int y1=round(roi_2.x*spatial_scale);\n"
|
|
" int x2=round(roi_3.x*spatial_scale);\n"
|
|
" int y2=round(roi_4.x*spatial_scale);\n"
|
|
"#elif defined ROI_C5H1W1\n"
|
|
" FLOAT4 roi_0=RI_F(roi,SAMPLER,(int2)(0,roi_batch_idx));\n"
|
|
" int input_batch=roi_0.x;\n"
|
|
" if(input_batch >= in_batch){\n"
|
|
" return;\n"
|
|
" }\n"
|
|
" FLOAT4 roi_1=RI_F(roi,SAMPLER,(int2)(1,roi_batch_idx));\n"
|
|
" int x1=round(roi_0.y*spatial_scale);\n"
|
|
" int y1=round(roi_0.z*spatial_scale);\n"
|
|
" int x2=round(roi_0.w*spatial_scale);\n"
|
|
" int y2=round(roi_1.x*spatial_scale);\n"
|
|
"#else\n"
|
|
" const int roi_batch_offset=roi_batch_idx*5;\n"
|
|
" FLOAT4 roi_0=RI_F(roi,SAMPLER,(int2)(0,roi_batch_offset));\n"
|
|
" int input_batch=roi_0.x;\n"
|
|
" if(input_batch >= in_batch){\n"
|
|
" return;\n"
|
|
" }\n"
|
|
" FLOAT4 roi_1=RI_F(roi,SAMPLER,(int2)(0,roi_batch_offset+1));\n"
|
|
" FLOAT4 roi_2=RI_F(roi,SAMPLER,(int2)(0,roi_batch_offset+2));\n"
|
|
" FLOAT4 roi_3=RI_F(roi,SAMPLER,(int2)(0,roi_batch_offset+3));\n"
|
|
" FLOAT4 roi_4=RI_F(roi,SAMPLER,(int2)(0,roi_batch_offset+4));\n"
|
|
" int x1=round(roi_1.x*spatial_scale);\n"
|
|
" int y1=round(roi_2.x*spatial_scale);\n"
|
|
" int x2=round(roi_3.x*spatial_scale);\n"
|
|
" int y2=round(roi_4.x*spatial_scale);\n"
|
|
"#endif\n"
|
|
" int roiW=max(x2-x1+1,1);\n"
|
|
" int roiH=max(y2-y1+1,1);\n"
|
|
" float binSizeW=(float)roiW/(float)out_width;\n"
|
|
" float binSizeH=(float)roiH/(float)out_height;\n"
|
|
" int hStart=min(max(y1+(int)floor(out_height_idx*binSizeH),0),in_height);\n"
|
|
" int hEnd=min(max(y1+(int)ceil((out_height_idx+1)*binSizeH),0),in_height);\n"
|
|
" int hLen=hEnd-hStart;\n"
|
|
" int wStart=min(max(x1+(int)floor(out_width_idx*binSizeW),0),in_width);\n"
|
|
" int wEnd=min(max(x1+(int)ceil((out_width_idx+1)*binSizeW),0),in_width);\n"
|
|
" int wLen=wEnd-wStart;\n"
|
|
" const int pos=mad24(out_channel_idx,out_width,out_width_idx);\n"
|
|
" const FLOAT4 zero_vec=(FLOAT4)(0);\n"
|
|
" if (wLen <= 0 || hLen <= 0) {\n"
|
|
" WI_F(output,(int2)(pos,out_hb_idx),zero_vec);\n"
|
|
" return;\n"
|
|
" }\n"
|
|
" FLOAT4 res=(FLOAT4)(MIN_VALUE);\n"
|
|
" const int in_height_start=hStart;\n"
|
|
" const int in_width_start=wStart;\n"
|
|
" const int in_channel_offset=mul24(out_channel_idx,in_width);\n"
|
|
" const int in_height_offset=mul24(input_batch,in_height);\n"
|
|
" const int batch_idx=mul24(input_batch,in_height);\n"
|
|
" for (int height=0; height<hLen; ++height) {\n"
|
|
" int in_height_idx=in_height_start+height;\n"
|
|
" for (int width=0; width<wLen; ++width) {\n"
|
|
" int in_width_idx=in_width_start+width;\n"
|
|
" FLOAT4 in=RI_F(input,SAMPLER,(int2)(in_channel_offset+in_width_idx,in_height_offset+in_height_idx));\n"
|
|
" res=fmax(res,in);\n"
|
|
" }\n"
|
|
" }\n"
|
|
" WI_F(output,(int2)(pos,out_hb_idx),res);\n"
|
|
"}\n"
|
|
;
|
|
}
|