Files
2026-07-13 13:33:03 +08:00

140 lines
7.5 KiB
Common Lisp

#ifdef MNN_SUPPORT_FP16
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
#endif
#define GLOBAL_SIZE_3_DIMS \
__private const int global_size_dim0, __private const int global_size_dim1, __private const int global_size_dim2,
#define DEAL_NON_UNIFORM_DIM3(input1, input2, input3) \
if (input1 >= global_size_dim0 || input2 >= global_size_dim1 || input3 >= global_size_dim2) { \
return; \
}
__kernel void nearest_buf(GLOBAL_SIZE_3_DIMS __global const FLOAT* input,
__global FLOAT* output,
__private const float height_scale,
__private const float width_scale,
__private const float height_offset,
__private const float width_offset,
__private const int input_height,
__private const int input_width,
__private const int out_height,
__private const int out_width,
__private const int batch) {
const int output_channel_block_idx = get_global_id(0);
const int output_width_block_idx = get_global_id(1);
const int output_batch_height_block_idx = get_global_id(2);
DEAL_NON_UNIFORM_DIM3(output_channel_block_idx, output_width_block_idx, output_batch_height_block_idx);
const int output_batch_idx = output_batch_height_block_idx / out_height;
const int output_height_idx = output_batch_height_block_idx % out_height;
const float in_h_idx = output_height_idx * height_scale + height_offset;
const float in_w_idx = output_width_block_idx * width_scale + width_offset;
#ifdef USE_ROUND
const int in_h_index = min(max(0, (int)floor(in_h_idx + 0.499f)), input_height-1);
const int in_w_index = min(max(0, (int)floor(in_w_idx + 0.499f)), input_width-1);
#else
const int in_h_index = min(max(0, (int)floor(in_h_idx)), input_height-1);
const int in_w_index = min(max(0, (int)floor(in_w_idx)), input_width-1);
#endif
const int inp_offset = ((output_batch_idx + output_channel_block_idx*batch) * input_height + in_h_index) * input_width + in_w_index;
FLOAT4 value = vload4(inp_offset, input);
const int out_offset = ((output_batch_idx + output_channel_block_idx*batch) * out_height + output_height_idx) * out_width + output_width_block_idx;
vstore4(value, out_offset, output);
}
__kernel void bilinear_buf(GLOBAL_SIZE_3_DIMS __global const FLOAT* input,
__global FLOAT* output,
__private const float height_scale,
__private const float width_scale,
__private const float height_offset,
__private const float width_offset,
__private const int input_height,
__private const int input_width,
__private const int out_height,
__private const int out_width,
__private const int batch) {
const int output_channel_block_idx = get_global_id(0);
const int output_width_block_idx = get_global_id(1);
const int output_batch_height_block_idx = get_global_id(2);
DEAL_NON_UNIFORM_DIM3(output_channel_block_idx, output_width_block_idx, output_batch_height_block_idx);
const int output_batch_idx = output_batch_height_block_idx / out_height;
const int output_height_idx = output_batch_height_block_idx % out_height;
const float in_h_idx = output_height_idx * height_scale + height_offset;
const float in_w_idx = output_width_block_idx * width_scale + width_offset;
const int in_h0_index = min(max(0, (int)floor(in_h_idx)), input_height-1);
const int in_w0_index = min(max(0, (int)floor(in_w_idx)), input_width-1);
const int in_h1_index = min(max(0, (int)floor(in_h_idx)+1), input_height-1);
const int in_w1_index = min(max(0, (int)floor(in_w_idx)+1), input_width-1);
float factor_w = (in_w_idx - (int)floor(in_w_idx));
float factor_h = (in_h_idx - (int)floor(in_h_idx));
const int inp_offset_base = (output_batch_idx + output_channel_block_idx*batch) * input_height;
const int inp_offset_00 = (inp_offset_base + in_h0_index) * input_width + in_w0_index;
const int inp_offset_01 = (inp_offset_base + in_h0_index) * input_width + in_w1_index;
const int inp_offset_10 = (inp_offset_base + in_h1_index) * input_width + in_w0_index;
const int inp_offset_11 = (inp_offset_base + in_h1_index) * input_width + in_w1_index;
FLOAT4 value_00 = vload4(inp_offset_00, input);
FLOAT4 value_01 = vload4(inp_offset_01, input);
FLOAT4 value_10 = vload4(inp_offset_10, input);
FLOAT4 value_11 = vload4(inp_offset_11, input);
FLOAT4 value = CONVERT_FLOAT4((float4)((1.0-factor_w)*(1.0-factor_h))*convert_float4(value_00) + (float4)(factor_w*(1.0-factor_h))*convert_float4(value_01) + (float4)((1.0-factor_w)*factor_h)*convert_float4(value_10) + (float4)(factor_w*factor_h)*convert_float4(value_11));
const int out_offset = ((output_batch_idx + output_channel_block_idx*batch) * out_height + output_height_idx) * out_width + output_width_block_idx;
vstore4(value, out_offset, output);
}
__kernel void nearest3D_buf(GLOBAL_SIZE_3_DIMS __global const FLOAT* input,
__global FLOAT* output,
__private const float depth_scale,
__private const float height_scale,
__private const float width_scale,
__private const float depth_offset,
__private const float height_offset,
__private const float width_offset,
__private const int input_depth,
__private const int input_height,
__private const int input_width,
__private const int out_depth,
__private const int out_height,
__private const int out_width,
__private const int batch) {
const int output_channel_block_idx = get_global_id(0);
const int output_height_width_block_idx = get_global_id(1);
const int output_batch_depth_block_idx = get_global_id(2);
DEAL_NON_UNIFORM_DIM3(output_channel_block_idx, output_height_width_block_idx, output_batch_depth_block_idx);
const int output_batch_idx = output_batch_depth_block_idx / out_depth;
const int output_depth_idx = output_batch_depth_block_idx % out_depth;
const int output_height_idx = output_height_width_block_idx / out_height;
const int output_width_idx = output_height_width_block_idx % out_height;
const float in_d_idx = output_depth_idx * depth_scale + depth_offset;
const float in_h_idx = output_height_idx * height_scale + height_offset;
const float in_w_idx = output_width_idx * width_scale + width_offset;
const int in_d_index = min(max(0, (int)floor(in_d_idx)), input_depth-1);
const int in_h_index = min(max(0, (int)floor(in_h_idx)), input_height-1);
const int in_w_index = min(max(0, (int)floor(in_w_idx)), input_width-1);
const int inp_offset = (((output_batch_idx + output_channel_block_idx*batch)
* input_depth + in_d_index) * input_height + in_h_index) * input_width + in_w_index;
const int out_offset = (((output_batch_idx + output_channel_block_idx*batch)
* out_depth + output_depth_idx) * out_height + output_height_idx) * out_width + output_width_idx;
FLOAT4 value = vload4(inp_offset, input);
vstore4(value, out_offset, output);
}