wehub-resource-sync
184 lines
9.9 KiB
Common Lisp
184 lines
9.9 KiB
Common Lisp
#ifdef MNN_SUPPORT_FP16
|
|
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
|
|
#endif
|
|
__constant sampler_t SAMPLER = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
|
|
__kernel void winogradTransformSource(__read_only image2d_t uInput, // 0
|
|
__write_only image2d_t uOutput, __private const int unitWidth,
|
|
__private const int unitHeight, // 3
|
|
__private const int padX, __private const int padY,
|
|
__private const int srcWidth, // 6
|
|
__private const int srcHeight, __private const int srcChannelC4,
|
|
__private const int batchOffset) {
|
|
int2 pos = (int2)(get_global_id(0), get_global_id(1));
|
|
if (pos.x < unitWidth*unitHeight && pos.y < srcChannelC4) {
|
|
int unitWidth_idx = pos.x % unitWidth;
|
|
int unitHeight_idx = pos.x / unitWidth;
|
|
int dstX = mad24(pos.y, unitWidth, unitWidth_idx);
|
|
|
|
{
|
|
int sxStart = unitWidth_idx * 2 - padX;
|
|
int syStart = unitHeight_idx * 2 - padY;
|
|
FLOAT4 S00;
|
|
FLOAT4 S10;
|
|
FLOAT4 S20;
|
|
FLOAT4 S30;
|
|
FLOAT4 S01;
|
|
FLOAT4 S11;
|
|
FLOAT4 S21;
|
|
FLOAT4 S31;
|
|
FLOAT4 S02;
|
|
FLOAT4 S12;
|
|
FLOAT4 S22;
|
|
FLOAT4 S32;
|
|
FLOAT4 S03;
|
|
FLOAT4 S13;
|
|
FLOAT4 S23;
|
|
FLOAT4 S33;
|
|
{
|
|
int sx = 0 + sxStart;
|
|
int sy = 0 + syStart;
|
|
int imageSx = select(sx + pos.y * srcWidth, -1, sx < 0 || sx >= srcWidth);
|
|
int imageSy = select(batchOffset * srcHeight + sy, -1, sy < 0 || sy >= srcHeight);
|
|
S00 = RI_F(uInput, SAMPLER, (int2)(imageSx, imageSy));
|
|
}
|
|
{
|
|
int sx = 1 + sxStart;
|
|
int sy = 0 + syStart;
|
|
int imageSx = select(sx + pos.y * srcWidth, -1, sx < 0 || sx >= srcWidth);
|
|
int imageSy = select(batchOffset * srcHeight + sy, -1, sy < 0 || sy >= srcHeight);
|
|
S10 = RI_F(uInput, SAMPLER, (int2)(imageSx, imageSy));
|
|
}
|
|
{
|
|
int sx = 2 + sxStart;
|
|
int sy = 0 + syStart;
|
|
int imageSx = select(sx + pos.y * srcWidth, -1, sx < 0 || sx >= srcWidth);
|
|
int imageSy = select(batchOffset * srcHeight + sy, -1, sy < 0 || sy >= srcHeight);
|
|
S20 = RI_F(uInput, SAMPLER, (int2)(imageSx, imageSy));
|
|
}
|
|
{
|
|
int sx = 3 + sxStart;
|
|
int sy = 0 + syStart;
|
|
int imageSx = select(sx + pos.y * srcWidth, -1, sx < 0 || sx >= srcWidth);
|
|
int imageSy = select(batchOffset * srcHeight + sy, -1, sy < 0 || sy >= srcHeight);
|
|
S30 = RI_F(uInput, SAMPLER, (int2)(imageSx, imageSy));
|
|
}
|
|
{
|
|
int sx = 0 + sxStart;
|
|
int sy = 1 + syStart;
|
|
int imageSx = select(sx + pos.y * srcWidth, -1, sx < 0 || sx >= srcWidth);
|
|
int imageSy = select(batchOffset * srcHeight + sy, -1, sy < 0 || sy >= srcHeight);
|
|
S01 = RI_F(uInput, SAMPLER, (int2)(imageSx, imageSy));
|
|
}
|
|
{
|
|
int sx = 1 + sxStart;
|
|
int sy = 1 + syStart;
|
|
int imageSx = select(sx + pos.y * srcWidth, -1, sx < 0 || sx >= srcWidth);
|
|
int imageSy = select(batchOffset * srcHeight + sy, -1, sy < 0 || sy >= srcHeight);
|
|
S11 = RI_F(uInput, SAMPLER, (int2)(imageSx, imageSy));
|
|
}
|
|
{
|
|
int sx = 2 + sxStart;
|
|
int sy = 1 + syStart;
|
|
int imageSx = select(sx + pos.y * srcWidth, -1, sx < 0 || sx >= srcWidth);
|
|
int imageSy = select(batchOffset * srcHeight + sy, -1, sy < 0 || sy >= srcHeight);
|
|
S21 = RI_F(uInput, SAMPLER, (int2)(imageSx, imageSy));
|
|
}
|
|
{
|
|
int sx = 3 + sxStart;
|
|
int sy = 1 + syStart;
|
|
int imageSx = select(sx + pos.y * srcWidth, -1, sx < 0 || sx >= srcWidth);
|
|
int imageSy = select(batchOffset * srcHeight + sy, -1, sy < 0 || sy >= srcHeight);
|
|
S31 = RI_F(uInput, SAMPLER, (int2)(imageSx, imageSy));
|
|
}
|
|
{
|
|
int sx = 0 + sxStart;
|
|
int sy = 2 + syStart;
|
|
int imageSx = select(sx + pos.y * srcWidth, -1, sx < 0 || sx >= srcWidth);
|
|
int imageSy = select(batchOffset * srcHeight + sy, -1, sy < 0 || sy >= srcHeight);
|
|
S02 = RI_F(uInput, SAMPLER, (int2)(imageSx, imageSy));
|
|
}
|
|
{
|
|
int sx = 1 + sxStart;
|
|
int sy = 2 + syStart;
|
|
int imageSx = select(sx + pos.y * srcWidth, -1, sx < 0 || sx >= srcWidth);
|
|
int imageSy = select(batchOffset * srcHeight + sy, -1, sy < 0 || sy >= srcHeight);
|
|
S12 = RI_F(uInput, SAMPLER, (int2)(imageSx, imageSy));
|
|
}
|
|
{
|
|
int sx = 2 + sxStart;
|
|
int sy = 2 + syStart;
|
|
int imageSx = select(sx + pos.y * srcWidth, -1, sx < 0 || sx >= srcWidth);
|
|
int imageSy = select(batchOffset * srcHeight + sy, -1, sy < 0 || sy >= srcHeight);
|
|
S22 = RI_F(uInput, SAMPLER, (int2)(imageSx, imageSy));
|
|
}
|
|
{
|
|
int sx = 3 + sxStart;
|
|
int sy = 2 + syStart;
|
|
int imageSx = select(sx + pos.y * srcWidth, -1, sx < 0 || sx >= srcWidth);
|
|
int imageSy = select(batchOffset * srcHeight + sy, -1, sy < 0 || sy >= srcHeight);
|
|
S32 = RI_F(uInput, SAMPLER, (int2)(imageSx, imageSy));
|
|
}
|
|
{
|
|
int sx = 0 + sxStart;
|
|
int sy = 3 + syStart;
|
|
int imageSx = select(sx + pos.y * srcWidth, -1, sx < 0 || sx >= srcWidth);
|
|
int imageSy = select(batchOffset * srcHeight + sy, -1, sy < 0 || sy >= srcHeight);
|
|
S03 = RI_F(uInput, SAMPLER, (int2)(imageSx, imageSy));
|
|
}
|
|
{
|
|
int sx = 1 + sxStart;
|
|
int sy = 3 + syStart;
|
|
int imageSx = select(sx + pos.y * srcWidth, -1, sx < 0 || sx >= srcWidth);
|
|
int imageSy = select(batchOffset * srcHeight + sy, -1, sy < 0 || sy >= srcHeight);
|
|
S13 = RI_F(uInput, SAMPLER, (int2)(imageSx, imageSy));
|
|
}
|
|
{
|
|
int sx = 2 + sxStart;
|
|
int sy = 3 + syStart;
|
|
int imageSx = select(sx + pos.y * srcWidth, -1, sx < 0 || sx >= srcWidth);
|
|
int imageSy = select(batchOffset * srcHeight + sy, -1, sy < 0 || sy >= srcHeight);
|
|
S23 = RI_F(uInput, SAMPLER, (int2)(imageSx, imageSy));
|
|
}
|
|
{
|
|
int sx = 3 + sxStart;
|
|
int sy = 3 + syStart;
|
|
int imageSx = select(sx + pos.y * srcWidth, -1, sx < 0 || sx >= srcWidth);
|
|
int imageSy = select(batchOffset * srcHeight + sy, -1, sy < 0 || sy >= srcHeight);
|
|
S33 = RI_F(uInput, SAMPLER, (int2)(imageSx, imageSy));
|
|
}
|
|
FLOAT4 m00 = +S00 - S02;
|
|
FLOAT4 m10 = +S10 - S12;
|
|
FLOAT4 m20 = +S20 - S22;
|
|
FLOAT4 m30 = +S30 - S32;
|
|
FLOAT4 m01 = +(FLOAT)0.5f * S01 + (FLOAT)0.5f * S02;
|
|
FLOAT4 m11 = +(FLOAT)0.5f * S11 + (FLOAT)0.5f * S12;
|
|
FLOAT4 m21 = +(FLOAT)0.5f * S21 + (FLOAT)0.5f * S22;
|
|
FLOAT4 m31 = +(FLOAT)0.5f * S31 + (FLOAT)0.5f * S32;
|
|
FLOAT4 m02 = -(FLOAT)0.5f * S01 + (FLOAT)0.5f * S02;
|
|
FLOAT4 m12 = -(FLOAT)0.5f * S11 + (FLOAT)0.5f * S12;
|
|
FLOAT4 m22 = -(FLOAT)0.5f * S21 + (FLOAT)0.5f * S22;
|
|
FLOAT4 m32 = -(FLOAT)0.5f * S31 + (FLOAT)0.5f * S32;
|
|
FLOAT4 m03 = -S01 + S03;
|
|
FLOAT4 m13 = -S11 + S13;
|
|
FLOAT4 m23 = -S21 + S23;
|
|
FLOAT4 m33 = -S31 + S33;
|
|
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 0), +m00 - m20);
|
|
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 1), +(FLOAT)0.5f * m10 + (FLOAT)0.5f * m20);
|
|
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 2), -(FLOAT)0.5f * m10 + (FLOAT)0.5f * m20);
|
|
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 3), -m10 + m30);
|
|
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 4), +m01 - m21);
|
|
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 5), +(FLOAT)0.5f * m11 + (FLOAT)0.5f * m21);
|
|
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 6), -(FLOAT)0.5f * m11 + (FLOAT)0.5f * m21);
|
|
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 7), -m11 + m31);
|
|
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 8), +m02 - m22);
|
|
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 9), +(FLOAT)0.5f * m12 + (FLOAT)0.5f * m22);
|
|
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 10), -(FLOAT)0.5f * m12 + (FLOAT)0.5f * m22);
|
|
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 11), -m12 + m32);
|
|
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 12), +m03 - m23);
|
|
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 13), +(FLOAT)0.5f * m13 + (FLOAT)0.5f * m23);
|
|
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 14), -(FLOAT)0.5f * m13 + (FLOAT)0.5f * m23);
|
|
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 15), -m13 + m33);
|
|
}
|
|
}
|
|
}
|