wehub-resource-sync
367 lines
16 KiB
Common Lisp
367 lines
16 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; \
|
|
}
|
|
|
|
#define MNN_DATA_FORMAT_NCHW 0
|
|
#define MNN_DATA_FORMAT_NHWC 1
|
|
#define MNN_DATA_FORMAT_NC4HW4 2
|
|
#define MNN_DATA_FORMAT_C4NHW4 3
|
|
|
|
#define __CAT(x, y) x##y
|
|
#define CAT(x, y) __CAT(x, y)
|
|
#define OUTPUT_TYPE2 CAT(OUTPUT_TYPE, 2)
|
|
#define OUTPUT_TYPE3 CAT(OUTPUT_TYPE, 3)
|
|
__constant sampler_t SAMPLER = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
|
|
|
|
#ifdef SHARED_TO_CL
|
|
__kernel void gl_to_cl(GLOBAL_SIZE_3_DIMS
|
|
__global uchar *input_ptr,
|
|
#ifdef USE_IMAGE
|
|
__write_only image2d_t output_ptr,
|
|
#else
|
|
__global OUTPUT_TYPE *output_ptr,
|
|
#endif
|
|
__private const int4 shape, // N C H W
|
|
__private const OffsetStride
|
|
) {
|
|
|
|
int wblock = get_global_id(0);
|
|
int cblock = get_global_id(1);
|
|
int nh = get_global_id(2);
|
|
|
|
DEAL_NON_UNIFORM_DIM3(wblock, cblock, nh);
|
|
const int w = wblock << 2;
|
|
const int h = nh % shape.z;
|
|
const int c = cblock << 2;
|
|
const int n = nh / shape.z;
|
|
|
|
int Inidx = c * OffsetStride + w; // c/4*w
|
|
int idx = c * shape.w + w; // c/4*w
|
|
int idy = nh; // n*h
|
|
const int offset = idy * OffsetStride * 4;
|
|
OUTPUT_TYPE4 in0 = CONVERT_OUTPUT4(vload4(Inidx, input_ptr + offset));
|
|
OUTPUT_TYPE4 in1 = w + 1 < OffsetStride ? CONVERT_OUTPUT4(vload4(Inidx + 1, input_ptr + offset)) : (OUTPUT_TYPE4)(0);
|
|
OUTPUT_TYPE4 in2 = w + 2 < OffsetStride ? CONVERT_OUTPUT4(vload4(Inidx + 2, input_ptr + offset)) : (OUTPUT_TYPE4)(0);
|
|
OUTPUT_TYPE4 in3 = w + 3 < OffsetStride ? CONVERT_OUTPUT4(vload4(Inidx + 3, input_ptr + offset)) : (OUTPUT_TYPE4)(0);
|
|
|
|
#ifdef USE_IMAGE
|
|
WI_DATA(output_ptr, (int2)(idx, idy), in0);
|
|
if(w + 1 >= shape.w) return;
|
|
WI_DATA(output_ptr, (int2)(idx+1, idy), in1);
|
|
if(w + 2 >= shape.w) return;
|
|
WI_DATA(output_ptr, (int2)(idx+2, idy), in2);
|
|
if(w + 3 >= shape.w) return;
|
|
WI_DATA(output_ptr, (int2)(idx+3, idy), in3);
|
|
#else
|
|
#if OUTPUT_FORMAT == MNN_DATA_FORMAT_NCHW
|
|
int output_offset = ((n * shape.y + c) * shape.z + h) * shape.w + w;
|
|
int stride = shape.z * shape.w;
|
|
int remain = shape.w - w;
|
|
if(remain >= 4){
|
|
vstore4((OUTPUT_TYPE4)(in0.x, in1.x, in2.x, in3.x), 0, output_ptr + output_offset);
|
|
if(c + 1 >= shape.y) return;
|
|
vstore4((OUTPUT_TYPE4)(in0.y, in1.y, in2.y, in3.y), 0, output_ptr + output_offset + stride);
|
|
if(c + 2 >= shape.y) return;
|
|
vstore4((OUTPUT_TYPE4)(in0.z, in1.z, in2.z, in3.z), 0, output_ptr + output_offset + stride + stride);
|
|
if(c + 3 >= shape.y) return;
|
|
vstore4((OUTPUT_TYPE4)(in0.w, in1.w, in2.w, in3.w), 0, output_ptr + output_offset + stride + stride + stride);
|
|
} else if(remain == 3){
|
|
vstore3((OUTPUT_TYPE3)(in0.x, in1.x, in2.x), 0, output_ptr + output_offset);
|
|
if(c + 1 >= shape.y) return;
|
|
vstore3((OUTPUT_TYPE3)(in0.y, in1.y, in2.y), 0, output_ptr + output_offset + stride);
|
|
if(c + 2 >= shape.y) return;
|
|
vstore3((OUTPUT_TYPE3)(in0.z, in1.z, in2.z), 0, output_ptr + output_offset + stride + stride);
|
|
if(c + 3 >= shape.y) return;
|
|
vstore3((OUTPUT_TYPE3)(in0.w, in1.w, in2.w), 0, output_ptr + output_offset + stride + stride + stride);
|
|
} else if(remain == 2){
|
|
vstore2((OUTPUT_TYPE2)(in0.x, in1.x), 0, output_ptr + output_offset);
|
|
if(c + 1 >= shape.y) return;
|
|
vstore2((OUTPUT_TYPE2)(in0.y, in1.y), 0, output_ptr + output_offset + stride);
|
|
if(c + 2 >= shape.y) return;
|
|
vstore2((OUTPUT_TYPE2)(in0.z, in1.z), 0, output_ptr + output_offset + stride + stride);
|
|
if(c + 3 >= shape.y) return;
|
|
vstore2((OUTPUT_TYPE2)(in0.w, in1.w), 0, output_ptr + output_offset + stride + stride + stride);
|
|
}else if(remain == 1){
|
|
output_ptr[output_offset] = in0.x;
|
|
if(c + 1 >= shape.y) return;
|
|
output_ptr[output_offset + stride] = in0.y;
|
|
if(c + 2 >= shape.y) return;
|
|
output_ptr[output_offset + stride + stride] = in0.z;
|
|
if(c + 3 >= shape.y) return;
|
|
output_ptr[output_offset + stride + stride + stride] = in0.w;
|
|
}
|
|
#elif OUTPUT_FORMAT == MNN_DATA_FORMAT_NHWC
|
|
int output_offset = ((n * shape.z + h) * shape.w + w) * shape.y + c;
|
|
int remain = shape.y - c;
|
|
if(remain >= 4){
|
|
vstore4(CONVERT_OUTPUT4(in0), 0, output_ptr + output_offset);
|
|
if(w + 1 >= shape.w) return;
|
|
vstore4(CONVERT_OUTPUT4(in1), 0, output_ptr + output_offset + shape.y);
|
|
if(w + 2 >= shape.w) return;
|
|
vstore4(CONVERT_OUTPUT4(in2), 0, output_ptr + output_offset + shape.y + shape.y);
|
|
if(w + 3 >= shape.w) return;
|
|
vstore4(CONVERT_OUTPUT4(in3), 0, output_ptr + output_offset + shape.y + shape.y + shape.y);
|
|
} else if(remain == 3){
|
|
vstore3((OUTPUT_TYPE3)(in0.x, in0.y, in0.z), 0, output_ptr + output_offset);
|
|
if(w + 1 >= shape.w) return;
|
|
vstore3((OUTPUT_TYPE3)(in1.x, in1.y, in1.z), 0, output_ptr + output_offset + shape.y);
|
|
if(w + 2 >= shape.w) return;
|
|
vstore3((OUTPUT_TYPE3)(in2.x, in2.y, in2.z), 0, output_ptr + output_offset + shape.y + shape.y);
|
|
if(w + 3 >= shape.w) return;
|
|
vstore3((OUTPUT_TYPE3)(in3.x, in3.y, in3.z), 0, output_ptr + output_offset + shape.y + shape.y + shape.y);
|
|
} else if(remain == 2){
|
|
vstore2((OUTPUT_TYPE2)(in0.x, in0.y), 0, output_ptr + output_offset);
|
|
if(w + 1 >= shape.w) return;
|
|
vstore2((OUTPUT_TYPE2)(in1.x, in1.y), 0, output_ptr + output_offset + shape.y);
|
|
if(w + 2 >= shape.w) return;
|
|
vstore2((OUTPUT_TYPE2)(in2.x, in2.y), 0, output_ptr + output_offset + shape.y + shape.y);
|
|
if(w + 3 >= shape.w) return;
|
|
vstore2((OUTPUT_TYPE2)(in3.x, in3.y), 0, output_ptr + output_offset + shape.y + shape.y + shape.y);
|
|
}else if(remain == 1){
|
|
output_ptr[output_offset] = in0.x;
|
|
if(w + 1 >= shape.w) return;
|
|
output_ptr[output_offset + shape.y] = in1.x;
|
|
if(w + 2 >= shape.w) return;
|
|
output_ptr[output_offset + shape.y + shape.y] = in1.x;
|
|
if(w + 3 >= shape.w) return;
|
|
output_ptr[output_offset + shape.y + shape.y + shape.y] = in1.x;
|
|
}
|
|
#elif OUTPUT_FORMAT == MNN_DATA_FORMAT_NC4HW4
|
|
int output_offset = (((cblock * shape.x + n) * shape.z + h) * shape.w + w) * 4;
|
|
vstore4(in0, 0, output_ptr + output_offset);
|
|
if(w + 1 >= shape.w) return;
|
|
vstore4(in1, 0, output_ptr + output_offset + 4);
|
|
if(w + 2 >= shape.w) return;
|
|
vstore4(in2, 0, output_ptr + output_offset + 8);
|
|
if(w + 3 >= shape.w) return;
|
|
vstore4(in3, 0, output_ptr + output_offset + 12);
|
|
#else
|
|
//not support
|
|
#endif
|
|
#endif
|
|
}
|
|
|
|
// only copy y channel
|
|
__kernel void yuv_to_cl(GLOBAL_SIZE_3_DIMS
|
|
__global uchar *input_ptr,
|
|
#ifdef USE_IMAGE
|
|
__write_only image2d_t output_ptr,
|
|
#else
|
|
__global OUTPUT_TYPE *output_ptr,
|
|
#endif
|
|
__private const int4 shape, // N C H W
|
|
__private const OffsetStride
|
|
) {
|
|
|
|
int wblock = get_global_id(0);
|
|
int cblock = get_global_id(1);
|
|
int nh = get_global_id(2);
|
|
|
|
DEAL_NON_UNIFORM_DIM3(wblock, cblock, nh);
|
|
const int w = wblock << 2;
|
|
const int h = nh % shape.z;
|
|
const int c = cblock << 2;
|
|
const int n = nh / shape.z;
|
|
|
|
int idx = c * shape.w + w; // c/4*w
|
|
int idy = nh; // n*h
|
|
const int offset = idy * OffsetStride + w;
|
|
OUTPUT_TYPE4 in0 = CONVERT_OUTPUT4(vload4(0, input_ptr + offset));
|
|
|
|
#ifdef USE_IMAGE
|
|
WI_DATA(output_ptr, (int2)(idx, idy), (OUTPUT_TYPE4)(in0.x, 0, 0, 0));
|
|
if(w + 1 >= shape.w) return;
|
|
WI_DATA(output_ptr, (int2)(idx+1, idy), (OUTPUT_TYPE4)(in0.y, 0, 0, 0));
|
|
if(w + 2 >= shape.w) return;
|
|
WI_DATA(output_ptr, (int2)(idx+2, idy), (OUTPUT_TYPE4)(in0.z, 0, 0, 0));
|
|
if(w + 3 >= shape.w) return;
|
|
WI_DATA(output_ptr, (int2)(idx+3, idy), (OUTPUT_TYPE4)(in0.w, 0, 0, 0));
|
|
#else
|
|
#if OUTPUT_FORMAT == MNN_DATA_FORMAT_NCHW
|
|
int output_offset = ((n * shape.y + c) * shape.z + h) * shape.w + w;
|
|
int remain = shape.w - w;
|
|
if(remain >= 4){
|
|
vstore4(in0, 0, output_ptr + output_offset);
|
|
} else if(remain == 3){
|
|
vstore3((OUTPUT_TYPE3)(in0.x, in0.y, in0.z), 0, output_ptr + output_offset);
|
|
} else if(remain == 2){
|
|
vstore2((OUTPUT_TYPE2)(in0.x, in0.y), 0, output_ptr + output_offset);
|
|
}else if(remain == 1){
|
|
output_ptr[output_offset] = in0.x;
|
|
}
|
|
#elif OUTPUT_FORMAT == MNN_DATA_FORMAT_NHWC
|
|
int output_offset = ((n * shape.z + h) * shape.w + w) * shape.y + c;
|
|
output_ptr[output_offset] = in0.x;
|
|
if(w + 1 >= shape.w) return;
|
|
output_ptr[output_offset + shape.y] = in0.y;
|
|
if(w + 2 >= shape.w) return;
|
|
output_ptr[output_offset + shape.y + shape.y] = in0.z;
|
|
if(w + 3 >= shape.w) return;
|
|
output_ptr[output_offset + shape.y + shape.y + shape.y] = in0.w;
|
|
#elif OUTPUT_FORMAT == MNN_DATA_FORMAT_NC4HW4
|
|
int output_offset = (((cblock * shape.x + n) * shape.z + h) * shape.w + w) * 4;
|
|
vstore4((OUTPUT_TYPE4)(in0.x, 0, 0, 0), 0, output_ptr + output_offset);
|
|
if(w + 1 >= shape.w) return;
|
|
vstore4((OUTPUT_TYPE4)(in0.y, 0, 0, 0), 0, output_ptr + output_offset + 4);
|
|
if(w + 2 >= shape.w) return;
|
|
vstore4((OUTPUT_TYPE4)(in0.z, 0, 0, 0), 0, output_ptr + output_offset + 8);
|
|
if(w + 3 >= shape.w) return;
|
|
vstore4((OUTPUT_TYPE4)(in0.w, 0, 0, 0), 0, output_ptr + output_offset + 12);
|
|
#else
|
|
//not support
|
|
#endif
|
|
#endif
|
|
}
|
|
|
|
#endif // SHARED_TO_CL
|
|
|
|
#ifdef CL_TO_SHARED
|
|
__kernel void cl_to_gl(GLOBAL_SIZE_3_DIMS
|
|
#ifdef USE_IMAGE
|
|
__read_only image2d_t input_ptr,
|
|
#else
|
|
__global INPUT_TYPE *input_ptr,
|
|
#endif
|
|
__global uchar *output_ptr,
|
|
__private const int4 shape, // N C H W
|
|
__private const OffsetStride
|
|
) {
|
|
|
|
int wblock = get_global_id(0);
|
|
int cblock = get_global_id(1);
|
|
int nh = get_global_id(2);
|
|
|
|
DEAL_NON_UNIFORM_DIM3(wblock, cblock, nh);
|
|
const int w = wblock << 2;
|
|
const int h = nh % shape.z;
|
|
const int c = cblock << 2;
|
|
const int n = nh / shape.z;
|
|
|
|
int idx = c * shape.w + w; // c/4*w
|
|
int idy = nh; // n*h
|
|
INPUT_TYPE4 in0, in1, in2, in3;
|
|
#ifdef USE_IMAGE
|
|
in0 = RI_DATA(input_ptr, SAMPLER, (int2)(idx, idy));
|
|
in1 = RI_DATA(input_ptr, SAMPLER, (int2)(idx+1, idy));
|
|
in2 = RI_DATA(input_ptr, SAMPLER, (int2)(idx+2, idy));
|
|
in3 = RI_DATA(input_ptr, SAMPLER, (int2)(idx+3, idy));
|
|
#else
|
|
#if INPUT_FORMAT == MNN_DATA_FORMAT_NCHW
|
|
int input_offset = ((n * shape.y + c) * shape.z + h) * shape.w + w;
|
|
int stride = shape.z * shape.w;
|
|
INPUT_TYPE4 tmp0, tmp1, tmp2, tmp3;
|
|
tmp0 = vload4(0, input_ptr + input_offset);
|
|
tmp1 = vload4(0, input_ptr + input_offset + stride);
|
|
tmp2 = vload4(0, input_ptr + input_offset + stride + stride);
|
|
tmp3 = vload4(0, input_ptr + input_offset + stride + stride + stride);
|
|
in0 = (INPUT_TYPE4)(tmp0.x, tmp1.x, tmp2.x, tmp3.x);
|
|
in1 = (INPUT_TYPE4)(tmp0.y, tmp1.y, tmp2.y, tmp3.y);
|
|
in2 = (INPUT_TYPE4)(tmp0.z, tmp1.z, tmp2.z, tmp3.z);
|
|
in3 = (INPUT_TYPE4)(tmp0.w, tmp1.w, tmp2.w, tmp3.w);
|
|
#elif INPUT_FORMAT == MNN_DATA_FORMAT_NHWC
|
|
int input_offset = ((n * shape.z + h) * shape.w + w) * shape.y + c;
|
|
in0 = vload4(0, input_ptr + input_offset);
|
|
in1 = vload4(0, input_ptr + input_offset + shape.y);
|
|
in2 = vload4(0, input_ptr + input_offset + shape.y + shape.y);
|
|
in3 = vload4(0, input_ptr + input_offset + shape.y + shape.y + shape.y);
|
|
#elif INPUT_FORMAT == MNN_DATA_FORMAT_NC4HW4
|
|
int input_offset = (((cblock * shape.x + n) * shape.z + h) * shape.w + w) * 4;
|
|
in0 = vload4(0, input_ptr + input_offset);
|
|
in1 = vload4(0, input_ptr + input_offset + 4);
|
|
in2 = vload4(0, input_ptr + input_offset + 8);
|
|
in3 = vload4(0, input_ptr + input_offset + 12);
|
|
#else
|
|
// not support
|
|
#endif
|
|
#endif
|
|
int OutIdx = c * OffsetStride + w; // c/4*w
|
|
const int offset = idy * OffsetStride * 4;
|
|
vstore4(convert_uchar4(in0), OutIdx, output_ptr + offset);
|
|
if(w + 1 >= OffsetStride) return;
|
|
vstore4(convert_uchar4(in1), OutIdx+1, output_ptr + offset);
|
|
if(w + 2 >= OffsetStride) return;
|
|
vstore4(convert_uchar4(in2), OutIdx+2, output_ptr + offset);
|
|
if(w + 3 >= OffsetStride) return;
|
|
vstore4(convert_uchar4(in3), OutIdx+3, output_ptr + offset);
|
|
}
|
|
|
|
// only copy y channel
|
|
__kernel void cl_to_yuv(GLOBAL_SIZE_3_DIMS
|
|
#ifdef USE_IMAGE
|
|
__read_only image2d_t input_ptr,
|
|
#else
|
|
__global INPUT_TYPE *input_ptr,
|
|
#endif
|
|
__global uchar *output_ptr,
|
|
__private const int4 shape, // N C H W
|
|
__private const OffsetStride
|
|
) {
|
|
|
|
int wblock = get_global_id(0);
|
|
int cblock = get_global_id(1);
|
|
int nh = get_global_id(2);
|
|
|
|
DEAL_NON_UNIFORM_DIM3(wblock, cblock, nh);
|
|
const int w = wblock << 2;
|
|
const int h = nh % shape.z;
|
|
const int c = cblock << 2;
|
|
const int n = nh / shape.z;
|
|
|
|
int idx = c * shape.w + w; // c/4*w
|
|
int idy = nh; // n*h
|
|
INPUT_TYPE4 in0, in1, in2, in3;
|
|
#ifdef USE_IMAGE
|
|
in0 = RI_DATA(input_ptr, SAMPLER, (int2)(idx, idy));
|
|
in1 = RI_DATA(input_ptr, SAMPLER, (int2)(idx+1, idy));
|
|
in2 = RI_DATA(input_ptr, SAMPLER, (int2)(idx+2, idy));
|
|
in3 = RI_DATA(input_ptr, SAMPLER, (int2)(idx+3, idy));
|
|
#else
|
|
#if INPUT_FORMAT == MNN_DATA_FORMAT_NCHW
|
|
int input_offset = ((n * shape.y + c) * shape.z + h) * shape.w + w;
|
|
int stride = shape.z * shape.w;
|
|
INPUT_TYPE4 tmp0, tmp1, tmp2, tmp3;
|
|
tmp0 = vload4(0, input_ptr + input_offset);
|
|
tmp1 = vload4(0, input_ptr + input_offset + stride);
|
|
tmp2 = vload4(0, input_ptr + input_offset + stride + stride);
|
|
tmp3 = vload4(0, input_ptr + input_offset + stride + stride + stride);
|
|
in0 = (INPUT_TYPE4)(tmp0.x, tmp1.x, tmp2.x, tmp3.x);
|
|
in1 = (INPUT_TYPE4)(tmp0.y, tmp1.y, tmp2.y, tmp3.y);
|
|
in2 = (INPUT_TYPE4)(tmp0.z, tmp1.z, tmp2.z, tmp3.z);
|
|
in3 = (INPUT_TYPE4)(tmp0.w, tmp1.w, tmp2.w, tmp3.w);
|
|
#elif INPUT_FORMAT == MNN_DATA_FORMAT_NHWC
|
|
int input_offset = ((n * shape.z + h) * shape.w + w) * shape.y + c;
|
|
in0 = vload4(0, input_ptr + input_offset);
|
|
in1 = vload4(0, input_ptr + input_offset + shape.y);
|
|
in2 = vload4(0, input_ptr + input_offset + shape.y + shape.y);
|
|
in3 = vload4(0, input_ptr + input_offset + shape.y + shape.y + shape.y);
|
|
#elif INPUT_FORMAT == MNN_DATA_FORMAT_NC4HW4
|
|
int input_offset = (((cblock * shape.x + n) * shape.z + h) * shape.w + w) * 4;
|
|
in0 = vload4(0, input_ptr + input_offset);
|
|
in1 = vload4(0, input_ptr + input_offset + 4);
|
|
in2 = vload4(0, input_ptr + input_offset + 8);
|
|
in3 = vload4(0, input_ptr + input_offset + 12);
|
|
#else
|
|
// not support
|
|
#endif
|
|
#endif
|
|
int remain = OffsetStride - w;
|
|
uchar4 out = convert_uchar4((INPUT_TYPE4)(in0.x, in1.x, in2.x, in3.x));
|
|
const int offset = idy * OffsetStride + w;
|
|
if(remain >= 4){
|
|
vstore4(out, 0, output_ptr + offset);
|
|
}else if(remain == 3){
|
|
vstore3(out.xyz, 0, output_ptr + offset);
|
|
}else if(remain == 2){
|
|
vstore2(out.xy, 0, output_ptr + offset);
|
|
}else if(remain == 1){
|
|
output_ptr[offset] = out.x;
|
|
}
|
|
}
|
|
#endif // CL_TO_SHARED
|