// Vendored PTX for nv12_resize.cu, JIT-compiled by the NVIDIA driver at
// runtime (no CUDA toolkit needed to build this crate). Targets PTX ISA 8.2 /
// sm_52 so it loads on any driver cudarc's cuda-12020 pin supports (>= 535)
// and any GPU that has NVDEC. Regenerate after editing the .cu:
// nvcc -ptx -arch=compute_52 -o nv12_resize.ptx nv12_resize.cu
// using a CUDA 12.2 nvcc (e.g. from nixpkgs 24.05: cudaPackages_12_2.cuda_nvcc
// with -ccbin gcc12 and -I cuda_cudart's include dir).
//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-33191640
// Cuda compilation tools, release 12.2, V12.2.140
// Based on NVVM 7.0.1
//
.version 8.2
.target sm_52
.address_size 64
// .globl resize_luma
.visible .entry resize_luma(
.param .u64 resize_luma_param_0,
.param .u32 resize_luma_param_1,
.param .u32 resize_luma_param_2,
.param .u32 resize_luma_param_3,
.param .u64 resize_luma_param_4,
.param .u32 resize_luma_param_5,
.param .u32 resize_luma_param_6,
.param .u32 resize_luma_param_7
)
{
.reg .pred %p<14>;
.reg .b32 %r<93>;
.reg .b64 %rd<21>;
ld.param.u64 %rd3, [resize_luma_param_0];
ld.param.u32 %r31, [resize_luma_param_1];
ld.param.u32 %r32, [resize_luma_param_2];
ld.param.u32 %r33, [resize_luma_param_3];
ld.param.u64 %rd2, [resize_luma_param_4];
ld.param.u32 %r34, [resize_luma_param_5];
ld.param.u32 %r35, [resize_luma_param_6];
ld.param.u32 %r36, [resize_luma_param_7];
cvta.to.global.u64 %rd1, %rd3;
mov.u32 %r37, %ntid.x;
mov.u32 %r38, %ctaid.x;
mov.u32 %r39, %tid.x;
mad.lo.s32 %r1, %r38, %r37, %r39;
mov.u32 %r40, %ntid.y;
mov.u32 %r41, %ctaid.y;
mov.u32 %r42, %tid.y;
mad.lo.s32 %r2, %r41, %r40, %r42;
setp.ge.u32 %p1, %r1, %r35;
setp.ge.u32 %p2, %r2, %r36;
or.pred %p3, %p1, %p2;
@%p3 bra $L__BB0_14;
mul.lo.s32 %r43, %r1, %r32;
div.u32 %r3, %r43, %r35;
add.s32 %r44, %r43, %r32;
div.u32 %r45, %r44, %r35;
setp.gt.u32 %p4, %r45, %r3;
add.s32 %r4, %r3, 1;
selp.b32 %r5, %r45, %r4, %p4;
mul.lo.s32 %r46, %r2, %r33;
div.u32 %r6, %r46, %r36;
add.s32 %r47, %r46, %r33;
div.u32 %r48, %r47, %r36;
setp.gt.u32 %p5, %r48, %r6;
add.s32 %r49, %r6, 1;
selp.b32 %r7, %r48, %r49, %p5;
setp.gt.u32 %p6, %r7, %r6;
@%p6 bra $L__BB0_3;
bra.uni $L__BB0_2;
$L__BB0_3:
sub.s32 %r52, %r5, %r3;
not.b32 %r53, %r3;
add.s32 %r8, %r5, %r53;
and.b32 %r9, %r52, 3;
add.s32 %r10, %r3, 2;
add.s32 %r11, %r3, 3;
mov.u32 %r91, 0;
mov.u32 %r83, %r6;
$L__BB0_4:
setp.le.u32 %p7, %r5, %r3;
@%p7 bra $L__BB0_12;
setp.eq.s32 %p8, %r9, 0;
mul.lo.s32 %r14, %r83, %r31;
mov.u32 %r85, %r3;
@%p8 bra $L__BB0_9;
setp.eq.s32 %p9, %r9, 1;
add.s32 %r55, %r3, %r14;
cvt.u64.u32 %rd4, %r55;
add.s64 %rd5, %rd1, %rd4;
ld.global.u8 %r56, [%rd5];
add.s32 %r91, %r91, %r56;
mov.u32 %r85, %r4;
@%p9 bra $L__BB0_9;
setp.eq.s32 %p10, %r9, 2;
add.s32 %r57, %r4, %r14;
cvt.u64.u32 %rd6, %r57;
add.s64 %rd7, %rd1, %rd6;
ld.global.u8 %r58, [%rd7];
add.s32 %r91, %r91, %r58;
mov.u32 %r85, %r10;
@%p10 bra $L__BB0_9;
add.s32 %r59, %r10, %r14;
cvt.u64.u32 %rd8, %r59;
add.s64 %rd9, %rd1, %rd8;
ld.global.u8 %r60, [%rd9];
add.s32 %r91, %r91, %r60;
mov.u32 %r85, %r11;
$L__BB0_9:
setp.lt.u32 %p11, %r8, 3;
@%p11 bra $L__BB0_12;
add.s32 %r61, %r85, %r14;
add.s32 %r88, %r61, 3;
$L__BB0_11:
add.s32 %r62, %r14, %r85;
cvt.u64.u32 %rd10, %r62;
add.s64 %rd11, %rd1, %rd10;
ld.global.u8 %r63, [%rd11];
add.s32 %r64, %r91, %r63;
add.s32 %r65, %r88, -2;
cvt.u64.u32 %rd12, %r65;
add.s64 %rd13, %rd1, %rd12;
ld.global.u8 %r66, [%rd13];
add.s32 %r67, %r64, %r66;
add.s32 %r68, %r88, -1;
cvt.u64.u32 %rd14, %r68;
add.s64 %rd15, %rd1, %rd14;
ld.global.u8 %r69, [%rd15];
add.s32 %r70, %r67, %r69;
cvt.u64.u32 %rd16, %r88;
add.s64 %rd17, %rd1, %rd16;
ld.global.u8 %r71, [%rd17];
add.s32 %r91, %r70, %r71;
add.s32 %r88, %r88, 4;
add.s32 %r85, %r85, 4;
setp.lt.u32 %p12, %r85, %r5;
@%p12 bra $L__BB0_11;
$L__BB0_12:
add.s32 %r83, %r83, 1;
setp.lt.u32 %p13, %r83, %r7;
@%p13 bra $L__BB0_4;
bra.uni $L__BB0_13;
$L__BB0_2:
mov.u32 %r91, 0;
$L__BB0_13:
sub.s32 %r72, %r7, %r6;
sub.s32 %r73, %r5, %r3;
mul.lo.s32 %r74, %r72, %r73;
shr.u32 %r75, %r74, 1;
add.s32 %r76, %r91, %r75;
div.u32 %r77, %r76, %r74;
mad.lo.s32 %r82, %r2, %r34, %r1;
cvt.u64.u32 %rd18, %r82;
cvta.to.global.u64 %rd19, %rd2;
add.s64 %rd20, %rd19, %rd18;
st.global.u8 [%rd20], %r77;
$L__BB0_14:
ret;
}
// .globl resize_chroma
.visible .entry resize_chroma(
.param .u64 resize_chroma_param_0,
.param .u32 resize_chroma_param_1,
.param .u32 resize_chroma_param_2,
.param .u32 resize_chroma_param_3,
.param .u64 resize_chroma_param_4,
.param .u32 resize_chroma_param_5,
.param .u32 resize_chroma_param_6,
.param .u32 resize_chroma_param_7
)
{
.reg .pred %p<14>;
.reg .b32 %r<128>;
.reg .b64 %rd<37>;
ld.param.u64 %rd3, [resize_chroma_param_0];
ld.param.u32 %r41, [resize_chroma_param_1];
ld.param.u32 %r42, [resize_chroma_param_2];
ld.param.u32 %r43, [resize_chroma_param_3];
ld.param.u64 %rd2, [resize_chroma_param_4];
ld.param.u32 %r44, [resize_chroma_param_5];
ld.param.u32 %r45, [resize_chroma_param_6];
ld.param.u32 %r46, [resize_chroma_param_7];
cvta.to.global.u64 %rd1, %rd3;
mov.u32 %r47, %ntid.x;
mov.u32 %r48, %ctaid.x;
mov.u32 %r49, %tid.x;
mad.lo.s32 %r1, %r48, %r47, %r49;
mov.u32 %r50, %ntid.y;
mov.u32 %r51, %ctaid.y;
mov.u32 %r52, %tid.y;
mad.lo.s32 %r2, %r51, %r50, %r52;
setp.ge.u32 %p1, %r1, %r45;
setp.ge.u32 %p2, %r2, %r46;
or.pred %p3, %p1, %p2;
@%p3 bra $L__BB1_13;
mul.lo.s32 %r53, %r1, %r42;
div.u32 %r3, %r53, %r45;
add.s32 %r54, %r53, %r42;
div.u32 %r55, %r54, %r45;
setp.gt.u32 %p4, %r55, %r3;
add.s32 %r4, %r3, 1;
selp.b32 %r5, %r55, %r4, %p4;
mul.lo.s32 %r56, %r2, %r43;
div.u32 %r6, %r56, %r46;
add.s32 %r57, %r56, %r43;
div.u32 %r58, %r57, %r46;
setp.gt.u32 %p5, %r58, %r6;
add.s32 %r59, %r6, 1;
selp.b32 %r7, %r58, %r59, %p5;
setp.gt.u32 %p6, %r7, %r6;
@%p6 bra $L__BB1_3;
bra.uni $L__BB1_2;
$L__BB1_3:
sub.s32 %r64, %r5, %r3;
not.b32 %r65, %r3;
add.s32 %r8, %r5, %r65;
and.b32 %r9, %r64, 3;
shl.b32 %r10, %r3, 1;
shl.b32 %r11, %r4, 1;
add.s32 %r12, %r3, 2;
shl.b32 %r13, %r12, 1;
add.s32 %r14, %r3, 3;
mov.u32 %r125, 0;
mov.u32 %r113, %r6;
mov.u32 %r124, %r125;
$L__BB1_4:
setp.le.u32 %p7, %r5, %r3;
@%p7 bra $L__BB1_11;
setp.eq.s32 %p8, %r9, 0;
mul.lo.s32 %r18, %r113, %r41;
mov.u32 %r116, %r3;
@%p8 bra $L__BB1_9;
setp.eq.s32 %p9, %r9, 1;
add.s32 %r67, %r10, %r18;
cvt.u64.u32 %rd4, %r67;
add.s64 %rd5, %rd1, %rd4;
ld.global.u8 %r68, [%rd5];
add.s32 %r124, %r124, %r68;
add.s32 %r69, %r67, 1;
cvt.u64.u32 %rd6, %r69;
add.s64 %rd7, %rd1, %rd6;
ld.global.u8 %r70, [%rd7];
add.s32 %r125, %r125, %r70;
mov.u32 %r116, %r4;
@%p9 bra $L__BB1_9;
setp.eq.s32 %p10, %r9, 2;
add.s32 %r71, %r11, %r18;
cvt.u64.u32 %rd8, %r71;
add.s64 %rd9, %rd1, %rd8;
ld.global.u8 %r72, [%rd9];
add.s32 %r124, %r124, %r72;
add.s32 %r73, %r71, 1;
cvt.u64.u32 %rd10, %r73;
add.s64 %rd11, %rd1, %rd10;
ld.global.u8 %r74, [%rd11];
add.s32 %r125, %r125, %r74;
mov.u32 %r116, %r12;
@%p10 bra $L__BB1_9;
add.s32 %r75, %r13, %r18;
cvt.u64.u32 %rd12, %r75;
add.s64 %rd13, %rd1, %rd12;
ld.global.u8 %r76, [%rd13];
add.s32 %r124, %r124, %r76;
add.s32 %r77, %r75, 1;
cvt.u64.u32 %rd14, %r77;
add.s64 %rd15, %rd1, %rd14;
ld.global.u8 %r78, [%rd15];
add.s32 %r125, %r125, %r78;
mov.u32 %r116, %r14;
$L__BB1_9:
setp.lt.u32 %p11, %r8, 3;
@%p11 bra $L__BB1_11;
$L__BB1_10:
shl.b32 %r79, %r116, 1;
add.s32 %r80, %r79, %r18;
cvt.u64.u32 %rd16, %r80;
add.s64 %rd17, %rd1, %rd16;
ld.global.u8 %r81, [%rd17];
add.s32 %r82, %r124, %r81;
add.s32 %r83, %r80, 1;
cvt.u64.u32 %rd18, %r83;
add.s64 %rd19, %rd1, %rd18;
ld.global.u8 %r84, [%rd19];
add.s32 %r85, %r125, %r84;
add.s32 %r86, %r80, 2;
cvt.u64.u32 %rd20, %r86;
add.s64 %rd21, %rd1, %rd20;
ld.global.u8 %r87, [%rd21];
add.s32 %r88, %r82, %r87;
add.s32 %r89, %r80, 3;
cvt.u64.u32 %rd22, %r89;
add.s64 %rd23, %rd1, %rd22;
ld.global.u8 %r90, [%rd23];
add.s32 %r91, %r85, %r90;
add.s32 %r92, %r80, 4;
cvt.u64.u32 %rd24, %r92;
add.s64 %rd25, %rd1, %rd24;
ld.global.u8 %r93, [%rd25];
add.s32 %r94, %r88, %r93;
add.s32 %r95, %r80, 5;
cvt.u64.u32 %rd26, %r95;
add.s64 %rd27, %rd1, %rd26;
ld.global.u8 %r96, [%rd27];
add.s32 %r97, %r91, %r96;
add.s32 %r98, %r80, 6;
cvt.u64.u32 %rd28, %r98;
add.s64 %rd29, %rd1, %rd28;
ld.global.u8 %r99, [%rd29];
add.s32 %r124, %r94, %r99;
add.s32 %r100, %r80, 7;
cvt.u64.u32 %rd30, %r100;
add.s64 %rd31, %rd1, %rd30;
ld.global.u8 %r101, [%rd31];
add.s32 %r125, %r97, %r101;
add.s32 %r116, %r116, 4;
setp.lt.u32 %p12, %r116, %r5;
@%p12 bra $L__BB1_10;
$L__BB1_11:
add.s32 %r113, %r113, 1;
setp.lt.u32 %p13, %r113, %r7;
@%p13 bra $L__BB1_4;
bra.uni $L__BB1_12;
$L__BB1_2:
mov.u32 %r124, 0;
mov.u32 %r125, %r124;
$L__BB1_12:
cvta.to.global.u64 %rd32, %rd2;
sub.s32 %r102, %r7, %r6;
sub.s32 %r103, %r5, %r3;
mul.lo.s32 %r104, %r102, %r103;
shr.u32 %r105, %r104, 1;
add.s32 %r106, %r124, %r105;
div.u32 %r107, %r106, %r104;
shl.b32 %r108, %r1, 1;
mad.lo.s32 %r109, %r2, %r44, %r108;
cvt.u64.u32 %rd33, %r109;
add.s64 %rd34, %rd32, %rd33;
st.global.u8 [%rd34], %r107;
add.s32 %r110, %r125, %r105;
div.u32 %r111, %r110, %r104;
add.s32 %r112, %r109, 1;
cvt.u64.u32 %rd35, %r112;
add.s64 %rd36, %rd32, %rd35;
st.global.u8 [%rd36], %r111;
$L__BB1_13:
ret;
}