wehub-resource-sync
149 lines
7.2 KiB
Common Lisp
149 lines
7.2 KiB
Common Lisp
#define GLOBAL_SIZE_3_DIMS \
|
|
__private const int global_size_dim0, __private const int global_size_dim1, __private const int global_size_dim2,
|
|
#ifdef MNN_SUPPORT_FP16
|
|
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
|
|
#endif
|
|
#define DEAL_NON_UNIFORM_DIM3(input1, input2, input3) \
|
|
if (input1 >= global_size_dim0 || input2 >= global_size_dim1 || input3 >= global_size_dim2) { \
|
|
return; \
|
|
}
|
|
|
|
__constant sampler_t SAMPLER = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
|
|
|
|
__kernel void deconv_2d(GLOBAL_SIZE_3_DIMS
|
|
#ifdef USE_BUFFER
|
|
__global FLOAT* input,
|
|
__global FLOAT* weights,
|
|
#ifdef BIAS
|
|
__global FLOAT* bias,
|
|
#endif
|
|
__global FLOAT* output, __private const int batch,
|
|
#else
|
|
__read_only image2d_t input,
|
|
__read_only image2d_t weights,
|
|
#ifdef BIAS
|
|
__read_only image2d_t bias,
|
|
#endif
|
|
__write_only image2d_t output,
|
|
#endif
|
|
__private const int2 input_shape,
|
|
__private const int2 output_shape,
|
|
__private const int2 stride_shape,
|
|
__private const int2 align_shape,
|
|
__private const int2 padding_shape,
|
|
__private const int2 kernel_shape,
|
|
__private const int kernel_size,
|
|
__private const int in_channel_blocks, __private const int out_channel_blocks) {
|
|
|
|
const int out_channel_blocks_idx = get_global_id(0);
|
|
const int out_w_idx = get_global_id(1);
|
|
const int out_batch_height_idx = get_global_id(2);
|
|
|
|
DEAL_NON_UNIFORM_DIM3(out_channel_blocks_idx, out_w_idx, out_batch_height_idx);
|
|
|
|
#ifdef BIAS
|
|
#ifdef USE_BUFFER
|
|
FLOAT4 out0 = vload4(out_channel_blocks_idx, bias);
|
|
#else
|
|
FLOAT4 out0 = RI_F(bias, SAMPLER, (int2)(out_channel_blocks_idx, 0));
|
|
#endif
|
|
#else
|
|
FLOAT4 out0 = (FLOAT4)0;
|
|
#endif
|
|
|
|
const int out_b_idx = out_batch_height_idx / output_shape.x;
|
|
const int out_h_idx = out_batch_height_idx % output_shape.x;
|
|
|
|
int kernel_start_x = max(0, (out_w_idx + align_shape.y) / stride_shape.y);
|
|
int kernel_start_y = max(0, (out_h_idx + align_shape.x) / stride_shape.x);
|
|
int deal_kernel_width = kernel_shape.y - mad24(kernel_start_x, stride_shape.y, padding_shape.y) + out_w_idx - 1;
|
|
int deal_kernel_height = kernel_shape.x - mad24(kernel_start_y, stride_shape.x, padding_shape.x) + out_h_idx - 1;
|
|
|
|
|
|
int kernel_x_0, kernel_x_1, kernel_x_2, kernel_x_3, kernel_y;
|
|
FLOAT4 in0;
|
|
FLOAT4 weights0, weights1, weights2, weights3;
|
|
for (int ic = 0; ic < in_channel_blocks; ic++) {
|
|
kernel_x_0 = ic << 2;
|
|
kernel_x_1 = kernel_x_0 + 1;
|
|
kernel_x_2 = kernel_x_0 + 2;
|
|
kernel_x_3 = kernel_x_0 + 3;
|
|
for (int k_y = deal_kernel_height, idx_h = kernel_start_y; k_y >= 0; k_y -= stride_shape.x, idx_h++) {
|
|
#ifdef USE_BUFFER
|
|
int in_width0 = kernel_start_x;
|
|
for (int k_x = deal_kernel_width; k_x >= 0; k_x -= stride_shape.y) {
|
|
kernel_y = mad24(k_y, kernel_shape.y, k_x);
|
|
kernel_y = mad24(out_channel_blocks_idx, kernel_size, kernel_y);
|
|
//weights NC4HW4 [1, 4*icC4, ocC4*kh*kw, 1] xic4
|
|
//index: [0, kernel_x_0, kernel_y, 0]
|
|
weights0 = vload4(kernel_x_0*(out_channel_blocks*kernel_shape.x*kernel_shape.y)+kernel_y, weights);
|
|
weights1 = vload4(kernel_x_1*(out_channel_blocks*kernel_shape.x*kernel_shape.y)+kernel_y, weights);
|
|
weights2 = vload4(kernel_x_2*(out_channel_blocks*kernel_shape.x*kernel_shape.y)+kernel_y, weights);
|
|
weights3 = vload4(kernel_x_3*(out_channel_blocks*kernel_shape.x*kernel_shape.y)+kernel_y, weights);
|
|
|
|
bool outBoundry = (idx_h < 0 || idx_h >= input_shape.x || kernel_start_x < 0 || in_width0 >= input_shape.y);
|
|
int inp_offset = (((out_b_idx + ic * batch) * input_shape.x + idx_h) * input_shape.y + in_width0) * 4;
|
|
in0 = outBoundry ? (FLOAT4)0 : vload4(0, input+inp_offset);
|
|
|
|
out0 = mad(in0.x, weights0, out0);
|
|
out0 = mad(in0.y, weights1, out0);
|
|
out0 = mad(in0.z, weights2, out0);
|
|
out0 = mad(in0.w, weights3, out0);
|
|
in_width0++;
|
|
}
|
|
#else
|
|
int in_idy = mad24(out_b_idx, input_shape.x, idx_h);
|
|
int in_hb_value = select(in_idy, -1, idx_h < 0 || idx_h >= input_shape.x);
|
|
int in_width0 = kernel_start_x;
|
|
for (int k_x = deal_kernel_width; k_x >= 0; k_x -= stride_shape.y) {
|
|
kernel_y = mad24(k_y, kernel_shape.y, k_x);
|
|
kernel_y = mad24(out_channel_blocks_idx, kernel_size, kernel_y);
|
|
weights0 = RI_F(weights, SAMPLER, (int2)(kernel_x_0, kernel_y));
|
|
weights1 = RI_F(weights, SAMPLER, (int2)(kernel_x_1, kernel_y));
|
|
weights2 = RI_F(weights, SAMPLER, (int2)(kernel_x_2, kernel_y));
|
|
weights3 = RI_F(weights, SAMPLER, (int2)(kernel_x_3, kernel_y));
|
|
|
|
int in_idx = mul24(ic, input_shape.y);
|
|
int in_width_value0 = in_width0; \
|
|
in_width_value0 = \
|
|
select(in_idx + in_width_value0, -1, (in_width_value0 < 0 || in_width_value0 >= input_shape.y)); \
|
|
in0 = RI_F(input, SAMPLER, (int2)(in_width_value0, in_hb_value));
|
|
|
|
out0 = mad(in0.x, weights0, out0);
|
|
out0 = mad(in0.y, weights1, out0);
|
|
out0 = mad(in0.z, weights2, out0);
|
|
out0 = mad(in0.w, weights3, out0);
|
|
in_width0++;
|
|
}
|
|
#endif
|
|
}
|
|
}
|
|
#ifdef RELU
|
|
out0 = fmax(out0, (FLOAT4)0);
|
|
#endif
|
|
|
|
#ifdef RELU6
|
|
out0 = clamp(out0, (FLOAT4)0, (FLOAT4)6);
|
|
#endif
|
|
|
|
#ifdef USE_BUFFER
|
|
const int out_offset = (((out_b_idx + out_channel_blocks_idx*batch)*output_shape.x + out_h_idx)*output_shape.y + out_w_idx)*4;
|
|
vstore4(out0, 0, output+out_offset);
|
|
#else
|
|
int out_image_width_idx = mad24(out_channel_blocks_idx, output_shape.y, out_w_idx);
|
|
WI_F(output, (int2)(out_image_width_idx, out_batch_height_idx), out0);
|
|
#endif
|
|
}
|
|
|
|
__kernel void iohw2oihw(__global const float* input_ptr, __global float* output_ptr, int plane_number, int input_channel, int output_channel) {
|
|
const int ic_index = get_global_id(0), oc_index = get_global_id(1);
|
|
if (ic_index >= input_channel || oc_index >= output_channel) {
|
|
return;
|
|
}
|
|
const int input_offset = (ic_index * output_channel + oc_index) * plane_number;
|
|
const int output_offset = (oc_index * input_channel + ic_index) * plane_number;
|
|
for (int i = 0; i < plane_number; ++i) {
|
|
output_ptr[output_offset + i] = input_ptr[input_offset + i];
|
|
}
|
|
}
|