wehub-resource-sync
231 lines
10 KiB
Common Lisp
231 lines
10 KiB
Common Lisp
#ifdef MNN_SUPPORT_FP16
|
|
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
|
|
#endif
|
|
#define READ_INPUT_IMAGE(i, base) \
|
|
int inOffset##i = inWidthOffset##i + base; \
|
|
inOffset##i = \
|
|
select(inCurIdx + inOffset##i, -1, (inOffset##i < 0 || inOffset##i >= inputShape.y)); \
|
|
inValue##i = RI_F(input, SAMPLER, (int2)(inOffset##i, inHeightIdx));
|
|
|
|
#define CALCULATE_OUTPUT(i) \
|
|
outValue##i = mad(inValue##i.x, weights0, outValue##i); \
|
|
outValue##i = mad(inValue##i.y, weights1, outValue##i); \
|
|
outValue##i = mad(inValue##i.z, weights2, outValue##i); \
|
|
outValue##i = mad(inValue##i.w, weights3, outValue##i);
|
|
|
|
#define GLOBAL_SIZE_2_DIMS __private const int global_size_dim0, __private const int global_size_dim1,
|
|
|
|
__constant sampler_t SAMPLER = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
|
|
|
|
#define DEAL_NON_UNIFORM_DIM2(input1, input2) \
|
|
if (input1 >= global_size_dim0 || input2 >= global_size_dim1) { \
|
|
return; \
|
|
}
|
|
|
|
__kernel
|
|
#ifdef SET_ATTRIBUTE
|
|
__attribute__((work_group_size_hint(16, 16, 1)))
|
|
#endif
|
|
void depthwise_conv2d_s1(GLOBAL_SIZE_2_DIMS __read_only image2d_t input, __read_only image2d_t filter,
|
|
#ifndef NO_BIAS
|
|
__read_only image2d_t bias,
|
|
#endif
|
|
__write_only image2d_t output,
|
|
__private const int2 inputShape,
|
|
__private const int inChannelBlocks,
|
|
__private const int2 outputShape,
|
|
__private const int2 filterShape,
|
|
__private const int2 paddingShape) {
|
|
|
|
const int outChannelWidthIdx = get_global_id(0);
|
|
const int outHeightBlockIdx = get_global_id(1);
|
|
DEAL_NON_UNIFORM_DIM2(outChannelWidthIdx, outHeightBlockIdx);
|
|
int ow4 = (outputShape.y + 3) / 4;
|
|
const int outChannelBlockIdx = outChannelWidthIdx / ow4;
|
|
const int outWidthBlockidx = outChannelWidthIdx % ow4;
|
|
|
|
const int inChannelBlockIdx = outChannelBlockIdx;
|
|
|
|
#ifndef NO_BIAS
|
|
FLOAT4 outValue0 = RI_F(bias, SAMPLER, (int2)(outChannelBlockIdx, 0));
|
|
#else
|
|
FLOAT4 outValue0 = (FLOAT4)(0.0f);
|
|
#endif
|
|
FLOAT4 outValue1 = outValue0;
|
|
FLOAT4 outValue2 = outValue0;
|
|
FLOAT4 outValue3 = outValue0;
|
|
|
|
const int outWidthBlockidx4 = outWidthBlockidx << 2;
|
|
const int inWidthOffset0 = outWidthBlockidx4 - paddingShape.y;
|
|
const int inWidthOffset1 = inWidthOffset0 + 1;
|
|
const int inWidthOffset2 = inWidthOffset0 + 2;
|
|
const int inWidthOffset3 = inWidthOffset0 + 3;
|
|
|
|
int heightIdx = outHeightBlockIdx % outputShape.x - paddingShape.x;
|
|
const int outBatchIdx = mul24((outHeightBlockIdx / outputShape.x), inputShape.x);
|
|
const int inCurIdx = mul24(inChannelBlockIdx, inputShape.y);
|
|
|
|
const int inWidthIdx0 = select(inCurIdx + inWidthOffset0, -1, (inWidthOffset0 < 0 || inWidthOffset0 >= inputShape.y));
|
|
const int inWidthIdx1 = select(inCurIdx + inWidthOffset1, -1, (inWidthOffset1 < 0 || inWidthOffset1 >= inputShape.y));
|
|
const int inWidthIdx2 = select(inCurIdx + inWidthOffset2, -1, (inWidthOffset2 < 0 || inWidthOffset2 >= inputShape.y));
|
|
|
|
FLOAT4 inValue0, inValue1, inValue2, inValue3;
|
|
for (int kh = 0; kh < filterShape.x; kh++) {
|
|
int inHeightIdx = select(heightIdx + outBatchIdx, -1, (heightIdx < 0 || heightIdx >= inputShape.x));
|
|
heightIdx++;
|
|
inValue1 = RI_F(input, SAMPLER, (int2)(inWidthIdx0, inHeightIdx));
|
|
inValue2 = RI_F(input, SAMPLER, (int2)(inWidthIdx1, inHeightIdx));
|
|
inValue3 = RI_F(input, SAMPLER, (int2)(inWidthIdx2, inHeightIdx));
|
|
for (int kw = 0; kw < filterShape.y; kw++) {
|
|
|
|
int filterIdx = mad24(kh, filterShape.y, kw);
|
|
inValue0 = inValue1;
|
|
inValue1 = inValue2;
|
|
inValue2 = inValue3;
|
|
|
|
int inWidthIdx = inWidthOffset3 + kw;
|
|
inWidthIdx = select(inCurIdx + inWidthIdx, -1, (inWidthIdx < 0 || inWidthIdx >= inputShape.y));
|
|
inValue3 = RI_F(input, SAMPLER, (int2)(inWidthIdx, inHeightIdx));
|
|
|
|
FLOAT4 weights = RI_F(filter, SAMPLER, (int2)(filterIdx, inChannelBlockIdx));
|
|
|
|
outValue0 = mad(inValue0, weights, outValue0);
|
|
outValue1 = mad(inValue1, weights, outValue1);
|
|
outValue2 = mad(inValue2, weights, outValue2);
|
|
outValue3 = mad(inValue3, weights, outValue3);
|
|
}
|
|
}
|
|
|
|
#ifdef RELU
|
|
outValue0 = fmax(outValue0, (FLOAT4)0);
|
|
outValue1 = fmax(outValue1, (FLOAT4)0);
|
|
outValue2 = fmax(outValue2, (FLOAT4)0);
|
|
outValue3 = fmax(outValue3, (FLOAT4)0);
|
|
#endif
|
|
|
|
#ifdef RELU6
|
|
outValue0 = clamp(outValue0, (FLOAT4)0, (FLOAT4)6);
|
|
outValue1 = clamp(outValue1, (FLOAT4)0, (FLOAT4)6);
|
|
outValue2 = clamp(outValue2, (FLOAT4)0, (FLOAT4)6);
|
|
outValue3 = clamp(outValue3, (FLOAT4)0, (FLOAT4)6);
|
|
#endif
|
|
|
|
const int remain = outputShape.y - outWidthBlockidx4;
|
|
int outWidthIdx = mul24(outChannelBlockIdx, outputShape.y) + outWidthBlockidx4;
|
|
if (remain >= 4) {
|
|
WI_F(output, (int2)(outWidthIdx, outHeightBlockIdx), outValue0);
|
|
WI_F(output, (int2)(outWidthIdx + 1, outHeightBlockIdx), outValue1);
|
|
WI_F(output, (int2)(outWidthIdx + 2, outHeightBlockIdx), outValue2);
|
|
WI_F(output, (int2)(outWidthIdx + 3, outHeightBlockIdx), outValue3);
|
|
} else if (remain == 3) {
|
|
WI_F(output, (int2)(outWidthIdx, outHeightBlockIdx), outValue0);
|
|
WI_F(output, (int2)(outWidthIdx + 1, outHeightBlockIdx), outValue1);
|
|
WI_F(output, (int2)(outWidthIdx + 2, outHeightBlockIdx), outValue2);
|
|
} else if (remain == 2) {
|
|
WI_F(output, (int2)(outWidthIdx, outHeightBlockIdx), outValue0);
|
|
WI_F(output, (int2)(outWidthIdx + 1, outHeightBlockIdx), outValue1);
|
|
} else if (remain == 1) {
|
|
WI_F(output, (int2)(outWidthIdx, outHeightBlockIdx), outValue0);
|
|
}
|
|
}
|
|
|
|
__kernel
|
|
#ifdef SET_ATTRIBUTE
|
|
__attribute__((work_group_size_hint(16, 16, 1)))
|
|
#endif
|
|
void depthwise_conv2d(GLOBAL_SIZE_2_DIMS __read_only image2d_t input, __read_only image2d_t filter,
|
|
#ifndef NO_BIAS
|
|
__read_only image2d_t bias,
|
|
#endif
|
|
__write_only image2d_t output,
|
|
__private const int2 inputShape,
|
|
__private const int inChannelBlocks, __private const int2 outputShape,
|
|
__private const int2 filterShape,
|
|
__private const int2 paddingShape,
|
|
__private const int2 dilationShape,
|
|
__private const int2 strideShape) {
|
|
|
|
const int outChannelWidthIdx = get_global_id(0);
|
|
const int outHeightIdx = get_global_id(1);
|
|
DEAL_NON_UNIFORM_DIM2(outChannelWidthIdx, outHeightIdx);
|
|
|
|
int ow4 = (outputShape.y + 3) / 4;
|
|
const int outChannelBlockIdx = outChannelWidthIdx / ow4;
|
|
const int outWidthBlockidx = outChannelWidthIdx % ow4;
|
|
|
|
const int inChannelBlockIdx = outChannelBlockIdx;
|
|
|
|
#ifndef NO_BIAS
|
|
FLOAT4 outValue0 = RI_F(bias, SAMPLER, (int2)(outChannelBlockIdx, 0));
|
|
#else
|
|
FLOAT4 outValue0 = (FLOAT4)(0.0f);
|
|
#endif
|
|
FLOAT4 outValue1 = outValue0;
|
|
FLOAT4 outValue2 = outValue0;
|
|
FLOAT4 outValue3 = outValue0;
|
|
|
|
const int inWidthOffset0 = mad24(outWidthBlockidx, strideShape.y << 2, -paddingShape.y);
|
|
const int inWidthOffset1 = inWidthOffset0 + strideShape.y;
|
|
const int inWidthOffset2 = inWidthOffset1 + strideShape.y;
|
|
const int inWidthOffset3 = inWidthOffset2 + strideShape.y;
|
|
int heightIdx = mad24(outHeightIdx % outputShape.x, strideShape.x, -paddingShape.x);
|
|
|
|
const int outBatchIdx = mul24((outHeightIdx / outputShape.x), inputShape.x);
|
|
|
|
const int inCurIdx = mul24(inChannelBlockIdx, inputShape.y);
|
|
for (int kh = 0; kh < filterShape.x; kh++) {
|
|
int inHeightIdx = select(heightIdx + outBatchIdx, -1, (heightIdx < 0 || heightIdx >= inputShape.x));
|
|
heightIdx += dilationShape.x;
|
|
for (int kw = 0; kw < filterShape.y; kw++) {
|
|
int filterIdx = mad24(kh, filterShape.y, kw);
|
|
FLOAT4 inValue0, inValue1, inValue2, inValue3;
|
|
int inWidthIdx = mul24(kw, dilationShape.y);
|
|
|
|
READ_INPUT_IMAGE(0, inWidthIdx);
|
|
READ_INPUT_IMAGE(1, inWidthIdx);
|
|
READ_INPUT_IMAGE(2, inWidthIdx);
|
|
READ_INPUT_IMAGE(3, inWidthIdx);
|
|
|
|
FLOAT4 weights = RI_F(filter, SAMPLER, (int2)(filterIdx, inChannelBlockIdx));
|
|
|
|
outValue0 = mad(inValue0, weights, outValue0);
|
|
outValue1 = mad(inValue1, weights, outValue1);
|
|
outValue2 = mad(inValue2, weights, outValue2);
|
|
outValue3 = mad(inValue3, weights, outValue3);
|
|
}
|
|
}
|
|
|
|
#ifdef RELU
|
|
outValue0 = fmax(outValue0, (FLOAT4)0);
|
|
outValue1 = fmax(outValue1, (FLOAT4)0);
|
|
outValue2 = fmax(outValue2, (FLOAT4)0);
|
|
outValue3 = fmax(outValue3, (FLOAT4)0);
|
|
#endif
|
|
|
|
#ifdef RELU6
|
|
outValue0 = clamp(outValue0, (FLOAT4)0, (FLOAT4)6);
|
|
outValue1 = clamp(outValue1, (FLOAT4)0, (FLOAT4)6);
|
|
outValue2 = clamp(outValue2, (FLOAT4)0, (FLOAT4)6);
|
|
outValue3 = clamp(outValue3, (FLOAT4)0, (FLOAT4)6);
|
|
#endif
|
|
|
|
const int outWidthBlockidx4 = outWidthBlockidx << 2;
|
|
const int remain = outputShape.y - outWidthBlockidx4;
|
|
int outWidthIdx = mul24(outChannelBlockIdx, outputShape.y) + outWidthBlockidx4;
|
|
if (remain >= 4) {
|
|
WI_F(output, (int2)(outWidthIdx, outHeightIdx), outValue0);
|
|
WI_F(output, (int2)(outWidthIdx + 1, outHeightIdx), outValue1);
|
|
WI_F(output, (int2)(outWidthIdx + 2, outHeightIdx), outValue2);
|
|
WI_F(output, (int2)(outWidthIdx + 3, outHeightIdx), outValue3);
|
|
} else if (remain == 3) {
|
|
WI_F(output, (int2)(outWidthIdx, outHeightIdx), outValue0);
|
|
WI_F(output, (int2)(outWidthIdx + 1, outHeightIdx), outValue1);
|
|
WI_F(output, (int2)(outWidthIdx + 2, outHeightIdx), outValue2);
|
|
} else if (remain == 2) {
|
|
WI_F(output, (int2)(outWidthIdx, outHeightIdx), outValue0);
|
|
WI_F(output, (int2)(outWidthIdx + 1, outHeightIdx), outValue1);
|
|
} else if (remain == 1) {
|
|
WI_F(output, (int2)(outWidthIdx, outHeightIdx), outValue0);
|
|
}
|
|
}
|