// Vendored PTX for rgba_to_nv12.cu, JIT-compiled by the NVIDIA driver at
// runtime (no CUDA toolkit needed to build this crate). Targets PTX ISA 8.2 /
// sm_52 like nv12_resize.ptx, so it loads on any driver cudarc's cuda-12020 pin
// supports (>= 535). Regenerate after editing the .cu:
// nvcc -ptx -arch=compute_52 -o rgba_to_nv12.ptx rgba_to_nv12.cu
// using a CUDA 12.2 nvcc (e.g. the docker.io/nvidia/cuda:12.2.2-devel-ubuntu22.04
// image, which needs no GPU to emit PTX).
//
// 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 rgba_to_nv12
.visible .entry rgba_to_nv12(
.param .u64 rgba_to_nv12_param_0,
.param .u32 rgba_to_nv12_param_1,
.param .u32 rgba_to_nv12_param_2,
.param .u32 rgba_to_nv12_param_3,
.param .align 16 .b8 rgba_to_nv12_param_4[16],
.param .align 16 .b8 rgba_to_nv12_param_5[16],
.param .align 16 .b8 rgba_to_nv12_param_6[16],
.param .u64 rgba_to_nv12_param_7,
.param .u32 rgba_to_nv12_param_8
)
{
.reg .pred %p<5>;
.reg .b16 %rs<57>;
.reg .f32 %f<130>;
.reg .b32 %r<36>;
.reg .b64 %rd<17>;
ld.param.u64 %rd6, [rgba_to_nv12_param_0];
ld.param.u32 %r8, [rgba_to_nv12_param_1];
ld.param.u32 %r5, [rgba_to_nv12_param_2];
ld.param.u32 %r6, [rgba_to_nv12_param_3];
ld.param.u64 %rd7, [rgba_to_nv12_param_7];
ld.param.u32 %r7, [rgba_to_nv12_param_8];
ld.param.f32 %f21, [rgba_to_nv12_param_6+12];
ld.param.f32 %f20, [rgba_to_nv12_param_6+8];
ld.param.f32 %f19, [rgba_to_nv12_param_6+4];
ld.param.f32 %f18, [rgba_to_nv12_param_6];
ld.param.f32 %f17, [rgba_to_nv12_param_5+12];
ld.param.f32 %f16, [rgba_to_nv12_param_5+8];
ld.param.f32 %f15, [rgba_to_nv12_param_5+4];
ld.param.f32 %f14, [rgba_to_nv12_param_5];
ld.param.f32 %f13, [rgba_to_nv12_param_4+12];
ld.param.f32 %f12, [rgba_to_nv12_param_4+8];
ld.param.f32 %f11, [rgba_to_nv12_param_4+4];
ld.param.f32 %f10, [rgba_to_nv12_param_4];
cvta.to.global.u64 %rd1, %rd7;
mov.u32 %r9, %ntid.x;
mov.u32 %r10, %ctaid.x;
mov.u32 %r11, %tid.x;
mad.lo.s32 %r1, %r10, %r9, %r11;
shl.b32 %r2, %r1, 1;
mov.u32 %r12, %ntid.y;
mov.u32 %r13, %ctaid.y;
mov.u32 %r14, %tid.y;
mad.lo.s32 %r3, %r13, %r12, %r14;
shl.b32 %r4, %r3, 1;
setp.ge.u32 %p1, %r2, %r8;
setp.ge.u32 %p2, %r4, %r5;
or.pred %p3, %p1, %p2;
@%p3 bra $L__BB0_5;
cvt.u64.u32 %rd2, %r2;
mul.lo.s32 %r15, %r4, %r7;
cvt.u64.u32 %rd8, %r15;
add.s64 %rd9, %rd8, %rd2;
add.s64 %rd3, %rd1, %rd9;
cvt.u64.u32 %rd10, %r7;
add.s64 %rd4, %rd3, %rd10;
add.s32 %r16, %r7, 1;
cvt.u64.u32 %rd11, %r16;
add.s64 %rd5, %rd3, %rd11;
setp.eq.s32 %p4, %r6, 0;
@%p4 bra $L__BB0_3;
shl.b32 %r17, %r1, 3;
suld.b.2d.v4.b8.trap {%rs1, %rs2, %rs3, %rs4}, [%rd6, {%r17, %r4}];
and.b16 %rs5, %rs1, 255;
and.b16 %rs6, %rs2, 255;
and.b16 %rs7, %rs3, 255;
cvt.rn.f32.u16 %f22, %rs7;
cvt.rn.f32.u16 %f23, %rs6;
cvt.rn.f32.u16 %f24, %rs5;
mul.f32 %f25, %f10, %f22;
fma.rn.f32 %f26, %f11, %f23, %f25;
fma.rn.f32 %f27, %f12, %f24, %f26;
add.f32 %f28, %f13, %f27;
mov.f32 %f29, 0f00000000;
max.f32 %f30, %f28, %f29;
mov.f32 %f31, 0f437F0000;
min.f32 %f32, %f30, %f31;
cvt.rni.s32.f32 %r18, %f32;
st.global.u8 [%rd3], %r18;
or.b32 %r19, %r17, 4;
suld.b.2d.v4.b8.trap {%rs8, %rs9, %rs10, %rs11}, [%rd6, {%r19, %r4}];
and.b16 %rs12, %rs8, 255;
and.b16 %rs13, %rs9, 255;
and.b16 %rs14, %rs10, 255;
cvt.rn.f32.u16 %f33, %rs14;
cvt.rn.f32.u16 %f34, %rs13;
cvt.rn.f32.u16 %f35, %rs12;
add.f32 %f36, %f22, %f33;
add.f32 %f37, %f23, %f34;
add.f32 %f38, %f24, %f35;
mul.f32 %f39, %f10, %f33;
fma.rn.f32 %f40, %f11, %f34, %f39;
fma.rn.f32 %f41, %f12, %f35, %f40;
add.f32 %f42, %f13, %f41;
max.f32 %f43, %f42, %f29;
min.f32 %f44, %f43, %f31;
cvt.rni.s32.f32 %r20, %f44;
st.global.u8 [%rd3+1], %r20;
add.s32 %r21, %r4, 1;
suld.b.2d.v4.b8.trap {%rs15, %rs16, %rs17, %rs18}, [%rd6, {%r17, %r21}];
and.b16 %rs19, %rs15, 255;
and.b16 %rs20, %rs16, 255;
and.b16 %rs21, %rs17, 255;
cvt.rn.f32.u16 %f45, %rs21;
cvt.rn.f32.u16 %f46, %rs20;
cvt.rn.f32.u16 %f47, %rs19;
add.f32 %f48, %f36, %f45;
add.f32 %f49, %f37, %f46;
add.f32 %f50, %f38, %f47;
mul.f32 %f51, %f10, %f45;
fma.rn.f32 %f52, %f11, %f46, %f51;
fma.rn.f32 %f53, %f12, %f47, %f52;
add.f32 %f54, %f13, %f53;
max.f32 %f55, %f54, %f29;
min.f32 %f56, %f55, %f31;
cvt.rni.s32.f32 %r22, %f56;
st.global.u8 [%rd4], %r22;
suld.b.2d.v4.b8.trap {%rs22, %rs23, %rs24, %rs25}, [%rd6, {%r19, %r21}];
and.b16 %rs26, %rs22, 255;
and.b16 %rs27, %rs23, 255;
and.b16 %rs28, %rs24, 255;
cvt.rn.f32.u16 %f57, %rs28;
cvt.rn.f32.u16 %f58, %rs27;
cvt.rn.f32.u16 %f59, %rs26;
add.f32 %f127, %f48, %f57;
add.f32 %f128, %f49, %f58;
add.f32 %f129, %f50, %f59;
mul.f32 %f60, %f10, %f57;
fma.rn.f32 %f61, %f11, %f58, %f60;
fma.rn.f32 %f62, %f12, %f59, %f61;
add.f32 %f63, %f13, %f62;
max.f32 %f64, %f63, %f29;
min.f32 %f65, %f64, %f31;
cvt.rni.s32.f32 %r23, %f65;
st.global.u8 [%rd5], %r23;
bra.uni $L__BB0_4;
$L__BB0_3:
shl.b32 %r24, %r1, 3;
suld.b.2d.v4.b8.trap {%rs29, %rs30, %rs31, %rs32}, [%rd6, {%r24, %r4}];
and.b16 %rs33, %rs29, 255;
and.b16 %rs34, %rs30, 255;
and.b16 %rs35, %rs31, 255;
cvt.rn.f32.u16 %f66, %rs33;
cvt.rn.f32.u16 %f67, %rs34;
cvt.rn.f32.u16 %f68, %rs35;
mul.f32 %f69, %f10, %f66;
fma.rn.f32 %f70, %f11, %f67, %f69;
fma.rn.f32 %f71, %f12, %f68, %f70;
add.f32 %f72, %f13, %f71;
mov.f32 %f73, 0f00000000;
max.f32 %f74, %f72, %f73;
mov.f32 %f75, 0f437F0000;
min.f32 %f76, %f74, %f75;
cvt.rni.s32.f32 %r25, %f76;
st.global.u8 [%rd3], %r25;
or.b32 %r26, %r24, 4;
suld.b.2d.v4.b8.trap {%rs36, %rs37, %rs38, %rs39}, [%rd6, {%r26, %r4}];
and.b16 %rs40, %rs36, 255;
and.b16 %rs41, %rs37, 255;
and.b16 %rs42, %rs38, 255;
cvt.rn.f32.u16 %f77, %rs40;
cvt.rn.f32.u16 %f78, %rs41;
cvt.rn.f32.u16 %f79, %rs42;
add.f32 %f80, %f66, %f77;
add.f32 %f81, %f67, %f78;
add.f32 %f82, %f68, %f79;
mul.f32 %f83, %f10, %f77;
fma.rn.f32 %f84, %f11, %f78, %f83;
fma.rn.f32 %f85, %f12, %f79, %f84;
add.f32 %f86, %f13, %f85;
max.f32 %f87, %f86, %f73;
min.f32 %f88, %f87, %f75;
cvt.rni.s32.f32 %r27, %f88;
st.global.u8 [%rd3+1], %r27;
add.s32 %r28, %r4, 1;
suld.b.2d.v4.b8.trap {%rs43, %rs44, %rs45, %rs46}, [%rd6, {%r24, %r28}];
and.b16 %rs47, %rs43, 255;
and.b16 %rs48, %rs44, 255;
and.b16 %rs49, %rs45, 255;
cvt.rn.f32.u16 %f89, %rs47;
cvt.rn.f32.u16 %f90, %rs48;
cvt.rn.f32.u16 %f91, %rs49;
add.f32 %f92, %f80, %f89;
add.f32 %f93, %f81, %f90;
add.f32 %f94, %f82, %f91;
mul.f32 %f95, %f10, %f89;
fma.rn.f32 %f96, %f11, %f90, %f95;
fma.rn.f32 %f97, %f12, %f91, %f96;
add.f32 %f98, %f13, %f97;
max.f32 %f99, %f98, %f73;
min.f32 %f100, %f99, %f75;
cvt.rni.s32.f32 %r29, %f100;
st.global.u8 [%rd4], %r29;
suld.b.2d.v4.b8.trap {%rs50, %rs51, %rs52, %rs53}, [%rd6, {%r26, %r28}];
and.b16 %rs54, %rs50, 255;
and.b16 %rs55, %rs51, 255;
and.b16 %rs56, %rs52, 255;
cvt.rn.f32.u16 %f101, %rs54;
cvt.rn.f32.u16 %f102, %rs55;
cvt.rn.f32.u16 %f103, %rs56;
add.f32 %f127, %f92, %f101;
add.f32 %f128, %f93, %f102;
add.f32 %f129, %f94, %f103;
mul.f32 %f104, %f10, %f101;
fma.rn.f32 %f105, %f11, %f102, %f104;
fma.rn.f32 %f106, %f12, %f103, %f105;
add.f32 %f107, %f13, %f106;
max.f32 %f108, %f107, %f73;
min.f32 %f109, %f108, %f75;
cvt.rni.s32.f32 %r30, %f109;
st.global.u8 [%rd5], %r30;
$L__BB0_4:
mul.lo.s32 %r31, %r7, %r5;
cvt.u64.u32 %rd12, %r31;
and.b32 %r32, %r3, 2147483647;
mul.lo.s32 %r33, %r32, %r7;
cvt.u64.u32 %rd13, %r33;
add.s64 %rd14, %rd13, %rd2;
add.s64 %rd15, %rd14, %rd12;
add.s64 %rd16, %rd1, %rd15;
mul.f32 %f110, %f127, 0f3E800000;
mul.f32 %f111, %f128, 0f3E800000;
mul.f32 %f112, %f15, %f111;
fma.rn.f32 %f113, %f14, %f110, %f112;
mul.f32 %f114, %f129, 0f3E800000;
fma.rn.f32 %f115, %f16, %f114, %f113;
add.f32 %f116, %f17, %f115;
mov.f32 %f117, 0f00000000;
max.f32 %f118, %f116, %f117;
mov.f32 %f119, 0f437F0000;
min.f32 %f120, %f118, %f119;
cvt.rni.s32.f32 %r34, %f120;
st.global.u8 [%rd16], %r34;
mul.f32 %f121, %f19, %f111;
fma.rn.f32 %f122, %f18, %f110, %f121;
fma.rn.f32 %f123, %f20, %f114, %f122;
add.f32 %f124, %f21, %f123;
max.f32 %f125, %f124, %f117;
min.f32 %f126, %f125, %f119;
cvt.rni.s32.f32 %r35, %f126;
st.global.u8 [%rd16+1], %r35;
$L__BB0_5:
ret;
}