Files
alibaba--mnn/source/backend/opencl/execution/cl/winogradTransformSource2_5_1.cl
2026-07-13 13:33:03 +08:00

420 lines
25 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 S40;
FLOAT4 S50;
FLOAT4 S01;
FLOAT4 S11;
FLOAT4 S21;
FLOAT4 S31;
FLOAT4 S41;
FLOAT4 S51;
FLOAT4 S02;
FLOAT4 S12;
FLOAT4 S22;
FLOAT4 S32;
FLOAT4 S42;
FLOAT4 S52;
FLOAT4 S03;
FLOAT4 S13;
FLOAT4 S23;
FLOAT4 S33;
FLOAT4 S43;
FLOAT4 S53;
FLOAT4 S04;
FLOAT4 S14;
FLOAT4 S24;
FLOAT4 S34;
FLOAT4 S44;
FLOAT4 S54;
FLOAT4 S05;
FLOAT4 S15;
FLOAT4 S25;
FLOAT4 S35;
FLOAT4 S45;
FLOAT4 S55;
{
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 = 4 + 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);
S40 = RI_F(uInput, SAMPLER, (int2)(imageSx, imageSy));
}
{
int sx = 5 + 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);
S50 = 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 = 4 + 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);
S41 = RI_F(uInput, SAMPLER, (int2)(imageSx, imageSy));
}
{
int sx = 5 + 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);
S51 = 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 = 4 + 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);
S42 = RI_F(uInput, SAMPLER, (int2)(imageSx, imageSy));
}
{
int sx = 5 + 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);
S52 = 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));
}
{
int sx = 4 + 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);
S43 = RI_F(uInput, SAMPLER, (int2)(imageSx, imageSy));
}
{
int sx = 5 + 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);
S53 = RI_F(uInput, SAMPLER, (int2)(imageSx, imageSy));
}
{
int sx = 0 + sxStart;
int sy = 4 + syStart;
int imageSx = select(sx + pos.y * srcWidth, -1, sx < 0 || sx >= srcWidth);
int imageSy = select(batchOffset * srcHeight + sy, -1, sy < 0 || sy >= srcHeight);
S04 = RI_F(uInput, SAMPLER, (int2)(imageSx, imageSy));
}
{
int sx = 1 + sxStart;
int sy = 4 + syStart;
int imageSx = select(sx + pos.y * srcWidth, -1, sx < 0 || sx >= srcWidth);
int imageSy = select(batchOffset * srcHeight + sy, -1, sy < 0 || sy >= srcHeight);
S14 = RI_F(uInput, SAMPLER, (int2)(imageSx, imageSy));
}
{
int sx = 2 + sxStart;
int sy = 4 + syStart;
int imageSx = select(sx + pos.y * srcWidth, -1, sx < 0 || sx >= srcWidth);
int imageSy = select(batchOffset * srcHeight + sy, -1, sy < 0 || sy >= srcHeight);
S24 = RI_F(uInput, SAMPLER, (int2)(imageSx, imageSy));
}
{
int sx = 3 + sxStart;
int sy = 4 + syStart;
int imageSx = select(sx + pos.y * srcWidth, -1, sx < 0 || sx >= srcWidth);
int imageSy = select(batchOffset * srcHeight + sy, -1, sy < 0 || sy >= srcHeight);
S34 = RI_F(uInput, SAMPLER, (int2)(imageSx, imageSy));
}
{
int sx = 4 + sxStart;
int sy = 4 + syStart;
int imageSx = select(sx + pos.y * srcWidth, -1, sx < 0 || sx >= srcWidth);
int imageSy = select(batchOffset * srcHeight + sy, -1, sy < 0 || sy >= srcHeight);
S44 = RI_F(uInput, SAMPLER, (int2)(imageSx, imageSy));
}
{
int sx = 5 + sxStart;
int sy = 4 + syStart;
int imageSx = select(sx + pos.y * srcWidth, -1, sx < 0 || sx >= srcWidth);
int imageSy = select(batchOffset * srcHeight + sy, -1, sy < 0 || sy >= srcHeight);
S54 = RI_F(uInput, SAMPLER, (int2)(imageSx, imageSy));
}
{
int sx = 0 + sxStart;
int sy = 5 + syStart;
int imageSx = select(sx + pos.y * srcWidth, -1, sx < 0 || sx >= srcWidth);
int imageSy = select(batchOffset * srcHeight + sy, -1, sy < 0 || sy >= srcHeight);
S05 = RI_F(uInput, SAMPLER, (int2)(imageSx, imageSy));
}
{
int sx = 1 + sxStart;
int sy = 5 + syStart;
int imageSx = select(sx + pos.y * srcWidth, -1, sx < 0 || sx >= srcWidth);
int imageSy = select(batchOffset * srcHeight + sy, -1, sy < 0 || sy >= srcHeight);
S15 = RI_F(uInput, SAMPLER, (int2)(imageSx, imageSy));
}
{
int sx = 2 + sxStart;
int sy = 5 + syStart;
int imageSx = select(sx + pos.y * srcWidth, -1, sx < 0 || sx >= srcWidth);
int imageSy = select(batchOffset * srcHeight + sy, -1, sy < 0 || sy >= srcHeight);
S25 = RI_F(uInput, SAMPLER, (int2)(imageSx, imageSy));
}
{
int sx = 3 + sxStart;
int sy = 5 + syStart;
int imageSx = select(sx + pos.y * srcWidth, -1, sx < 0 || sx >= srcWidth);
int imageSy = select(batchOffset * srcHeight + sy, -1, sy < 0 || sy >= srcHeight);
S35 = RI_F(uInput, SAMPLER, (int2)(imageSx, imageSy));
}
{
int sx = 4 + sxStart;
int sy = 5 + syStart;
int imageSx = select(sx + pos.y * srcWidth, -1, sx < 0 || sx >= srcWidth);
int imageSy = select(batchOffset * srcHeight + sy, -1, sy < 0 || sy >= srcHeight);
S45 = RI_F(uInput, SAMPLER, (int2)(imageSx, imageSy));
}
{
int sx = 5 + sxStart;
int sy = 5 + syStart;
int imageSx = select(sx + pos.y * srcWidth, -1, sx < 0 || sx >= srcWidth);
int imageSy = select(batchOffset * srcHeight + sy, -1, sy < 0 || sy >= srcHeight);
S55 = RI_F(uInput, SAMPLER, (int2)(imageSx, imageSy));
}
FLOAT4 m00 = +S00 - (FLOAT)1.25 * S02 + (FLOAT)0.25 * S04;
FLOAT4 m10 = +S10 - (FLOAT)1.25 * S12 + (FLOAT)0.25 * S14;
FLOAT4 m20 = +S20 - (FLOAT)1.25 * S22 + (FLOAT)0.25 * S24;
FLOAT4 m30 = +S30 - (FLOAT)1.25 * S32 + (FLOAT)0.25 * S34;
FLOAT4 m40 = +S40 - (FLOAT)1.25 * S42 + (FLOAT)0.25 * S44;
FLOAT4 m50 = +S50 - (FLOAT)1.25 * S52 + (FLOAT)0.25 * S54;
FLOAT4 m01 = +(FLOAT)0.666667 * S01 + (FLOAT)0.666667 * S02 - (FLOAT)0.166667 * S03 - (FLOAT)0.166667 * S04;
FLOAT4 m11 = +(FLOAT)0.666667 * S11 + (FLOAT)0.666667 * S12 - (FLOAT)0.166667 * S13 - (FLOAT)0.166667 * S14;
FLOAT4 m21 = +(FLOAT)0.666667 * S21 + (FLOAT)0.666667 * S22 - (FLOAT)0.166667 * S23 - (FLOAT)0.166667 * S24;
FLOAT4 m31 = +(FLOAT)0.666667 * S31 + (FLOAT)0.666667 * S32 - (FLOAT)0.166667 * S33 - (FLOAT)0.166667 * S34;
FLOAT4 m41 = +(FLOAT)0.666667 * S41 + (FLOAT)0.666667 * S42 - (FLOAT)0.166667 * S43 - (FLOAT)0.166667 * S44;
FLOAT4 m51 = +(FLOAT)0.666667 * S51 + (FLOAT)0.666667 * S52 - (FLOAT)0.166667 * S53 - (FLOAT)0.166667 * S54;
FLOAT4 m02 = -(FLOAT)0.666667 * S01 + (FLOAT)0.666667 * S02 + (FLOAT)0.166667 * S03 - (FLOAT)0.166667 * S04;
FLOAT4 m12 = -(FLOAT)0.666667 * S11 + (FLOAT)0.666667 * S12 + (FLOAT)0.166667 * S13 - (FLOAT)0.166667 * S14;
FLOAT4 m22 = -(FLOAT)0.666667 * S21 + (FLOAT)0.666667 * S22 + (FLOAT)0.166667 * S23 - (FLOAT)0.166667 * S24;
FLOAT4 m32 = -(FLOAT)0.666667 * S31 + (FLOAT)0.666667 * S32 + (FLOAT)0.166667 * S33 - (FLOAT)0.166667 * S34;
FLOAT4 m42 = -(FLOAT)0.666667 * S41 + (FLOAT)0.666667 * S42 + (FLOAT)0.166667 * S43 - (FLOAT)0.166667 * S44;
FLOAT4 m52 = -(FLOAT)0.666667 * S51 + (FLOAT)0.666667 * S52 + (FLOAT)0.166667 * S53 - (FLOAT)0.166667 * S54;
FLOAT4 m03 =
-(FLOAT)0.0833333 * S01 - (FLOAT)0.0416667 * S02 + (FLOAT)0.0833333 * S03 + (FLOAT)0.0416667 * S04;
FLOAT4 m13 =
-(FLOAT)0.0833333 * S11 - (FLOAT)0.0416667 * S12 + (FLOAT)0.0833333 * S13 + (FLOAT)0.0416667 * S14;
FLOAT4 m23 =
-(FLOAT)0.0833333 * S21 - (FLOAT)0.0416667 * S22 + (FLOAT)0.0833333 * S23 + (FLOAT)0.0416667 * S24;
FLOAT4 m33 =
-(FLOAT)0.0833333 * S31 - (FLOAT)0.0416667 * S32 + (FLOAT)0.0833333 * S33 + (FLOAT)0.0416667 * S34;
FLOAT4 m43 =
-(FLOAT)0.0833333 * S41 - (FLOAT)0.0416667 * S42 + (FLOAT)0.0833333 * S43 + (FLOAT)0.0416667 * S44;
FLOAT4 m53 =
-(FLOAT)0.0833333 * S51 - (FLOAT)0.0416667 * S52 + (FLOAT)0.0833333 * S53 + (FLOAT)0.0416667 * S54;
FLOAT4 m04 =
+(FLOAT)0.0833333 * S01 - (FLOAT)0.0416667 * S02 - (FLOAT)0.0833333 * S03 + (FLOAT)0.0416667 * S04;
FLOAT4 m14 =
+(FLOAT)0.0833333 * S11 - (FLOAT)0.0416667 * S12 - (FLOAT)0.0833333 * S13 + (FLOAT)0.0416667 * S14;
FLOAT4 m24 =
+(FLOAT)0.0833333 * S21 - (FLOAT)0.0416667 * S22 - (FLOAT)0.0833333 * S23 + (FLOAT)0.0416667 * S24;
FLOAT4 m34 =
+(FLOAT)0.0833333 * S31 - (FLOAT)0.0416667 * S32 - (FLOAT)0.0833333 * S33 + (FLOAT)0.0416667 * S34;
FLOAT4 m44 =
+(FLOAT)0.0833333 * S41 - (FLOAT)0.0416667 * S42 - (FLOAT)0.0833333 * S43 + (FLOAT)0.0416667 * S44;
FLOAT4 m54 =
+(FLOAT)0.0833333 * S51 - (FLOAT)0.0416667 * S52 - (FLOAT)0.0833333 * S53 + (FLOAT)0.0416667 * S54;
FLOAT4 m05 = +(FLOAT)4.0 * S01 - (FLOAT)5.0 * S03 + S05;
FLOAT4 m15 = +(FLOAT)4.0 * S11 - (FLOAT)5.0 * S13 + S15;
FLOAT4 m25 = +(FLOAT)4.0 * S21 - (FLOAT)5.0 * S23 + S25;
FLOAT4 m35 = +(FLOAT)4.0 * S31 - (FLOAT)5.0 * S33 + S35;
FLOAT4 m45 = +(FLOAT)4.0 * S41 - (FLOAT)5.0 * S43 + S45;
FLOAT4 m55 = +(FLOAT)4.0 * S51 - (FLOAT)5.0 * S53 + S55;
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 0), +m00 - (FLOAT)1.25 * m20 + (FLOAT)0.25 * m40);
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 1),
+(FLOAT)0.666667 * m10 + (FLOAT)0.666667 * m20 - (FLOAT)0.166667 * m30 - (FLOAT)0.166667 * m40);
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 2),
-(FLOAT)0.666667 * m10 + (FLOAT)0.666667 * m20 + (FLOAT)0.166667 * m30 - (FLOAT)0.166667 * m40);
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 3),
-(FLOAT)0.0833333 * m10 - (FLOAT)0.0416667 * m20 + (FLOAT)0.0833333 * m30 + (FLOAT)0.0416667 * m40);
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 4),
+(FLOAT)0.0833333 * m10 - (FLOAT)0.0416667 * m20 - (FLOAT)0.0833333 * m30 + (FLOAT)0.0416667 * m40);
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 5), +(FLOAT)4.0 * m10 - (FLOAT)5.0 * m30 + m50);
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 6), +m01 - (FLOAT)1.25 * m21 + (FLOAT)0.25 * m41);
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 7),
+(FLOAT)0.666667 * m11 + (FLOAT)0.666667 * m21 - (FLOAT)0.166667 * m31 - (FLOAT)0.166667 * m41);
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 8),
-(FLOAT)0.666667 * m11 + (FLOAT)0.666667 * m21 + (FLOAT)0.166667 * m31 - (FLOAT)0.166667 * m41);
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 9),
-(FLOAT)0.0833333 * m11 - (FLOAT)0.0416667 * m21 + (FLOAT)0.0833333 * m31 + (FLOAT)0.0416667 * m41);
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 10),
+(FLOAT)0.0833333 * m11 - (FLOAT)0.0416667 * m21 - (FLOAT)0.0833333 * m31 + (FLOAT)0.0416667 * m41);
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 11), +(FLOAT)4.0 * m11 - (FLOAT)5.0 * m31 + m51);
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 12), +m02 - (FLOAT)1.25 * m22 + (FLOAT)0.25 * m42);
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 13),
+(FLOAT)0.666667 * m12 + (FLOAT)0.666667 * m22 - (FLOAT)0.166667 * m32 - (FLOAT)0.166667 * m42);
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 14),
-(FLOAT)0.666667 * m12 + (FLOAT)0.666667 * m22 + (FLOAT)0.166667 * m32 - (FLOAT)0.166667 * m42);
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 15),
-(FLOAT)0.0833333 * m12 - (FLOAT)0.0416667 * m22 + (FLOAT)0.0833333 * m32 + (FLOAT)0.0416667 * m42);
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 16),
+(FLOAT)0.0833333 * m12 - (FLOAT)0.0416667 * m22 - (FLOAT)0.0833333 * m32 + (FLOAT)0.0416667 * m42);
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 17), +(FLOAT)4.0 * m12 - (FLOAT)5.0 * m32 + m52);
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 18), +m03 - (FLOAT)1.25 * m23 + (FLOAT)0.25 * m43);
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 19),
+(FLOAT)0.666667 * m13 + (FLOAT)0.666667 * m23 - (FLOAT)0.166667 * m33 - (FLOAT)0.166667 * m43);
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 20),
-(FLOAT)0.666667 * m13 + (FLOAT)0.666667 * m23 + (FLOAT)0.166667 * m33 - (FLOAT)0.166667 * m43);
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 21),
-(FLOAT)0.0833333 * m13 - (FLOAT)0.0416667 * m23 + (FLOAT)0.0833333 * m33 + (FLOAT)0.0416667 * m43);
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 22),
+(FLOAT)0.0833333 * m13 - (FLOAT)0.0416667 * m23 - (FLOAT)0.0833333 * m33 + (FLOAT)0.0416667 * m43);
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 23), +(FLOAT)4.0 * m13 - (FLOAT)5.0 * m33 + m53);
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 24), +m04 - (FLOAT)1.25 * m24 + (FLOAT)0.25 * m44);
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 25),
+(FLOAT)0.666667 * m14 + (FLOAT)0.666667 * m24 - (FLOAT)0.166667 * m34 - (FLOAT)0.166667 * m44);
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 26),
-(FLOAT)0.666667 * m14 + (FLOAT)0.666667 * m24 + (FLOAT)0.166667 * m34 - (FLOAT)0.166667 * m44);
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 27),
-(FLOAT)0.0833333 * m14 - (FLOAT)0.0416667 * m24 + (FLOAT)0.0833333 * m34 + (FLOAT)0.0416667 * m44);
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 28),
+(FLOAT)0.0833333 * m14 - (FLOAT)0.0416667 * m24 - (FLOAT)0.0833333 * m34 + (FLOAT)0.0416667 * m44);
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 29), +(FLOAT)4.0 * m14 - (FLOAT)5.0 * m34 + m54);
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 30), +m05 - (FLOAT)1.25 * m25 + (FLOAT)0.25 * m45);
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 31),
+(FLOAT)0.666667 * m15 + (FLOAT)0.666667 * m25 - (FLOAT)0.166667 * m35 - (FLOAT)0.166667 * m45);
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 32),
-(FLOAT)0.666667 * m15 + (FLOAT)0.666667 * m25 + (FLOAT)0.166667 * m35 - (FLOAT)0.166667 * m45);
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 33),
-(FLOAT)0.0833333 * m15 - (FLOAT)0.0416667 * m25 + (FLOAT)0.0833333 * m35 + (FLOAT)0.0416667 * m45);
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 34),
+(FLOAT)0.0833333 * m15 - (FLOAT)0.0416667 * m25 - (FLOAT)0.0833333 * m35 + (FLOAT)0.0416667 * m45);
WI_F(uOutput, (int2)(dstX, unitHeight_idx + unitHeight * 35), +(FLOAT)4.0 * m15 - (FLOAT)5.0 * m35 + m55);
}
}
}