const ARGB_TO_NV12_PTX: &[u8] = b"//\n// Generated by NVIDIA NVVM Compiler\n//\n// Compiler Build ID: CL-34097967\n// Cuda compilation tools, release 12.4, V12.4.131\n// Based on NVVM 7.0.1\n//\n\n.version 3.1\n.target sm_30\n.address_size 64\n\n\t// .globl\targb_to_nv12\n// _ZZ9band_hashE4sums has been demoted\n\n.visible .entry argb_to_nv12(\n\t.param .u64 argb_to_nv12_param_0,\n\t.param .u32 argb_to_nv12_param_1,\n\t.param .u64 argb_to_nv12_param_2,\n\t.param .u32 argb_to_nv12_param_3,\n\t.param .u32 argb_to_nv12_param_4,\n\t.param .u32 argb_to_nv12_param_5,\n\t.param .u32 argb_to_nv12_param_6\n)\n{\n\t.reg .pred \t%p<5>;\n\t.reg .b16 \t%rs<25>;\n\t.reg .f32 \t%f<66>;\n\t.reg .b32 \t%r<30>;\n\t.reg .b64 \t%rd<55>;\n\n\n\tld.param.u64 \t%rd1, [argb_to_nv12_param_0];\n\tld.param.u32 \t%r4, [argb_to_nv12_param_1];\n\tld.param.u64 \t%rd2, [argb_to_nv12_param_2];\n\tld.param.u32 \t%r5, [argb_to_nv12_param_3];\n\tld.param.u32 \t%r6, [argb_to_nv12_param_4];\n\tld.param.u32 \t%r7, [argb_to_nv12_param_5];\n\tld.param.u32 \t%r8, [argb_to_nv12_param_6];\n\tmov.u32 \t%r9, %ntid.x;\n\tmov.u32 \t%r10, %ctaid.x;\n\tmov.u32 \t%r11, %tid.x;\n\tmad.lo.s32 \t%r12, %r10, %r9, %r11;\n\tmov.u32 \t%r13, %ntid.y;\n\tmov.u32 \t%r14, %ctaid.y;\n\tmov.u32 \t%r15, %tid.y;\n\tmad.lo.s32 \t%r1, %r14, %r13, %r15;\n\tshl.b32 \t%r2, %r12, 1;\n\tsetp.ge.s32 \t%p1, %r2, %r6;\n\tshl.b32 \t%r3, %r1, 1;\n\tsetp.ge.s32 \t%p2, %r3, %r7;\n\tor.pred \t%p3, %p1, %p2;\n\t@%p3 bra \t$L__BB0_2;\n\n\tcvta.to.global.u64 \t%rd3, %rd1;\n\tcvta.to.global.u64 \t%rd4, %rd2;\n\tcvt.s64.s32 \t%rd5, %r5;\n\tcvt.s64.s32 \t%rd6, %r1;\n\tcvt.s64.s32 \t%rd7, %r7;\n\tadd.s64 \t%rd8, %rd7, %rd6;\n\tmul.lo.s64 \t%rd9, %rd8, %rd5;\n\tcvt.s64.s32 \t%rd10, %r2;\n\tadd.s64 \t%rd11, %rd9, %rd10;\n\tadd.s64 \t%rd12, %rd4, %rd11;\n\tsetp.eq.s32 \t%p4, %r8, 0;\n\tselp.b64 \t%rd13, 0, 2, %p4;\n\tselp.b64 \t%rd14, 2, 0, %p4;\n\tadd.s32 \t%r16, %r6, -1;\n\tmin.s32 \t%r17, %r2, %r16;\n\tcvt.s64.s32 \t%rd15, %r17;\n\tmul.wide.s32 \t%rd16, %r17, 4;\n\tadd.s32 \t%r18, %r7, -1;\n\tmin.s32 \t%r19, %r3, %r18;\n\tmul.wide.s32 \t%rd17, %r19, %r4;\n\tmul.wide.s32 \t%rd18, %r19, %r5;\n\tadd.s64 \t%rd19, %rd16, %rd17;\n\tadd.s64 \t%rd20, %rd19, %rd13;\n\tadd.s64 \t%rd21, %rd3, %rd20;\n\tld.global.u8 \t%rs1, [%rd21];\n\tcvt.rn.f32.u16 \t%f1, %rs1;\n\tadd.s64 \t%rd22, %rd3, %rd19;\n\tld.global.u8 \t%rs3, [%rd22+1];\n\tcvt.rn.f32.u16 \t%f2, %rs3;\n\tadd.s64 \t%rd23, %rd19, %rd14;\n\tadd.s64 \t%rd24, %rd3, %rd23;\n\tld.global.u8 \t%rs5, [%rd24];\n\tcvt.rn.f32.u16 \t%f3, %rs5;\n\tmul.f32 \t%f4, %f3, 0f3E59B3D0;\n\tfma.rn.f32 \t%f5, %f2, 0f3F371759, %f4;\n\tfma.rn.f32 \t%f6, %f1, 0f3D93DD98, %f5;\n\tfma.rn.f32 \t%f7, %f6, 0f3F5BDBDC, 0f41800000;\n\tmov.f32 \t%f8, 0f00000000;\n\tmax.f32 \t%f9, %f7, %f8;\n\tmov.f32 \t%f10, 0f437F0000;\n\tmin.f32 \t%f11, %f9, %f10;\n\tcvt.rni.s32.f32 \t%r20, %f11;\n\tadd.s64 \t%rd25, %rd18, %rd15;\n\tadd.s64 \t%rd26, %rd4, %rd25;\n\tst.global.u8 \t[%rd26], %r20;\n\tadd.s32 \t%r21, %r2, 1;\n\tmin.s32 \t%r22, %r21, %r16;\n\tcvt.s64.s32 \t%rd27, %r22;\n\tmul.wide.s32 \t%rd28, %r22, 4;\n\tadd.s64 \t%rd29, %rd28, %rd17;\n\tadd.s64 \t%rd30, %rd29, %rd13;\n\tadd.s64 \t%rd31, %rd3, %rd30;\n\tld.global.u8 \t%rs7, [%rd31];\n\tcvt.rn.f32.u16 \t%f12, %rs7;\n\tadd.s64 \t%rd32, %rd3, %rd29;\n\tld.global.u8 \t%rs9, [%rd32+1];\n\tcvt.rn.f32.u16 \t%f13, %rs9;\n\tadd.s64 \t%rd33, %rd29, %rd14;\n\tadd.s64 \t%rd34, %rd3, %rd33;\n\tld.global.u8 \t%rs11, [%rd34];\n\tcvt.rn.f32.u16 \t%f14, %rs11;\n\tadd.f32 \t%f15, %f3, %f14;\n\tadd.f32 \t%f16, %f2, %f13;\n\tadd.f32 \t%f17, %f1, %f12;\n\tmul.f32 \t%f18, %f14, 0f3E59B3D0;\n\tfma.rn.f32 \t%f19, %f13, 0f3F371759, %f18;\n\tfma.rn.f32 \t%f20, %f12, 0f3D93DD98, %f19;\n\tfma.rn.f32 \t%f21, %f20, 0f3F5BDBDC, 0f41800000;\n\tmax.f32 \t%f22, %f21, %f8;\n\tmin.f32 \t%f23, %f22, %f10;\n\tcvt.rni.s32.f32 \t%r23, %f23;\n\tadd.s64 \t%rd35, %rd18, %rd27;\n\tadd.s64 \t%rd36, %rd4, %rd35;\n\tst.global.u8 \t[%rd36], %r23;\n\tadd.s32 \t%r24, %r3, 1;\n\tmin.s32 \t%r25, %r24, %r18;\n\tmul.wide.s32 \t%rd37, %r25, %r4;\n\tmul.wide.s32 \t%rd38, %r25, %r5;\n\tadd.s64 \t%rd39, %rd16, %rd37;\n\tadd.s64 \t%rd40, %rd39, %rd13;\n\tadd.s64 \t%rd41, %rd3, %rd40;\n\tld.global.u8 \t%rs13, [%rd41];\n\tcvt.rn.f32.u16 \t%f24, %rs13;\n\tadd.s64 \t%rd42, %rd3, %rd39;\n\tld.global.u8 \t%rs15, [%rd42+1];\n\tcvt.rn.f32.u16 \t%f25, %rs15;\n\tadd.s64 \t%rd43, %rd39, %rd14;\n\tadd.s64 \t%rd44, %rd3, %rd43;\n\tld.global.u8 \t%rs17, [%rd44];\n\tcvt.rn.f32.u16 \t%f26, %rs17;\n\tadd.f32 \t%f27, %f15, %f26;\n\tadd.f32 \t%f28, %f16, %f25;\n\tadd.f32 \t%f29, %f17, %f24;\n\tmul.f32 \t%f30, %f26, 0f3E59B3D0;\n\tfma.rn.f32 \t%f31, %f25, 0f3F371759, %f30;\n\tfma.rn.f32 \t%f32, %f24, 0f3D93DD98, %f31;\n\tfma.rn.f32 \t%f33, %f32, 0f3F5BDBDC, 0f41800000;\n\tmax.f32 \t%f34, %f33, %f8;\n\tmin.f32 \t%f35, %f34, %f10;\n\tcvt.rni.s32.f32 \t%r26, %f35;\n\tadd.s64 \t%rd45, %rd38, %rd15;\n\tadd.s64 \t%rd46, %rd4, %rd45;\n\tst.global.u8 \t[%rd46], %r26;\n\tadd.s64 \t%rd47, %rd28, %rd37;\n\tadd.s64 \t%rd48, %rd47, %rd13;\n\tadd.s64 \t%rd49, %rd3, %rd48;\n\tld.global.u8 \t%rs19, [%rd49];\n\tcvt.rn.f32.u16 \t%f36, %rs19;\n\tadd.s64 \t%rd50, %rd3, %rd47;\n\tld.global.u8 \t%rs21, [%rd50+1];\n\tcvt.rn.f32.u16 \t%f37, %rs21;\n\tadd.s64 \t%rd51, %rd47, %rd14;\n\tadd.s64 \t%rd52, %rd3, %rd51;\n\tld.global.u8 \t%rs23, [%rd52];\n\tcvt.rn.f32.u16 \t%f38, %rs23;\n\tadd.f32 \t%f39, %f27, %f38;\n\tadd.f32 \t%f40, %f28, %f37;\n\tadd.f32 \t%f41, %f29, %f36;\n\tmul.f32 \t%f42, %f38, 0f3E59B3D0;\n\tfma.rn.f32 \t%f43, %f37, 0f3F371759, %f42;\n\tfma.rn.f32 \t%f44, %f36, 0f3D93DD98, %f43;\n\tfma.rn.f32 \t%f45, %f44, 0f3F5BDBDC, 0f41800000;\n\tmax.f32 \t%f46, %f45, %f8;\n\tmin.f32 \t%f47, %f46, %f10;\n\tcvt.rni.s32.f32 \t%r27, %f47;\n\tadd.s64 \t%rd53, %rd38, %rd27;\n\tadd.s64 \t%rd54, %rd4, %rd53;\n\tst.global.u8 \t[%rd54], %r27;\n\tmul.f32 \t%f48, %f39, 0f3E800000;\n\tmul.f32 \t%f49, %f40, 0f3E800000;\n\tmul.f32 \t%f50, %f41, 0f3E800000;\n\tmul.f32 \t%f51, %f49, 0f3F371759;\n\tfma.rn.f32 \t%f52, %f48, 0f3E59B3D0, %f51;\n\tfma.rn.f32 \t%f53, %f50, 0f3D93DD98, %f52;\n\tsub.f32 \t%f54, %f50, %f53;\n\tmul.f32 \t%f55, %f54, 0f3F60E0E1;\n\tdiv.rn.f32 \t%f56, %f55, 0f3FED844D;\n\tadd.f32 \t%f57, %f56, 0f43000000;\n\tmax.f32 \t%f58, %f57, %f8;\n\tmin.f32 \t%f59, %f58, %f10;\n\tcvt.rni.s32.f32 \t%r28, %f59;\n\tst.global.u8 \t[%rd12], %r28;\n\tsub.f32 \t%f60, %f48, %f53;\n\tmul.f32 \t%f61, %f60, 0f3F60E0E1;\n\tdiv.rn.f32 \t%f62, %f61, 0f3FC9930C;\n\tadd.f32 \t%f63, %f62, 0f43000000;\n\tmax.f32 \t%f64, %f63, %f8;\n\tmin.f32 \t%f65, %f64, %f10;\n\tcvt.rni.s32.f32 \t%r29, %f65;\n\tst.global.u8 \t[%rd12+1], %r29;\n\n$L__BB0_2:\n\tret;\n\n}\n\t// .globl\targb_tex_to_nv12\n.visible .entry argb_tex_to_nv12(\n\t.param .u64 argb_tex_to_nv12_param_0,\n\t.param .u64 argb_tex_to_nv12_param_1,\n\t.param .u32 argb_tex_to_nv12_param_2,\n\t.param .u32 argb_tex_to_nv12_param_3,\n\t.param .u32 argb_tex_to_nv12_param_4,\n\t.param .u32 argb_tex_to_nv12_param_5\n)\n{\n\t.reg .pred \t%p<5>;\n\t.reg .b16 \t%rs<49>;\n\t.reg .f32 \t%f<129>;\n\t.reg .b32 \t%r<94>;\n\t.reg .b64 \t%rd<24>;\n\n\n\tld.param.u64 \t%rd6, [argb_tex_to_nv12_param_0];\n\tld.param.u64 \t%rd7, [argb_tex_to_nv12_param_1];\n\tld.param.u32 \t%r5, [argb_tex_to_nv12_param_2];\n\tld.param.u32 \t%r6, [argb_tex_to_nv12_param_3];\n\tld.param.u32 \t%r7, [argb_tex_to_nv12_param_4];\n\tld.param.u32 \t%r8, [argb_tex_to_nv12_param_5];\n\tmov.u32 \t%r9, %ntid.x;\n\tmov.u32 \t%r10, %ctaid.x;\n\tmov.u32 \t%r11, %tid.x;\n\tmad.lo.s32 \t%r12, %r10, %r9, %r11;\n\tmov.u32 \t%r13, %ntid.y;\n\tmov.u32 \t%r14, %ctaid.y;\n\tmov.u32 \t%r15, %tid.y;\n\tmad.lo.s32 \t%r1, %r14, %r13, %r15;\n\tshl.b32 \t%r2, %r12, 1;\n\tsetp.ge.s32 \t%p1, %r2, %r6;\n\tshl.b32 \t%r3, %r1, 1;\n\tsetp.ge.s32 \t%p2, %r3, %r7;\n\tor.pred \t%p3, %p1, %p2;\n\t@%p3 bra \t$L__BB1_5;\n\n\tcvta.to.global.u64 \t%rd8, %rd7;\n\tcvt.s64.s32 \t%rd9, %r5;\n\tcvt.s64.s32 \t%rd10, %r1;\n\tcvt.s64.s32 \t%rd11, %r7;\n\tadd.s64 \t%rd12, %rd11, %rd10;\n\tmul.lo.s64 \t%rd13, %rd12, %rd9;\n\tcvt.s64.s32 \t%rd14, %r2;\n\tadd.s64 \t%rd15, %rd13, %rd14;\n\tadd.s64 \t%rd1, %rd8, %rd15;\n\tadd.s32 \t%r16, %r7, -1;\n\tmin.s32 \t%r17, %r3, %r16;\n\tadd.s32 \t%r18, %r6, -1;\n\tmin.s32 \t%r19, %r2, %r18;\n\tcvt.s64.s32 \t%rd16, %r19;\n\tmul.wide.s32 \t%rd17, %r17, %r5;\n\tadd.s64 \t%rd18, %rd17, %rd16;\n\tadd.s64 \t%rd2, %rd8, %rd18;\n\tadd.s32 \t%r20, %r2, 1;\n\tmin.s32 \t%r21, %r20, %r18;\n\tcvt.s64.s32 \t%rd19, %r21;\n\tadd.s64 \t%rd20, %rd17, %rd19;\n\tadd.s64 \t%rd3, %rd8, %rd20;\n\tadd.s32 \t%r22, %r3, 1;\n\tmin.s32 \t%r4, %r22, %r16;\n\tmul.wide.s32 \t%rd21, %r4, %r5;\n\tadd.s64 \t%rd22, %rd21, %rd16;\n\tadd.s64 \t%rd4, %rd8, %rd22;\n\tadd.s64 \t%rd23, %rd21, %rd19;\n\tadd.s64 \t%rd5, %rd8, %rd23;\n\tsetp.eq.s32 \t%p4, %r8, 0;\n\t@%p4 bra \t$L__BB1_3;\n\n\tcvt.rn.f32.s32 \t%f10, %r17;\n\tcvt.rn.f32.s32 \t%f11, %r19;\n\ttex.2d.v4.u32.f32 \t{%r32, %r33, %r34, %r35}, [%rd6, {%f11, %f10}];\n\tcvt.u16.u32 \t%rs1, %r32;\n\tand.b16 \t%rs2, %rs1, 255;\n\tcvt.u16.u32 \t%rs3, %r33;\n\tand.b16 \t%rs4, %rs3, 255;\n\tcvt.u16.u32 \t%rs5, %r34;\n\tand.b16 \t%rs6, %rs5, 255;\n\tcvt.rn.f32.u16 \t%f12, %rs6;\n\tcvt.rn.f32.u16 \t%f13, %rs4;\n\tcvt.rn.f32.u16 \t%f14, %rs2;\n\tmul.f32 \t%f15, %f14, 0f3E59B3D0;\n\tfma.rn.f32 \t%f16, %f13, 0f3F371759, %f15;\n\tfma.rn.f32 \t%f17, %f12, 0f3D93DD98, %f16;\n\tfma.rn.f32 \t%f18, %f17, 0f3F5BDBDC, 0f41800000;\n\tmov.f32 \t%f19, 0f00000000;\n\tmax.f32 \t%f20, %f18, %f19;\n\tmov.f32 \t%f21, 0f437F0000;\n\tmin.f32 \t%f22, %f20, %f21;\n\tcvt.rni.s32.f32 \t%r36, %f22;\n\tst.global.u8 \t[%rd2], %r36;\n\tor.b32 \t%r37, %r2, 1;\n\tmin.s32 \t%r38, %r37, %r18;\n\tcvt.rn.f32.s32 \t%f23, %r38;\n\ttex.2d.v4.u32.f32 \t{%r39, %r40, %r41, %r42}, [%rd6, {%f23, %f10}];\n\tcvt.u16.u32 \t%rs7, %r39;\n\tand.b16 \t%rs8, %rs7, 255;\n\tcvt.u16.u32 \t%rs9, %r40;\n\tand.b16 \t%rs10, %rs9, 255;\n\tcvt.u16.u32 \t%rs11, %r41;\n\tand.b16 \t%rs12, %rs11, 255;\n\tcvt.rn.f32.u16 \t%f24, %rs12;\n\tcvt.rn.f32.u16 \t%f25, %rs10;\n\tcvt.rn.f32.u16 \t%f26, %rs8;\n\tadd.f32 \t%f27, %f14, %f26;\n\tadd.f32 \t%f28, %f13, %f25;\n\tadd.f32 \t%f29, %f12, %f24;\n\tmul.f32 \t%f30, %f26, 0f3E59B3D0;\n\tfma.rn.f32 \t%f31, %f25, 0f3F371759, %f30;\n\tfma.rn.f32 \t%f32, %f24, 0f3D93DD98, %f31;\n\tfma.rn.f32 \t%f33, %f32, 0f3F5BDBDC, 0f41800000;\n\tmax.f32 \t%f34, %f33, %f19;\n\tmin.f32 \t%f35, %f34, %f21;\n\tcvt.rni.s32.f32 \t%r43, %f35;\n\tst.global.u8 \t[%rd3], %r43;\n\tcvt.rn.f32.s32 \t%f36, %r4;\n\ttex.2d.v4.u32.f32 \t{%r46, %r47, %r48, %r49}, [%rd6, {%f11, %f36}];\n\tcvt.u16.u32 \t%rs13, %r46;\n\tand.b16 \t%rs14, %rs13, 255;\n\tcvt.u16.u32 \t%rs15, %r47;\n\tand.b16 \t%rs16, %rs15, 255;\n\tcvt.u16.u32 \t%rs17, %r48;\n\tand.b16 \t%rs18, %rs17, 255;\n\tcvt.rn.f32.u16 \t%f37, %rs18;\n\tcvt.rn.f32.u16 \t%f38, %rs16;\n\tcvt.rn.f32.u16 \t%f39, %rs14;\n\tadd.f32 \t%f40, %f27, %f39;\n\tadd.f32 \t%f41, %f28, %f38;\n\tadd.f32 \t%f42, %f29, %f37;\n\tmul.f32 \t%f43, %f39, 0f3E59B3D0;\n\tfma.rn.f32 \t%f44, %f38, 0f3F371759, %f43;\n\tfma.rn.f32 \t%f45, %f37, 0f3D93DD98, %f44;\n\tfma.rn.f32 \t%f46, %f45, 0f3F5BDBDC, 0f41800000;\n\tmax.f32 \t%f47, %f46, %f19;\n\tmin.f32 \t%f48, %f47, %f21;\n\tcvt.rni.s32.f32 \t%r50, %f48;\n\tst.global.u8 \t[%rd4], %r50;\n\ttex.2d.v4.u32.f32 \t{%r51, %r52, %r53, %r54}, [%rd6, {%f23, %f36}];\n\tcvt.u16.u32 \t%rs19, %r51;\n\tand.b16 \t%rs20, %rs19, 255;\n\tcvt.u16.u32 \t%rs21, %r52;\n\tand.b16 \t%rs22, %rs21, 255;\n\tcvt.u16.u32 \t%rs23, %r53;\n\tand.b16 \t%rs24, %rs23, 255;\n\tcvt.rn.f32.u16 \t%f49, %rs24;\n\tcvt.rn.f32.u16 \t%f50, %rs22;\n\tcvt.rn.f32.u16 \t%f51, %rs20;\n\tadd.f32 \t%f126, %f40, %f51;\n\tadd.f32 \t%f127, %f41, %f50;\n\tadd.f32 \t%f128, %f42, %f49;\n\tmul.f32 \t%f52, %f51, 0f3E59B3D0;\n\tfma.rn.f32 \t%f53, %f50, 0f3F371759, %f52;\n\tfma.rn.f32 \t%f54, %f49, 0f3D93DD98, %f53;\n\tfma.rn.f32 \t%f55, %f54, 0f3F5BDBDC, 0f41800000;\n\tmax.f32 \t%f56, %f55, %f19;\n\tmin.f32 \t%f57, %f56, %f21;\n\tcvt.rni.s32.f32 \t%r55, %f57;\n\tst.global.u8 \t[%rd5], %r55;\n\tbra.uni \t$L__BB1_4;\n\n$L__BB1_3:\n\tcvt.rn.f32.s32 \t%f58, %r17;\n\tcvt.rn.f32.s32 \t%f59, %r19;\n\ttex.2d.v4.u32.f32 \t{%r70, %r71, %r72, %r73}, [%rd6, {%f59, %f58}];\n\tcvt.u16.u32 \t%rs25, %r70;\n\tand.b16 \t%rs26, %rs25, 255;\n\tcvt.u16.u32 \t%rs27, %r71;\n\tand.b16 \t%rs28, %rs27, 255;\n\tcvt.u16.u32 \t%rs29, %r72;\n\tand.b16 \t%rs30, %rs29, 255;\n\tcvt.rn.f32.u16 \t%f60, %rs26;\n\tcvt.rn.f32.u16 \t%f61, %rs28;\n\tcvt.rn.f32.u16 \t%f62, %rs30;\n\tmul.f32 \t%f63, %f62, 0f3E59B3D0;\n\tfma.rn.f32 \t%f64, %f61, 0f3F371759, %f63;\n\tfma.rn.f32 \t%f65, %f60, 0f3D93DD98, %f64;\n\tfma.rn.f32 \t%f66, %f65, 0f3F5BDBDC, 0f41800000;\n\tmov.f32 \t%f67, 0f00000000;\n\tmax.f32 \t%f68, %f66, %f67;\n\tmov.f32 \t%f69, 0f437F0000;\n\tmin.f32 \t%f70, %f68, %f69;\n\tcvt.rni.s32.f32 \t%r74, %f70;\n\tst.global.u8 \t[%rd2], %r74;\n\tor.b32 \t%r75, %r2, 1;\n\tmin.s32 \t%r76, %r75, %r18;\n\tcvt.rn.f32.s32 \t%f71, %r76;\n\ttex.2d.v4.u32.f32 \t{%r77, %r78, %r79, %r80}, [%rd6, {%f71, %f58}];\n\tcvt.u16.u32 \t%rs31, %r77;\n\tand.b16 \t%rs32, %rs31, 255;\n\tcvt.u16.u32 \t%rs33, %r78;\n\tand.b16 \t%rs34, %rs33, 255;\n\tcvt.u16.u32 \t%rs35, %r79;\n\tand.b16 \t%rs36, %rs35, 255;\n\tcvt.rn.f32.u16 \t%f72, %rs32;\n\tcvt.rn.f32.u16 \t%f73, %rs34;\n\tcvt.rn.f32.u16 \t%f74, %rs36;\n\tadd.f32 \t%f75, %f62, %f74;\n\tadd.f32 \t%f76, %f61, %f73;\n\tadd.f32 \t%f77, %f60, %f72;\n\tmul.f32 \t%f78, %f74, 0f3E59B3D0;\n\tfma.rn.f32 \t%f79, %f73, 0f3F371759, %f78;\n\tfma.rn.f32 \t%f80, %f72, 0f3D93DD98, %f79;\n\tfma.rn.f32 \t%f81, %f80, 0f3F5BDBDC, 0f41800000;\n\tmax.f32 \t%f82, %f81, %f67;\n\tmin.f32 \t%f83, %f82, %f69;\n\tcvt.rni.s32.f32 \t%r81, %f83;\n\tst.global.u8 \t[%rd3], %r81;\n\tcvt.rn.f32.s32 \t%f84, %r4;\n\ttex.2d.v4.u32.f32 \t{%r82, %r83, %r84, %r85}, [%rd6, {%f59, %f84}];\n\tcvt.u16.u32 \t%rs37, %r82;\n\tand.b16 \t%rs38, %rs37, 255;\n\tcvt.u16.u32 \t%rs39, %r83;\n\tand.b16 \t%rs40, %rs39, 255;\n\tcvt.u16.u32 \t%rs41, %r84;\n\tand.b16 \t%rs42, %rs41, 255;\n\tcvt.rn.f32.u16 \t%f85, %rs38;\n\tcvt.rn.f32.u16 \t%f86, %rs40;\n\tcvt.rn.f32.u16 \t%f87, %rs42;\n\tadd.f32 \t%f88, %f75, %f87;\n\tadd.f32 \t%f89, %f76, %f86;\n\tadd.f32 \t%f90, %f77, %f85;\n\tmul.f32 \t%f91, %f87, 0f3E59B3D0;\n\tfma.rn.f32 \t%f92, %f86, 0f3F371759, %f91;\n\tfma.rn.f32 \t%f93, %f85, 0f3D93DD98, %f92;\n\tfma.rn.f32 \t%f94, %f93, 0f3F5BDBDC, 0f41800000;\n\tmax.f32 \t%f95, %f94, %f67;\n\tmin.f32 \t%f96, %f95, %f69;\n\tcvt.rni.s32.f32 \t%r86, %f96;\n\tst.global.u8 \t[%rd4], %r86;\n\ttex.2d.v4.u32.f32 \t{%r87, %r88, %r89, %r90}, [%rd6, {%f71, %f84}];\n\tcvt.u16.u32 \t%rs43, %r87;\n\tand.b16 \t%rs44, %rs43, 255;\n\tcvt.u16.u32 \t%rs45, %r88;\n\tand.b16 \t%rs46, %rs45, 255;\n\tcvt.u16.u32 \t%rs47, %r89;\n\tand.b16 \t%rs48, %rs47, 255;\n\tcvt.rn.f32.u16 \t%f97, %rs44;\n\tcvt.rn.f32.u16 \t%f98, %rs46;\n\tcvt.rn.f32.u16 \t%f99, %rs48;\n\tadd.f32 \t%f126, %f88, %f99;\n\tadd.f32 \t%f127, %f89, %f98;\n\tadd.f32 \t%f128, %f90, %f97;\n\tmul.f32 \t%f100, %f99, 0f3E59B3D0;\n\tfma.rn.f32 \t%f101, %f98, 0f3F371759, %f100;\n\tfma.rn.f32 \t%f102, %f97, 0f3D93DD98, %f101;\n\tfma.rn.f32 \t%f103, %f102, 0f3F5BDBDC, 0f41800000;\n\tmax.f32 \t%f104, %f103, %f67;\n\tmin.f32 \t%f105, %f104, %f69;\n\tcvt.rni.s32.f32 \t%r91, %f105;\n\tst.global.u8 \t[%rd5], %r91;\n\n$L__BB1_4:\n\tmul.f32 \t%f106, %f126, 0f3E800000;\n\tmul.f32 \t%f107, %f127, 0f3E800000;\n\tmul.f32 \t%f108, %f107, 0f3F371759;\n\tfma.rn.f32 \t%f109, %f106, 0f3E59B3D0, %f108;\n\tmul.f32 \t%f110, %f128, 0f3E800000;\n\tfma.rn.f32 \t%f111, %f110, 0f3D93DD98, %f109;\n\tsub.f32 \t%f112, %f110, %f111;\n\tmul.f32 \t%f113, %f112, 0f3F60E0E1;\n\tdiv.rn.f32 \t%f114, %f113, 0f3FED844D;\n\tadd.f32 \t%f115, %f114, 0f43000000;\n\tmov.f32 \t%f116, 0f00000000;\n\tmax.f32 \t%f117, %f115, %f116;\n\tmov.f32 \t%f118, 0f437F0000;\n\tmin.f32 \t%f119, %f117, %f118;\n\tcvt.rni.s32.f32 \t%r92, %f119;\n\tst.global.u8 \t[%rd1], %r92;\n\tsub.f32 \t%f120, %f106, %f111;\n\tmul.f32 \t%f121, %f120, 0f3F60E0E1;\n\tdiv.rn.f32 \t%f122, %f121, 0f3FC9930C;\n\tadd.f32 \t%f123, %f122, 0f43000000;\n\tmax.f32 \t%f124, %f123, %f116;\n\tmin.f32 \t%f125, %f124, %f118;\n\tcvt.rni.s32.f32 \t%r93, %f125;\n\tst.global.u8 \t[%rd1+1], %r93;\n\n$L__BB1_5:\n\tret;\n\n}\n\t// .globl\targb_to_p010\n.visible .entry argb_to_p010(\n\t.param .u64 argb_to_p010_param_0,\n\t.param .u32 argb_to_p010_param_1,\n\t.param .u64 argb_to_p010_param_2,\n\t.param .u32 argb_to_p010_param_3,\n\t.param .u32 argb_to_p010_param_4,\n\t.param .u32 argb_to_p010_param_5,\n\t.param .u32 argb_to_p010_param_6\n)\n{\n\t.reg .pred \t%p<5>;\n\t.reg .b16 \t%rs<37>;\n\t.reg .f32 \t%f<66>;\n\t.reg .b32 \t%r<30>;\n\t.reg .b64 \t%rd<53>;\n\n\n\tld.param.u64 \t%rd1, [argb_to_p010_param_0];\n\tld.param.u32 \t%r4, [argb_to_p010_param_1];\n\tld.param.u64 \t%rd2, [argb_to_p010_param_2];\n\tld.param.u32 \t%r5, [argb_to_p010_param_3];\n\tld.param.u32 \t%r6, [argb_to_p010_param_4];\n\tld.param.u32 \t%r7, [argb_to_p010_param_5];\n\tld.param.u32 \t%r8, [argb_to_p010_param_6];\n\tmov.u32 \t%r9, %ntid.x;\n\tmov.u32 \t%r10, %ctaid.x;\n\tmov.u32 \t%r11, %tid.x;\n\tmad.lo.s32 \t%r12, %r10, %r9, %r11;\n\tmov.u32 \t%r13, %ntid.y;\n\tmov.u32 \t%r14, %ctaid.y;\n\tmov.u32 \t%r15, %tid.y;\n\tmad.lo.s32 \t%r1, %r14, %r13, %r15;\n\tshl.b32 \t%r2, %r12, 1;\n\tsetp.ge.s32 \t%p1, %r2, %r6;\n\tshl.b32 \t%r3, %r1, 1;\n\tsetp.ge.s32 \t%p2, %r3, %r7;\n\tor.pred \t%p3, %p1, %p2;\n\t@%p3 bra \t$L__BB2_2;\n\n\tcvta.to.global.u64 \t%rd3, %rd1;\n\tcvta.to.global.u64 \t%rd4, %rd2;\n\tcvt.s64.s32 \t%rd5, %r5;\n\tcvt.s64.s32 \t%rd6, %r1;\n\tcvt.s64.s32 \t%rd7, %r7;\n\tadd.s64 \t%rd8, %rd7, %rd6;\n\tmul.lo.s64 \t%rd9, %rd8, %rd5;\n\tadd.s64 \t%rd10, %rd4, %rd9;\n\tmul.wide.s32 \t%rd11, %r2, 2;\n\tadd.s64 \t%rd12, %rd10, %rd11;\n\tsetp.eq.s32 \t%p4, %r8, 0;\n\tselp.b64 \t%rd13, 0, 2, %p4;\n\tselp.b64 \t%rd14, 2, 0, %p4;\n\tadd.s32 \t%r16, %r6, -1;\n\tmin.s32 \t%r17, %r2, %r16;\n\tmul.wide.s32 \t%rd15, %r17, 4;\n\tadd.s32 \t%r18, %r7, -1;\n\tmin.s32 \t%r19, %r3, %r18;\n\tmul.wide.s32 \t%rd16, %r19, %r5;\n\tadd.s64 \t%rd17, %rd4, %rd16;\n\tmul.wide.s32 \t%rd18, %r19, %r4;\n\tadd.s64 \t%rd19, %rd15, %rd18;\n\tadd.s64 \t%rd20, %rd19, %rd13;\n\tadd.s64 \t%rd21, %rd3, %rd20;\n\tld.global.u8 \t%rs1, [%rd21];\n\tcvt.rn.f32.u16 \t%f1, %rs1;\n\tadd.s64 \t%rd22, %rd3, %rd19;\n\tld.global.u8 \t%rs3, [%rd22+1];\n\tcvt.rn.f32.u16 \t%f2, %rs3;\n\tadd.s64 \t%rd23, %rd19, %rd14;\n\tadd.s64 \t%rd24, %rd3, %rd23;\n\tld.global.u8 \t%rs5, [%rd24];\n\tcvt.rn.f32.u16 \t%f3, %rs5;\n\tmul.f32 \t%f4, %f3, 0f3E59B3D0;\n\tfma.rn.f32 \t%f5, %f2, 0f3F371759, %f4;\n\tfma.rn.f32 \t%f6, %f1, 0f3D93DD98, %f5;\n\tfma.rn.f32 \t%f7, %f6, 0f405BDBDC, 0f42800000;\n\tmov.f32 \t%f8, 0f00000000;\n\tmax.f32 \t%f9, %f7, %f8;\n\tmov.f32 \t%f10, 0f447FC000;\n\tmin.f32 \t%f11, %f9, %f10;\n\tcvt.rni.s32.f32 \t%r20, %f11;\n\tcvt.u16.u32 \t%rs7, %r20;\n\tshl.b16 \t%rs8, %rs7, 6;\n\tmul.wide.s32 \t%rd25, %r17, 2;\n\tadd.s64 \t%rd26, %rd17, %rd25;\n\tst.global.u16 \t[%rd26], %rs8;\n\tadd.s32 \t%r21, %r2, 1;\n\tmin.s32 \t%r22, %r21, %r16;\n\tmul.wide.s32 \t%rd27, %r22, 4;\n\tadd.s64 \t%rd28, %rd27, %rd18;\n\tadd.s64 \t%rd29, %rd28, %rd13;\n\tadd.s64 \t%rd30, %rd3, %rd29;\n\tld.global.u8 \t%rs9, [%rd30];\n\tcvt.rn.f32.u16 \t%f12, %rs9;\n\tadd.s64 \t%rd31, %rd3, %rd28;\n\tld.global.u8 \t%rs11, [%rd31+1];\n\tcvt.rn.f32.u16 \t%f13, %rs11;\n\tadd.s64 \t%rd32, %rd28, %rd14;\n\tadd.s64 \t%rd33, %rd3, %rd32;\n\tld.global.u8 \t%rs13, [%rd33];\n\tcvt.rn.f32.u16 \t%f14, %rs13;\n\tadd.f32 \t%f15, %f3, %f14;\n\tadd.f32 \t%f16, %f2, %f13;\n\tadd.f32 \t%f17, %f1, %f12;\n\tmul.f32 \t%f18, %f14, 0f3E59B3D0;\n\tfma.rn.f32 \t%f19, %f13, 0f3F371759, %f18;\n\tfma.rn.f32 \t%f20, %f12, 0f3D93DD98, %f19;\n\tfma.rn.f32 \t%f21, %f20, 0f405BDBDC, 0f42800000;\n\tmax.f32 \t%f22, %f21, %f8;\n\tmin.f32 \t%f23, %f22, %f10;\n\tcvt.rni.s32.f32 \t%r23, %f23;\n\tcvt.u16.u32 \t%rs15, %r23;\n\tshl.b16 \t%rs16, %rs15, 6;\n\tmul.wide.s32 \t%rd34, %r22, 2;\n\tadd.s64 \t%rd35, %rd17, %rd34;\n\tst.global.u16 \t[%rd35], %rs16;\n\tadd.s32 \t%r24, %r3, 1;\n\tmin.s32 \t%r25, %r24, %r18;\n\tmul.wide.s32 \t%rd36, %r25, %r5;\n\tadd.s64 \t%rd37, %rd4, %rd36;\n\tmul.wide.s32 \t%rd38, %r25, %r4;\n\tadd.s64 \t%rd39, %rd15, %rd38;\n\tadd.s64 \t%rd40, %rd39, %rd13;\n\tadd.s64 \t%rd41, %rd3, %rd40;\n\tld.global.u8 \t%rs17, [%rd41];\n\tcvt.rn.f32.u16 \t%f24, %rs17;\n\tadd.s64 \t%rd42, %rd3, %rd39;\n\tld.global.u8 \t%rs19, [%rd42+1];\n\tcvt.rn.f32.u16 \t%f25, %rs19;\n\tadd.s64 \t%rd43, %rd39, %rd14;\n\tadd.s64 \t%rd44, %rd3, %rd43;\n\tld.global.u8 \t%rs21, [%rd44];\n\tcvt.rn.f32.u16 \t%f26, %rs21;\n\tadd.f32 \t%f27, %f15, %f26;\n\tadd.f32 \t%f28, %f16, %f25;\n\tadd.f32 \t%f29, %f17, %f24;\n\tmul.f32 \t%f30, %f26, 0f3E59B3D0;\n\tfma.rn.f32 \t%f31, %f25, 0f3F371759, %f30;\n\tfma.rn.f32 \t%f32, %f24, 0f3D93DD98, %f31;\n\tfma.rn.f32 \t%f33, %f32, 0f405BDBDC, 0f42800000;\n\tmax.f32 \t%f34, %f33, %f8;\n\tmin.f32 \t%f35, %f34, %f10;\n\tcvt.rni.s32.f32 \t%r26, %f35;\n\tcvt.u16.u32 \t%rs23, %r26;\n\tshl.b16 \t%rs24, %rs23, 6;\n\tadd.s64 \t%rd45, %rd37, %rd25;\n\tst.global.u16 \t[%rd45], %rs24;\n\tadd.s64 \t%rd46, %rd27, %rd38;\n\tadd.s64 \t%rd47, %rd46, %rd13;\n\tadd.s64 \t%rd48, %rd3, %rd47;\n\tld.global.u8 \t%rs25, [%rd48];\n\tcvt.rn.f32.u16 \t%f36, %rs25;\n\tadd.s64 \t%rd49, %rd3, %rd46;\n\tld.global.u8 \t%rs27, [%rd49+1];\n\tcvt.rn.f32.u16 \t%f37, %rs27;\n\tadd.s64 \t%rd50, %rd46, %rd14;\n\tadd.s64 \t%rd51, %rd3, %rd50;\n\tld.global.u8 \t%rs29, [%rd51];\n\tcvt.rn.f32.u16 \t%f38, %rs29;\n\tadd.f32 \t%f39, %f27, %f38;\n\tadd.f32 \t%f40, %f28, %f37;\n\tadd.f32 \t%f41, %f29, %f36;\n\tmul.f32 \t%f42, %f38, 0f3E59B3D0;\n\tfma.rn.f32 \t%f43, %f37, 0f3F371759, %f42;\n\tfma.rn.f32 \t%f44, %f36, 0f3D93DD98, %f43;\n\tfma.rn.f32 \t%f45, %f44, 0f405BDBDC, 0f42800000;\n\tmax.f32 \t%f46, %f45, %f8;\n\tmin.f32 \t%f47, %f46, %f10;\n\tcvt.rni.s32.f32 \t%r27, %f47;\n\tcvt.u16.u32 \t%rs31, %r27;\n\tshl.b16 \t%rs32, %rs31, 6;\n\tadd.s64 \t%rd52, %rd37, %rd34;\n\tst.global.u16 \t[%rd52], %rs32;\n\tmul.f32 \t%f48, %f39, 0f3E800000;\n\tmul.f32 \t%f49, %f40, 0f3E800000;\n\tmul.f32 \t%f50, %f41, 0f3E800000;\n\tmul.f32 \t%f51, %f49, 0f3F371759;\n\tfma.rn.f32 \t%f52, %f48, 0f3E59B3D0, %f51;\n\tfma.rn.f32 \t%f53, %f50, 0f3D93DD98, %f52;\n\tsub.f32 \t%f54, %f50, %f53;\n\tmul.f32 \t%f55, %f54, 0f4060E0E1;\n\tdiv.rn.f32 \t%f56, %f55, 0f3FED844D;\n\tadd.f32 \t%f57, %f56, 0f44000000;\n\tmax.f32 \t%f58, %f57, %f8;\n\tmin.f32 \t%f59, %f58, %f10;\n\tcvt.rni.s32.f32 \t%r28, %f59;\n\tcvt.u16.u32 \t%rs33, %r28;\n\tshl.b16 \t%rs34, %rs33, 6;\n\tst.global.u16 \t[%rd12], %rs34;\n\tsub.f32 \t%f60, %f48, %f53;\n\tmul.f32 \t%f61, %f60, 0f4060E0E1;\n\tdiv.rn.f32 \t%f62, %f61, 0f3FC9930C;\n\tadd.f32 \t%f63, %f62, 0f44000000;\n\tmax.f32 \t%f64, %f63, %f8;\n\tmin.f32 \t%f65, %f64, %f10;\n\tcvt.rni.s32.f32 \t%r29, %f65;\n\tcvt.u16.u32 \t%rs35, %r29;\n\tshl.b16 \t%rs36, %rs35, 6;\n\tst.global.u16 \t[%rd12+2], %rs36;\n\n$L__BB2_2:\n\tret;\n\n}\n\t// .globl\targb_tex_to_p010\n.visible .entry argb_tex_to_p010(\n\t.param .u64 argb_tex_to_p010_param_0,\n\t.param .u64 argb_tex_to_p010_param_1,\n\t.param .u32 argb_tex_to_p010_param_2,\n\t.param .u32 argb_tex_to_p010_param_3,\n\t.param .u32 argb_tex_to_p010_param_4,\n\t.param .u32 argb_tex_to_p010_param_5\n)\n{\n\t.reg .pred \t%p<5>;\n\t.reg .b16 \t%rs<69>;\n\t.reg .f32 \t%f<129>;\n\t.reg .b32 \t%r<94>;\n\t.reg .b64 \t%rd<22>;\n\n\n\tld.param.u64 \t%rd6, [argb_tex_to_p010_param_0];\n\tld.param.u64 \t%rd7, [argb_tex_to_p010_param_1];\n\tld.param.u32 \t%r5, [argb_tex_to_p010_param_2];\n\tld.param.u32 \t%r6, [argb_tex_to_p010_param_3];\n\tld.param.u32 \t%r7, [argb_tex_to_p010_param_4];\n\tld.param.u32 \t%r8, [argb_tex_to_p010_param_5];\n\tmov.u32 \t%r9, %ntid.x;\n\tmov.u32 \t%r10, %ctaid.x;\n\tmov.u32 \t%r11, %tid.x;\n\tmad.lo.s32 \t%r12, %r10, %r9, %r11;\n\tmov.u32 \t%r13, %ntid.y;\n\tmov.u32 \t%r14, %ctaid.y;\n\tmov.u32 \t%r15, %tid.y;\n\tmad.lo.s32 \t%r1, %r14, %r13, %r15;\n\tshl.b32 \t%r2, %r12, 1;\n\tsetp.ge.s32 \t%p1, %r2, %r6;\n\tshl.b32 \t%r3, %r1, 1;\n\tsetp.ge.s32 \t%p2, %r3, %r7;\n\tor.pred \t%p3, %p1, %p2;\n\t@%p3 bra \t$L__BB3_5;\n\n\tcvta.to.global.u64 \t%rd8, %rd7;\n\tcvt.s64.s32 \t%rd9, %r5;\n\tcvt.s64.s32 \t%rd10, %r1;\n\tcvt.s64.s32 \t%rd11, %r7;\n\tadd.s64 \t%rd12, %rd11, %rd10;\n\tmul.lo.s64 \t%rd13, %rd12, %rd9;\n\tadd.s64 \t%rd14, %rd8, %rd13;\n\tmul.wide.s32 \t%rd15, %r2, 2;\n\tadd.s64 \t%rd1, %rd14, %rd15;\n\tadd.s32 \t%r16, %r7, -1;\n\tmin.s32 \t%r17, %r3, %r16;\n\tmul.wide.s32 \t%rd16, %r17, %r5;\n\tadd.s64 \t%rd17, %rd8, %rd16;\n\tadd.s32 \t%r18, %r6, -1;\n\tmin.s32 \t%r19, %r2, %r18;\n\tmul.wide.s32 \t%rd18, %r19, 2;\n\tadd.s64 \t%rd2, %rd17, %rd18;\n\tadd.s32 \t%r20, %r2, 1;\n\tmin.s32 \t%r21, %r20, %r18;\n\tmul.wide.s32 \t%rd19, %r21, 2;\n\tadd.s64 \t%rd3, %rd17, %rd19;\n\tadd.s32 \t%r22, %r3, 1;\n\tmin.s32 \t%r4, %r22, %r16;\n\tmul.wide.s32 \t%rd20, %r4, %r5;\n\tadd.s64 \t%rd21, %rd8, %rd20;\n\tadd.s64 \t%rd4, %rd21, %rd18;\n\tadd.s64 \t%rd5, %rd21, %rd19;\n\tsetp.eq.s32 \t%p4, %r8, 0;\n\t@%p4 bra \t$L__BB3_3;\n\n\tcvt.rn.f32.s32 \t%f10, %r17;\n\tcvt.rn.f32.s32 \t%f11, %r19;\n\ttex.2d.v4.u32.f32 \t{%r32, %r33, %r34, %r35}, [%rd6, {%f11, %f10}];\n\tcvt.u16.u32 \t%rs1, %r32;\n\tand.b16 \t%rs2, %rs1, 255;\n\tcvt.u16.u32 \t%rs3, %r33;\n\tand.b16 \t%rs4, %rs3, 255;\n\tcvt.u16.u32 \t%rs5, %r34;\n\tand.b16 \t%rs6, %rs5, 255;\n\tcvt.rn.f32.u16 \t%f12, %rs6;\n\tcvt.rn.f32.u16 \t%f13, %rs4;\n\tcvt.rn.f32.u16 \t%f14, %rs2;\n\tmul.f32 \t%f15, %f14, 0f3E59B3D0;\n\tfma.rn.f32 \t%f16, %f13, 0f3F371759, %f15;\n\tfma.rn.f32 \t%f17, %f12, 0f3D93DD98, %f16;\n\tfma.rn.f32 \t%f18, %f17, 0f405BDBDC, 0f42800000;\n\tmov.f32 \t%f19, 0f00000000;\n\tmax.f32 \t%f20, %f18, %f19;\n\tmov.f32 \t%f21, 0f447FC000;\n\tmin.f32 \t%f22, %f20, %f21;\n\tcvt.rni.s32.f32 \t%r36, %f22;\n\tcvt.u16.u32 \t%rs7, %r36;\n\tshl.b16 \t%rs8, %rs7, 6;\n\tst.global.u16 \t[%rd2], %rs8;\n\tor.b32 \t%r37, %r2, 1;\n\tmin.s32 \t%r38, %r37, %r18;\n\tcvt.rn.f32.s32 \t%f23, %r38;\n\ttex.2d.v4.u32.f32 \t{%r39, %r40, %r41, %r42}, [%rd6, {%f23, %f10}];\n\tcvt.u16.u32 \t%rs9, %r39;\n\tand.b16 \t%rs10, %rs9, 255;\n\tcvt.u16.u32 \t%rs11, %r40;\n\tand.b16 \t%rs12, %rs11, 255;\n\tcvt.u16.u32 \t%rs13, %r41;\n\tand.b16 \t%rs14, %rs13, 255;\n\tcvt.rn.f32.u16 \t%f24, %rs14;\n\tcvt.rn.f32.u16 \t%f25, %rs12;\n\tcvt.rn.f32.u16 \t%f26, %rs10;\n\tadd.f32 \t%f27, %f14, %f26;\n\tadd.f32 \t%f28, %f13, %f25;\n\tadd.f32 \t%f29, %f12, %f24;\n\tmul.f32 \t%f30, %f26, 0f3E59B3D0;\n\tfma.rn.f32 \t%f31, %f25, 0f3F371759, %f30;\n\tfma.rn.f32 \t%f32, %f24, 0f3D93DD98, %f31;\n\tfma.rn.f32 \t%f33, %f32, 0f405BDBDC, 0f42800000;\n\tmax.f32 \t%f34, %f33, %f19;\n\tmin.f32 \t%f35, %f34, %f21;\n\tcvt.rni.s32.f32 \t%r43, %f35;\n\tcvt.u16.u32 \t%rs15, %r43;\n\tshl.b16 \t%rs16, %rs15, 6;\n\tst.global.u16 \t[%rd3], %rs16;\n\tcvt.rn.f32.s32 \t%f36, %r4;\n\ttex.2d.v4.u32.f32 \t{%r46, %r47, %r48, %r49}, [%rd6, {%f11, %f36}];\n\tcvt.u16.u32 \t%rs17, %r46;\n\tand.b16 \t%rs18, %rs17, 255;\n\tcvt.u16.u32 \t%rs19, %r47;\n\tand.b16 \t%rs20, %rs19, 255;\n\tcvt.u16.u32 \t%rs21, %r48;\n\tand.b16 \t%rs22, %rs21, 255;\n\tcvt.rn.f32.u16 \t%f37, %rs22;\n\tcvt.rn.f32.u16 \t%f38, %rs20;\n\tcvt.rn.f32.u16 \t%f39, %rs18;\n\tadd.f32 \t%f40, %f27, %f39;\n\tadd.f32 \t%f41, %f28, %f38;\n\tadd.f32 \t%f42, %f29, %f37;\n\tmul.f32 \t%f43, %f39, 0f3E59B3D0;\n\tfma.rn.f32 \t%f44, %f38, 0f3F371759, %f43;\n\tfma.rn.f32 \t%f45, %f37, 0f3D93DD98, %f44;\n\tfma.rn.f32 \t%f46, %f45, 0f405BDBDC, 0f42800000;\n\tmax.f32 \t%f47, %f46, %f19;\n\tmin.f32 \t%f48, %f47, %f21;\n\tcvt.rni.s32.f32 \t%r50, %f48;\n\tcvt.u16.u32 \t%rs23, %r50;\n\tshl.b16 \t%rs24, %rs23, 6;\n\tst.global.u16 \t[%rd4], %rs24;\n\ttex.2d.v4.u32.f32 \t{%r51, %r52, %r53, %r54}, [%rd6, {%f23, %f36}];\n\tcvt.u16.u32 \t%rs25, %r51;\n\tand.b16 \t%rs26, %rs25, 255;\n\tcvt.u16.u32 \t%rs27, %r52;\n\tand.b16 \t%rs28, %rs27, 255;\n\tcvt.u16.u32 \t%rs29, %r53;\n\tand.b16 \t%rs30, %rs29, 255;\n\tcvt.rn.f32.u16 \t%f49, %rs30;\n\tcvt.rn.f32.u16 \t%f50, %rs28;\n\tcvt.rn.f32.u16 \t%f51, %rs26;\n\tadd.f32 \t%f126, %f40, %f51;\n\tadd.f32 \t%f127, %f41, %f50;\n\tadd.f32 \t%f128, %f42, %f49;\n\tmul.f32 \t%f52, %f51, 0f3E59B3D0;\n\tfma.rn.f32 \t%f53, %f50, 0f3F371759, %f52;\n\tfma.rn.f32 \t%f54, %f49, 0f3D93DD98, %f53;\n\tfma.rn.f32 \t%f55, %f54, 0f405BDBDC, 0f42800000;\n\tmax.f32 \t%f56, %f55, %f19;\n\tmin.f32 \t%f57, %f56, %f21;\n\tcvt.rni.s32.f32 \t%r55, %f57;\n\tcvt.u16.u32 \t%rs31, %r55;\n\tshl.b16 \t%rs32, %rs31, 6;\n\tst.global.u16 \t[%rd5], %rs32;\n\tbra.uni \t$L__BB3_4;\n\n$L__BB3_3:\n\tcvt.rn.f32.s32 \t%f58, %r17;\n\tcvt.rn.f32.s32 \t%f59, %r19;\n\ttex.2d.v4.u32.f32 \t{%r70, %r71, %r72, %r73}, [%rd6, {%f59, %f58}];\n\tcvt.u16.u32 \t%rs33, %r70;\n\tand.b16 \t%rs34, %rs33, 255;\n\tcvt.u16.u32 \t%rs35, %r71;\n\tand.b16 \t%rs36, %rs35, 255;\n\tcvt.u16.u32 \t%rs37, %r72;\n\tand.b16 \t%rs38, %rs37, 255;\n\tcvt.rn.f32.u16 \t%f60, %rs34;\n\tcvt.rn.f32.u16 \t%f61, %rs36;\n\tcvt.rn.f32.u16 \t%f62, %rs38;\n\tmul.f32 \t%f63, %f62, 0f3E59B3D0;\n\tfma.rn.f32 \t%f64, %f61, 0f3F371759, %f63;\n\tfma.rn.f32 \t%f65, %f60, 0f3D93DD98, %f64;\n\tfma.rn.f32 \t%f66, %f65, 0f405BDBDC, 0f42800000;\n\tmov.f32 \t%f67, 0f00000000;\n\tmax.f32 \t%f68, %f66, %f67;\n\tmov.f32 \t%f69, 0f447FC000;\n\tmin.f32 \t%f70, %f68, %f69;\n\tcvt.rni.s32.f32 \t%r74, %f70;\n\tcvt.u16.u32 \t%rs39, %r74;\n\tshl.b16 \t%rs40, %rs39, 6;\n\tst.global.u16 \t[%rd2], %rs40;\n\tor.b32 \t%r75, %r2, 1;\n\tmin.s32 \t%r76, %r75, %r18;\n\tcvt.rn.f32.s32 \t%f71, %r76;\n\ttex.2d.v4.u32.f32 \t{%r77, %r78, %r79, %r80}, [%rd6, {%f71, %f58}];\n\tcvt.u16.u32 \t%rs41, %r77;\n\tand.b16 \t%rs42, %rs41, 255;\n\tcvt.u16.u32 \t%rs43, %r78;\n\tand.b16 \t%rs44, %rs43, 255;\n\tcvt.u16.u32 \t%rs45, %r79;\n\tand.b16 \t%rs46, %rs45, 255;\n\tcvt.rn.f32.u16 \t%f72, %rs42;\n\tcvt.rn.f32.u16 \t%f73, %rs44;\n\tcvt.rn.f32.u16 \t%f74, %rs46;\n\tadd.f32 \t%f75, %f62, %f74;\n\tadd.f32 \t%f76, %f61, %f73;\n\tadd.f32 \t%f77, %f60, %f72;\n\tmul.f32 \t%f78, %f74, 0f3E59B3D0;\n\tfma.rn.f32 \t%f79, %f73, 0f3F371759, %f78;\n\tfma.rn.f32 \t%f80, %f72, 0f3D93DD98, %f79;\n\tfma.rn.f32 \t%f81, %f80, 0f405BDBDC, 0f42800000;\n\tmax.f32 \t%f82, %f81, %f67;\n\tmin.f32 \t%f83, %f82, %f69;\n\tcvt.rni.s32.f32 \t%r81, %f83;\n\tcvt.u16.u32 \t%rs47, %r81;\n\tshl.b16 \t%rs48, %rs47, 6;\n\tst.global.u16 \t[%rd3], %rs48;\n\tcvt.rn.f32.s32 \t%f84, %r4;\n\ttex.2d.v4.u32.f32 \t{%r82, %r83, %r84, %r85}, [%rd6, {%f59, %f84}];\n\tcvt.u16.u32 \t%rs49, %r82;\n\tand.b16 \t%rs50, %rs49, 255;\n\tcvt.u16.u32 \t%rs51, %r83;\n\tand.b16 \t%rs52, %rs51, 255;\n\tcvt.u16.u32 \t%rs53, %r84;\n\tand.b16 \t%rs54, %rs53, 255;\n\tcvt.rn.f32.u16 \t%f85, %rs50;\n\tcvt.rn.f32.u16 \t%f86, %rs52;\n\tcvt.rn.f32.u16 \t%f87, %rs54;\n\tadd.f32 \t%f88, %f75, %f87;\n\tadd.f32 \t%f89, %f76, %f86;\n\tadd.f32 \t%f90, %f77, %f85;\n\tmul.f32 \t%f91, %f87, 0f3E59B3D0;\n\tfma.rn.f32 \t%f92, %f86, 0f3F371759, %f91;\n\tfma.rn.f32 \t%f93, %f85, 0f3D93DD98, %f92;\n\tfma.rn.f32 \t%f94, %f93, 0f405BDBDC, 0f42800000;\n\tmax.f32 \t%f95, %f94, %f67;\n\tmin.f32 \t%f96, %f95, %f69;\n\tcvt.rni.s32.f32 \t%r86, %f96;\n\tcvt.u16.u32 \t%rs55, %r86;\n\tshl.b16 \t%rs56, %rs55, 6;\n\tst.global.u16 \t[%rd4], %rs56;\n\ttex.2d.v4.u32.f32 \t{%r87, %r88, %r89, %r90}, [%rd6, {%f71, %f84}];\n\tcvt.u16.u32 \t%rs57, %r87;\n\tand.b16 \t%rs58, %rs57, 255;\n\tcvt.u16.u32 \t%rs59, %r88;\n\tand.b16 \t%rs60, %rs59, 255;\n\tcvt.u16.u32 \t%rs61, %r89;\n\tand.b16 \t%rs62, %rs61, 255;\n\tcvt.rn.f32.u16 \t%f97, %rs58;\n\tcvt.rn.f32.u16 \t%f98, %rs60;\n\tcvt.rn.f32.u16 \t%f99, %rs62;\n\tadd.f32 \t%f126, %f88, %f99;\n\tadd.f32 \t%f127, %f89, %f98;\n\tadd.f32 \t%f128, %f90, %f97;\n\tmul.f32 \t%f100, %f99, 0f3E59B3D0;\n\tfma.rn.f32 \t%f101, %f98, 0f3F371759, %f100;\n\tfma.rn.f32 \t%f102, %f97, 0f3D93DD98, %f101;\n\tfma.rn.f32 \t%f103, %f102, 0f405BDBDC, 0f42800000;\n\tmax.f32 \t%f104, %f103, %f67;\n\tmin.f32 \t%f105, %f104, %f69;\n\tcvt.rni.s32.f32 \t%r91, %f105;\n\tcvt.u16.u32 \t%rs63, %r91;\n\tshl.b16 \t%rs64, %rs63, 6;\n\tst.global.u16 \t[%rd5], %rs64;\n\n$L__BB3_4:\n\tmul.f32 \t%f106, %f126, 0f3E800000;\n\tmul.f32 \t%f107, %f127, 0f3E800000;\n\tmul.f32 \t%f108, %f107, 0f3F371759;\n\tfma.rn.f32 \t%f109, %f106, 0f3E59B3D0, %f108;\n\tmul.f32 \t%f110, %f128, 0f3E800000;\n\tfma.rn.f32 \t%f111, %f110, 0f3D93DD98, %f109;\n\tsub.f32 \t%f112, %f110, %f111;\n\tmul.f32 \t%f113, %f112, 0f4060E0E1;\n\tdiv.rn.f32 \t%f114, %f113, 0f3FED844D;\n\tadd.f32 \t%f115, %f114, 0f44000000;\n\tmov.f32 \t%f116, 0f00000000;\n\tmax.f32 \t%f117, %f115, %f116;\n\tmov.f32 \t%f118, 0f447FC000;\n\tmin.f32 \t%f119, %f117, %f118;\n\tcvt.rni.s32.f32 \t%r92, %f119;\n\tcvt.u16.u32 \t%rs65, %r92;\n\tshl.b16 \t%rs66, %rs65, 6;\n\tst.global.u16 \t[%rd1], %rs66;\n\tsub.f32 \t%f120, %f106, %f111;\n\tmul.f32 \t%f121, %f120, 0f4060E0E1;\n\tdiv.rn.f32 \t%f122, %f121, 0f3FC9930C;\n\tadd.f32 \t%f123, %f122, 0f44000000;\n\tmax.f32 \t%f124, %f123, %f116;\n\tmin.f32 \t%f125, %f124, %f118;\n\tcvt.rni.s32.f32 \t%r93, %f125;\n\tcvt.u16.u32 \t%rs67, %r93;\n\tshl.b16 \t%rs68, %rs67, 6;\n\tst.global.u16 \t[%rd1+2], %rs68;\n\n$L__BB3_5:\n\tret;\n\n}\n\t// .globl\targb_to_yuv444p10\n.visible .entry argb_to_yuv444p10(\n\t.param .u64 argb_to_yuv444p10_param_0,\n\t.param .u32 argb_to_yuv444p10_param_1,\n\t.param .u64 argb_to_yuv444p10_param_2,\n\t.param .u32 argb_to_yuv444p10_param_3,\n\t.param .u32 argb_to_yuv444p10_param_4,\n\t.param .u32 argb_to_yuv444p10_param_5,\n\t.param .u32 argb_to_yuv444p10_param_6\n)\n{\n\t.reg .pred \t%p<8>;\n\t.reg .b16 \t%rs<49>;\n\t.reg .f32 \t%f<93>;\n\t.reg .b32 \t%r<30>;\n\t.reg .b64 \t%rd<55>;\n\n\n\tld.param.u64 \t%rd18, [argb_to_yuv444p10_param_0];\n\tld.param.u32 \t%r5, [argb_to_yuv444p10_param_1];\n\tld.param.u64 \t%rd19, [argb_to_yuv444p10_param_2];\n\tld.param.u32 \t%r6, [argb_to_yuv444p10_param_3];\n\tld.param.u32 \t%r7, [argb_to_yuv444p10_param_4];\n\tld.param.u32 \t%r8, [argb_to_yuv444p10_param_5];\n\tld.param.u32 \t%r9, [argb_to_yuv444p10_param_6];\n\tcvta.to.global.u64 \t%rd1, %rd19;\n\tcvta.to.global.u64 \t%rd2, %rd18;\n\tmov.u32 \t%r10, %ntid.x;\n\tmov.u32 \t%r11, %ctaid.x;\n\tmov.u32 \t%r12, %tid.x;\n\tmad.lo.s32 \t%r13, %r11, %r10, %r12;\n\tmov.u32 \t%r14, %ntid.y;\n\tmov.u32 \t%r15, %ctaid.y;\n\tmov.u32 \t%r16, %tid.y;\n\tmad.lo.s32 \t%r17, %r15, %r14, %r16;\n\tshl.b32 \t%r1, %r13, 1;\n\tsetp.ge.s32 \t%p1, %r1, %r7;\n\tshl.b32 \t%r2, %r17, 1;\n\tsetp.ge.s32 \t%p2, %r2, %r8;\n\tor.pred \t%p3, %p1, %p2;\n\t@%p3 bra \t$L__BB4_6;\n\n\tcvt.s64.s32 \t%rd3, %r6;\n\tsetp.eq.s32 \t%p4, %r9, 0;\n\tselp.b64 \t%rd4, 0, 2, %p4;\n\tselp.b64 \t%rd5, 2, 0, %p4;\n\tmul.wide.s32 \t%rd20, %r2, %r6;\n\tadd.s64 \t%rd21, %rd1, %rd20;\n\tcvt.s64.s32 \t%rd6, %r5;\n\tmul.wide.s32 \t%rd7, %r2, %r5;\n\tcvt.s64.s32 \t%rd8, %r1;\n\tmul.wide.s32 \t%rd9, %r1, 4;\n\tadd.s64 \t%rd22, %rd9, %rd7;\n\tadd.s64 \t%rd23, %rd22, %rd4;\n\tadd.s64 \t%rd24, %rd2, %rd23;\n\tld.global.u8 \t%rs1, [%rd24];\n\tcvt.rn.f32.u16 \t%f1, %rs1;\n\tadd.s64 \t%rd25, %rd2, %rd22;\n\tld.global.u8 \t%rs3, [%rd25+1];\n\tcvt.rn.f32.u16 \t%f2, %rs3;\n\tadd.s64 \t%rd26, %rd22, %rd5;\n\tadd.s64 \t%rd27, %rd2, %rd26;\n\tld.global.u8 \t%rs5, [%rd27];\n\tcvt.rn.f32.u16 \t%f3, %rs5;\n\tmul.f32 \t%f4, %f3, 0f3E59B3D0;\n\tfma.rn.f32 \t%f5, %f2, 0f3F371759, %f4;\n\tfma.rn.f32 \t%f6, %f1, 0f3D93DD98, %f5;\n\tsub.f32 \t%f7, %f1, %f6;\n\tmul.f32 \t%f8, %f7, 0f4060E0E1;\n\tdiv.rn.f32 \t%f9, %f8, 0f3FED844D;\n\tadd.f32 \t%f10, %f9, 0f44000000;\n\tmov.f32 \t%f11, 0f00000000;\n\tmax.f32 \t%f12, %f10, %f11;\n\tmov.f32 \t%f13, 0f447FC000;\n\tmin.f32 \t%f14, %f12, %f13;\n\tcvt.rni.s32.f32 \t%r18, %f14;\n\tcvt.u16.u32 \t%rs7, %r18;\n\tshl.b16 \t%rs8, %rs7, 6;\n\tsub.f32 \t%f15, %f3, %f6;\n\tmul.f32 \t%f16, %f15, 0f4060E0E1;\n\tdiv.rn.f32 \t%f17, %f16, 0f3FC9930C;\n\tadd.f32 \t%f18, %f17, 0f44000000;\n\tmax.f32 \t%f19, %f18, %f11;\n\tmin.f32 \t%f20, %f19, %f13;\n\tcvt.rni.s32.f32 \t%r19, %f20;\n\tcvt.u16.u32 \t%rs9, %r19;\n\tshl.b16 \t%rs10, %rs9, 6;\n\tfma.rn.f32 \t%f21, %f6, 0f405BDBDC, 0f42800000;\n\tmax.f32 \t%f22, %f21, %f11;\n\tmin.f32 \t%f23, %f22, %f13;\n\tcvt.rni.s32.f32 \t%r20, %f23;\n\tcvt.u16.u32 \t%rs11, %r20;\n\tshl.b16 \t%rs12, %rs11, 6;\n\tmul.wide.s32 \t%rd28, %r1, 2;\n\tadd.s64 \t%rd10, %rd21, %rd28;\n\tst.global.u16 \t[%rd10], %rs12;\n\tcvt.s64.s32 \t%rd11, %r8;\n\tmul.wide.s32 \t%rd29, %r8, %r6;\n\tadd.s64 \t%rd12, %rd10, %rd29;\n\tst.global.u16 \t[%rd12], %rs8;\n\tadd.s64 \t%rd13, %rd12, %rd29;\n\tst.global.u16 \t[%rd13], %rs10;\n\tadd.s32 \t%r3, %r1, 1;\n\tsetp.ge.s32 \t%p5, %r3, %r7;\n\t@%p5 bra \t$L__BB4_3;\n\n\tmul.wide.s32 \t%rd30, %r3, 4;\n\tadd.s64 \t%rd31, %rd30, %rd7;\n\tadd.s64 \t%rd32, %rd31, %rd4;\n\tadd.s64 \t%rd33, %rd2, %rd32;\n\tld.global.u8 \t%rs13, [%rd33];\n\tcvt.rn.f32.u16 \t%f24, %rs13;\n\tadd.s64 \t%rd34, %rd2, %rd31;\n\tld.global.u8 \t%rs15, [%rd34+1];\n\tcvt.rn.f32.u16 \t%f25, %rs15;\n\tadd.s64 \t%rd35, %rd31, %rd5;\n\tadd.s64 \t%rd36, %rd2, %rd35;\n\tld.global.u8 \t%rs17, [%rd36];\n\tcvt.rn.f32.u16 \t%f26, %rs17;\n\tmul.f32 \t%f27, %f26, 0f3E59B3D0;\n\tfma.rn.f32 \t%f28, %f25, 0f3F371759, %f27;\n\tfma.rn.f32 \t%f29, %f24, 0f3D93DD98, %f28;\n\tsub.f32 \t%f30, %f24, %f29;\n\tmul.f32 \t%f31, %f30, 0f4060E0E1;\n\tdiv.rn.f32 \t%f32, %f31, 0f3FED844D;\n\tadd.f32 \t%f33, %f32, 0f44000000;\n\tmax.f32 \t%f35, %f33, %f11;\n\tmin.f32 \t%f37, %f35, %f13;\n\tcvt.rni.s32.f32 \t%r21, %f37;\n\tcvt.u16.u32 \t%rs19, %r21;\n\tshl.b16 \t%rs20, %rs19, 6;\n\tsub.f32 \t%f38, %f26, %f29;\n\tmul.f32 \t%f39, %f38, 0f4060E0E1;\n\tdiv.rn.f32 \t%f40, %f39, 0f3FC9930C;\n\tadd.f32 \t%f41, %f40, 0f44000000;\n\tmax.f32 \t%f42, %f41, %f11;\n\tmin.f32 \t%f43, %f42, %f13;\n\tcvt.rni.s32.f32 \t%r22, %f43;\n\tcvt.u16.u32 \t%rs21, %r22;\n\tshl.b16 \t%rs22, %rs21, 6;\n\tfma.rn.f32 \t%f44, %f29, 0f405BDBDC, 0f42800000;\n\tmax.f32 \t%f45, %f44, %f11;\n\tmin.f32 \t%f46, %f45, %f13;\n\tcvt.rni.s32.f32 \t%r23, %f46;\n\tcvt.u16.u32 \t%rs23, %r23;\n\tshl.b16 \t%rs24, %rs23, 6;\n\tst.global.u16 \t[%rd10+2], %rs24;\n\tst.global.u16 \t[%rd12+2], %rs20;\n\tst.global.u16 \t[%rd13+2], %rs22;\n\n$L__BB4_3:\n\tadd.s32 \t%r4, %r2, 1;\n\tsetp.ge.s32 \t%p6, %r4, %r8;\n\t@%p6 bra \t$L__BB4_6;\n\n\tcvt.s64.s32 \t%rd37, %r4;\n\tmul.lo.s64 \t%rd38, %rd37, %rd3;\n\tadd.s64 \t%rd39, %rd1, %rd38;\n\tmul.lo.s64 \t%rd14, %rd37, %rd6;\n\tadd.s64 \t%rd40, %rd9, %rd14;\n\tadd.s64 \t%rd41, %rd40, %rd4;\n\tadd.s64 \t%rd42, %rd2, %rd41;\n\tld.global.u8 \t%rs25, [%rd42];\n\tcvt.rn.f32.u16 \t%f47, %rs25;\n\tadd.s64 \t%rd43, %rd2, %rd40;\n\tld.global.u8 \t%rs27, [%rd43+1];\n\tcvt.rn.f32.u16 \t%f48, %rs27;\n\tadd.s64 \t%rd44, %rd40, %rd5;\n\tadd.s64 \t%rd45, %rd2, %rd44;\n\tld.global.u8 \t%rs29, [%rd45];\n\tcvt.rn.f32.u16 \t%f49, %rs29;\n\tmul.f32 \t%f50, %f49, 0f3E59B3D0;\n\tfma.rn.f32 \t%f51, %f48, 0f3F371759, %f50;\n\tfma.rn.f32 \t%f52, %f47, 0f3D93DD98, %f51;\n\tsub.f32 \t%f53, %f47, %f52;\n\tmul.f32 \t%f54, %f53, 0f4060E0E1;\n\tdiv.rn.f32 \t%f55, %f54, 0f3FED844D;\n\tadd.f32 \t%f56, %f55, 0f44000000;\n\tmov.f32 \t%f57, 0f00000000;\n\tmax.f32 \t%f58, %f56, %f57;\n\tmov.f32 \t%f59, 0f447FC000;\n\tmin.f32 \t%f60, %f58, %f59;\n\tcvt.rni.s32.f32 \t%r24, %f60;\n\tcvt.u16.u32 \t%rs31, %r24;\n\tshl.b16 \t%rs32, %rs31, 6;\n\tsub.f32 \t%f61, %f49, %f52;\n\tmul.f32 \t%f62, %f61, 0f4060E0E1;\n\tdiv.rn.f32 \t%f63, %f62, 0f3FC9930C;\n\tadd.f32 \t%f64, %f63, 0f44000000;\n\tmax.f32 \t%f65, %f64, %f57;\n\tmin.f32 \t%f66, %f65, %f59;\n\tcvt.rni.s32.f32 \t%r25, %f66;\n\tcvt.u16.u32 \t%rs33, %r25;\n\tshl.b16 \t%rs34, %rs33, 6;\n\tfma.rn.f32 \t%f67, %f52, 0f405BDBDC, 0f42800000;\n\tmax.f32 \t%f68, %f67, %f57;\n\tmin.f32 \t%f69, %f68, %f59;\n\tcvt.rni.s32.f32 \t%r26, %f69;\n\tcvt.u16.u32 \t%rs35, %r26;\n\tshl.b16 \t%rs36, %rs35, 6;\n\tshl.b64 \t%rd46, %rd8, 1;\n\tadd.s64 \t%rd15, %rd39, %rd46;\n\tst.global.u16 \t[%rd15], %rs36;\n\tmul.lo.s64 \t%rd47, %rd11, %rd3;\n\tadd.s64 \t%rd16, %rd15, %rd47;\n\tst.global.u16 \t[%rd16], %rs32;\n\tadd.s64 \t%rd17, %rd16, %rd47;\n\tst.global.u16 \t[%rd17], %rs34;\n\t@%p5 bra \t$L__BB4_6;\n\n\tmul.wide.s32 \t%rd48, %r3, 4;\n\tadd.s64 \t%rd49, %rd48, %rd14;\n\tadd.s64 \t%rd50, %rd49, %rd4;\n\tadd.s64 \t%rd51, %rd2, %rd50;\n\tld.global.u8 \t%rs37, [%rd51];\n\tcvt.rn.f32.u16 \t%f70, %rs37;\n\tadd.s64 \t%rd52, %rd2, %rd49;\n\tld.global.u8 \t%rs39, [%rd52+1];\n\tcvt.rn.f32.u16 \t%f71, %rs39;\n\tadd.s64 \t%rd53, %rd49, %rd5;\n\tadd.s64 \t%rd54, %rd2, %rd53;\n\tld.global.u8 \t%rs41, [%rd54];\n\tcvt.rn.f32.u16 \t%f72, %rs41;\n\tmul.f32 \t%f73, %f72, 0f3E59B3D0;\n\tfma.rn.f32 \t%f74, %f71, 0f3F371759, %f73;\n\tfma.rn.f32 \t%f75, %f70, 0f3D93DD98, %f74;\n\tsub.f32 \t%f76, %f70, %f75;\n\tmul.f32 \t%f77, %f76, 0f4060E0E1;\n\tdiv.rn.f32 \t%f78, %f77, 0f3FED844D;\n\tadd.f32 \t%f79, %f78, 0f44000000;\n\tmax.f32 \t%f81, %f79, %f57;\n\tmin.f32 \t%f83, %f81, %f59;\n\tcvt.rni.s32.f32 \t%r27, %f83;\n\tcvt.u16.u32 \t%rs43, %r27;\n\tshl.b16 \t%rs44, %rs43, 6;\n\tsub.f32 \t%f84, %f72, %f75;\n\tmul.f32 \t%f85, %f84, 0f4060E0E1;\n\tdiv.rn.f32 \t%f86, %f85, 0f3FC9930C;\n\tadd.f32 \t%f87, %f86, 0f44000000;\n\tmax.f32 \t%f88, %f87, %f57;\n\tmin.f32 \t%f89, %f88, %f59;\n\tcvt.rni.s32.f32 \t%r28, %f89;\n\tcvt.u16.u32 \t%rs45, %r28;\n\tshl.b16 \t%rs46, %rs45, 6;\n\tfma.rn.f32 \t%f90, %f75, 0f405BDBDC, 0f42800000;\n\tmax.f32 \t%f91, %f90, %f57;\n\tmin.f32 \t%f92, %f91, %f59;\n\tcvt.rni.s32.f32 \t%r29, %f92;\n\tcvt.u16.u32 \t%rs47, %r29;\n\tshl.b16 \t%rs48, %rs47, 6;\n\tst.global.u16 \t[%rd15+2], %rs48;\n\tst.global.u16 \t[%rd16+2], %rs44;\n\tst.global.u16 \t[%rd17+2], %rs46;\n\n$L__BB4_6:\n\tret;\n\n}\n\t// .globl\targb_tex_to_yuv444p10\n.visible .entry argb_tex_to_yuv444p10(\n\t.param .u64 argb_tex_to_yuv444p10_param_0,\n\t.param .u64 argb_tex_to_yuv444p10_param_1,\n\t.param .u32 argb_tex_to_yuv444p10_param_2,\n\t.param .u32 argb_tex_to_yuv444p10_param_3,\n\t.param .u32 argb_tex_to_yuv444p10_param_4,\n\t.param .u32 argb_tex_to_yuv444p10_param_5\n)\n{\n\t.reg .pred \t%p<11>;\n\t.reg .b16 \t%rs<97>;\n\t.reg .f32 \t%f<193>;\n\t.reg .b32 \t%r<75>;\n\t.reg .b64 \t%rd<56>;\n\n\n\tld.param.u64 \t%rd11, [argb_tex_to_yuv444p10_param_0];\n\tld.param.u64 \t%rd12, [argb_tex_to_yuv444p10_param_1];\n\tld.param.u32 \t%r7, [argb_tex_to_yuv444p10_param_2];\n\tld.param.u32 \t%r8, [argb_tex_to_yuv444p10_param_3];\n\tld.param.u32 \t%r9, [argb_tex_to_yuv444p10_param_4];\n\tld.param.u32 \t%r10, [argb_tex_to_yuv444p10_param_5];\n\tcvta.to.global.u64 \t%rd1, %rd12;\n\tmov.u32 \t%r11, %ntid.x;\n\tmov.u32 \t%r12, %ctaid.x;\n\tmov.u32 \t%r13, %tid.x;\n\tmad.lo.s32 \t%r14, %r12, %r11, %r13;\n\tmov.u32 \t%r15, %ntid.y;\n\tmov.u32 \t%r16, %ctaid.y;\n\tmov.u32 \t%r17, %tid.y;\n\tmad.lo.s32 \t%r18, %r16, %r15, %r17;\n\tshl.b32 \t%r1, %r14, 1;\n\tsetp.ge.s32 \t%p1, %r1, %r8;\n\tshl.b32 \t%r2, %r18, 1;\n\tsetp.ge.s32 \t%p2, %r2, %r9;\n\tor.pred \t%p3, %p1, %p2;\n\t@%p3 bra \t$L__BB5_12;\n\n\tcvt.s64.s32 \t%rd2, %r7;\n\tcvt.s64.s32 \t%rd3, %r9;\n\tmul.wide.s32 \t%rd13, %r2, %r7;\n\tadd.s64 \t%rd4, %rd1, %rd13;\n\tcvt.rn.f32.s32 \t%f1, %r2;\n\tcvt.rn.f32.s32 \t%f2, %r1;\n\tsetp.eq.s32 \t%p4, %r10, 0;\n\t@%p4 bra \t$L__BB5_7;\n\n\ttex.2d.v4.u32.f32 \t{%r19, %r20, %r21, %r22}, [%rd11, {%f2, %f1}];\n\tcvt.u16.u32 \t%rs1, %r19;\n\tand.b16 \t%rs2, %rs1, 255;\n\tcvt.u16.u32 \t%rs3, %r20;\n\tand.b16 \t%rs4, %rs3, 255;\n\tcvt.u16.u32 \t%rs5, %r21;\n\tand.b16 \t%rs6, %rs5, 255;\n\tcvt.rn.f32.u16 \t%f5, %rs6;\n\tcvt.rn.f32.u16 \t%f6, %rs4;\n\tcvt.rn.f32.u16 \t%f7, %rs2;\n\tmul.f32 \t%f8, %f7, 0f3E59B3D0;\n\tfma.rn.f32 \t%f9, %f6, 0f3F371759, %f8;\n\tfma.rn.f32 \t%f10, %f5, 0f3D93DD98, %f9;\n\tsub.f32 \t%f11, %f5, %f10;\n\tmul.f32 \t%f12, %f11, 0f4060E0E1;\n\tdiv.rn.f32 \t%f13, %f12, 0f3FED844D;\n\tadd.f32 \t%f14, %f13, 0f44000000;\n\tmov.f32 \t%f15, 0f00000000;\n\tmax.f32 \t%f16, %f14, %f15;\n\tmov.f32 \t%f17, 0f447FC000;\n\tmin.f32 \t%f18, %f16, %f17;\n\tcvt.rni.s32.f32 \t%r23, %f18;\n\tcvt.u16.u32 \t%rs7, %r23;\n\tshl.b16 \t%rs8, %rs7, 6;\n\tsub.f32 \t%f19, %f7, %f10;\n\tmul.f32 \t%f20, %f19, 0f4060E0E1;\n\tdiv.rn.f32 \t%f21, %f20, 0f3FC9930C;\n\tadd.f32 \t%f22, %f21, 0f44000000;\n\tmax.f32 \t%f23, %f22, %f15;\n\tmin.f32 \t%f24, %f23, %f17;\n\tcvt.rni.s32.f32 \t%r24, %f24;\n\tcvt.u16.u32 \t%rs9, %r24;\n\tshl.b16 \t%rs10, %rs9, 6;\n\tfma.rn.f32 \t%f25, %f10, 0f405BDBDC, 0f42800000;\n\tmax.f32 \t%f26, %f25, %f15;\n\tmin.f32 \t%f27, %f26, %f17;\n\tcvt.rni.s32.f32 \t%r25, %f27;\n\tcvt.u16.u32 \t%rs11, %r25;\n\tshl.b16 \t%rs12, %rs11, 6;\n\tcvt.s64.s32 \t%rd5, %r1;\n\tmul.wide.s32 \t%rd14, %r1, 2;\n\tadd.s64 \t%rd6, %rd4, %rd14;\n\tst.global.u16 \t[%rd6], %rs12;\n\tmul.lo.s64 \t%rd15, %rd3, %rd2;\n\tadd.s64 \t%rd16, %rd6, %rd15;\n\tst.global.u16 \t[%rd16], %rs8;\n\tadd.s64 \t%rd17, %rd16, %rd15;\n\tst.global.u16 \t[%rd17], %rs10;\n\tadd.s32 \t%r3, %r1, 1;\n\tsetp.ge.s32 \t%p5, %r3, %r8;\n\t@%p5 bra \t$L__BB5_4;\n\n\tcvt.rn.f32.s32 \t%f28, %r3;\n\ttex.2d.v4.u32.f32 \t{%r26, %r27, %r28, %r29}, [%rd11, {%f28, %f1}];\n\tcvt.u16.u32 \t%rs13, %r26;\n\tand.b16 \t%rs14, %rs13, 255;\n\tcvt.u16.u32 \t%rs15, %r27;\n\tand.b16 \t%rs16, %rs15, 255;\n\tcvt.u16.u32 \t%rs17, %r28;\n\tand.b16 \t%rs18, %rs17, 255;\n\tcvt.rn.f32.u16 \t%f29, %rs18;\n\tcvt.rn.f32.u16 \t%f30, %rs16;\n\tcvt.rn.f32.u16 \t%f31, %rs14;\n\tmul.f32 \t%f32, %f31, 0f3E59B3D0;\n\tfma.rn.f32 \t%f33, %f30, 0f3F371759, %f32;\n\tfma.rn.f32 \t%f34, %f29, 0f3D93DD98, %f33;\n\tsub.f32 \t%f35, %f29, %f34;\n\tmul.f32 \t%f36, %f35, 0f4060E0E1;\n\tdiv.rn.f32 \t%f37, %f36, 0f3FED844D;\n\tadd.f32 \t%f38, %f37, 0f44000000;\n\tmax.f32 \t%f40, %f38, %f15;\n\tmin.f32 \t%f42, %f40, %f17;\n\tcvt.rni.s32.f32 \t%r30, %f42;\n\tcvt.u16.u32 \t%rs19, %r30;\n\tshl.b16 \t%rs20, %rs19, 6;\n\tsub.f32 \t%f43, %f31, %f34;\n\tmul.f32 \t%f44, %f43, 0f4060E0E1;\n\tdiv.rn.f32 \t%f45, %f44, 0f3FC9930C;\n\tadd.f32 \t%f46, %f45, 0f44000000;\n\tmax.f32 \t%f47, %f46, %f15;\n\tmin.f32 \t%f48, %f47, %f17;\n\tcvt.rni.s32.f32 \t%r31, %f48;\n\tcvt.u16.u32 \t%rs21, %r31;\n\tshl.b16 \t%rs22, %rs21, 6;\n\tfma.rn.f32 \t%f49, %f34, 0f405BDBDC, 0f42800000;\n\tmax.f32 \t%f50, %f49, %f15;\n\tmin.f32 \t%f51, %f50, %f17;\n\tcvt.rni.s32.f32 \t%r32, %f51;\n\tcvt.u16.u32 \t%rs23, %r32;\n\tshl.b16 \t%rs24, %rs23, 6;\n\tadd.s64 \t%rd18, %rd6, 2;\n\tst.global.u16 \t[%rd6+2], %rs24;\n\tadd.s64 \t%rd20, %rd18, %rd15;\n\tst.global.u16 \t[%rd16+2], %rs20;\n\tadd.s64 \t%rd22, %rd20, %rd15;\n\tst.global.u16 \t[%rd22], %rs22;\n\n$L__BB5_4:\n\tadd.s32 \t%r4, %r2, 1;\n\tsetp.ge.s32 \t%p6, %r4, %r9;\n\t@%p6 bra \t$L__BB5_12;\n\n\tcvt.s64.s32 \t%rd23, %r4;\n\tmul.lo.s64 \t%rd24, %rd23, %rd2;\n\tadd.s64 \t%rd25, %rd1, %rd24;\n\tcvt.rn.f32.s32 \t%f3, %r4;\n\ttex.2d.v4.u32.f32 \t{%r33, %r34, %r35, %r36}, [%rd11, {%f2, %f3}];\n\tcvt.u16.u32 \t%rs25, %r33;\n\tand.b16 \t%rs26, %rs25, 255;\n\tcvt.u16.u32 \t%rs27, %r34;\n\tand.b16 \t%rs28, %rs27, 255;\n\tcvt.u16.u32 \t%rs29, %r35;\n\tand.b16 \t%rs30, %rs29, 255;\n\tcvt.rn.f32.u16 \t%f52, %rs30;\n\tcvt.rn.f32.u16 \t%f53, %rs28;\n\tcvt.rn.f32.u16 \t%f54, %rs26;\n\tmul.f32 \t%f55, %f54, 0f3E59B3D0;\n\tfma.rn.f32 \t%f56, %f53, 0f3F371759, %f55;\n\tfma.rn.f32 \t%f57, %f52, 0f3D93DD98, %f56;\n\tsub.f32 \t%f58, %f52, %f57;\n\tmul.f32 \t%f59, %f58, 0f4060E0E1;\n\tdiv.rn.f32 \t%f60, %f59, 0f3FED844D;\n\tadd.f32 \t%f61, %f60, 0f44000000;\n\tmov.f32 \t%f62, 0f00000000;\n\tmax.f32 \t%f63, %f61, %f62;\n\tmov.f32 \t%f64, 0f447FC000;\n\tmin.f32 \t%f65, %f63, %f64;\n\tcvt.rni.s32.f32 \t%r37, %f65;\n\tcvt.u16.u32 \t%rs31, %r37;\n\tshl.b16 \t%rs32, %rs31, 6;\n\tsub.f32 \t%f66, %f54, %f57;\n\tmul.f32 \t%f67, %f66, 0f4060E0E1;\n\tdiv.rn.f32 \t%f68, %f67, 0f3FC9930C;\n\tadd.f32 \t%f69, %f68, 0f44000000;\n\tmax.f32 \t%f70, %f69, %f62;\n\tmin.f32 \t%f71, %f70, %f64;\n\tcvt.rni.s32.f32 \t%r38, %f71;\n\tcvt.u16.u32 \t%rs33, %r38;\n\tshl.b16 \t%rs34, %rs33, 6;\n\tfma.rn.f32 \t%f72, %f57, 0f405BDBDC, 0f42800000;\n\tmax.f32 \t%f73, %f72, %f62;\n\tmin.f32 \t%f74, %f73, %f64;\n\tcvt.rni.s32.f32 \t%r39, %f74;\n\tcvt.u16.u32 \t%rs35, %r39;\n\tshl.b16 \t%rs36, %rs35, 6;\n\tshl.b64 \t%rd26, %rd5, 1;\n\tadd.s64 \t%rd7, %rd25, %rd26;\n\tst.global.u16 \t[%rd7], %rs36;\n\tadd.s64 \t%rd28, %rd7, %rd15;\n\tst.global.u16 \t[%rd28], %rs32;\n\tadd.s64 \t%rd29, %rd28, %rd15;\n\tst.global.u16 \t[%rd29], %rs34;\n\t@%p5 bra \t$L__BB5_12;\n\n\tcvt.rn.f32.s32 \t%f75, %r3;\n\ttex.2d.v4.u32.f32 \t{%r40, %r41, %r42, %r43}, [%rd11, {%f75, %f3}];\n\tcvt.u16.u32 \t%rs37, %r40;\n\tand.b16 \t%rs38, %rs37, 255;\n\tcvt.u16.u32 \t%rs39, %r41;\n\tand.b16 \t%rs40, %rs39, 255;\n\tcvt.u16.u32 \t%rs41, %r42;\n\tand.b16 \t%rs42, %rs41, 255;\n\tcvt.rn.f32.u16 \t%f76, %rs42;\n\tcvt.rn.f32.u16 \t%f77, %rs40;\n\tcvt.rn.f32.u16 \t%f78, %rs38;\n\tmul.f32 \t%f79, %f78, 0f3E59B3D0;\n\tfma.rn.f32 \t%f80, %f77, 0f3F371759, %f79;\n\tfma.rn.f32 \t%f81, %f76, 0f3D93DD98, %f80;\n\tsub.f32 \t%f82, %f76, %f81;\n\tmul.f32 \t%f83, %f82, 0f4060E0E1;\n\tdiv.rn.f32 \t%f84, %f83, 0f3FED844D;\n\tadd.f32 \t%f85, %f84, 0f44000000;\n\tmax.f32 \t%f87, %f85, %f62;\n\tmin.f32 \t%f89, %f87, %f64;\n\tcvt.rni.s32.f32 \t%r44, %f89;\n\tcvt.u16.u32 \t%rs43, %r44;\n\tshl.b16 \t%rs44, %rs43, 6;\n\tsub.f32 \t%f90, %f78, %f81;\n\tmul.f32 \t%f91, %f90, 0f4060E0E1;\n\tdiv.rn.f32 \t%f92, %f91, 0f3FC9930C;\n\tadd.f32 \t%f93, %f92, 0f44000000;\n\tmax.f32 \t%f94, %f93, %f62;\n\tmin.f32 \t%f95, %f94, %f64;\n\tcvt.rni.s32.f32 \t%r45, %f95;\n\tcvt.u16.u32 \t%rs45, %r45;\n\tshl.b16 \t%rs46, %rs45, 6;\n\tfma.rn.f32 \t%f96, %f81, 0f405BDBDC, 0f42800000;\n\tmax.f32 \t%f97, %f96, %f62;\n\tmin.f32 \t%f98, %f97, %f64;\n\tcvt.rni.s32.f32 \t%r46, %f98;\n\tcvt.u16.u32 \t%rs47, %r46;\n\tshl.b16 \t%rs48, %rs47, 6;\n\tadd.s64 \t%rd30, %rd7, 2;\n\tst.global.u16 \t[%rd7+2], %rs48;\n\tadd.s64 \t%rd32, %rd30, %rd15;\n\tst.global.u16 \t[%rd28+2], %rs44;\n\tadd.s64 \t%rd34, %rd32, %rd15;\n\tst.global.u16 \t[%rd34], %rs46;\n\tbra.uni \t$L__BB5_12;\n\n$L__BB5_7:\n\ttex.2d.v4.u32.f32 \t{%r47, %r48, %r49, %r50}, [%rd11, {%f2, %f1}];\n\tcvt.u16.u32 \t%rs49, %r47;\n\tand.b16 \t%rs50, %rs49, 255;\n\tcvt.u16.u32 \t%rs51, %r48;\n\tand.b16 \t%rs52, %rs51, 255;\n\tcvt.u16.u32 \t%rs53, %r49;\n\tand.b16 \t%rs54, %rs53, 255;\n\tcvt.rn.f32.u16 \t%f99, %rs50;\n\tcvt.rn.f32.u16 \t%f100, %rs52;\n\tcvt.rn.f32.u16 \t%f101, %rs54;\n\tmul.f32 \t%f102, %f101, 0f3E59B3D0;\n\tfma.rn.f32 \t%f103, %f100, 0f3F371759, %f102;\n\tfma.rn.f32 \t%f104, %f99, 0f3D93DD98, %f103;\n\tsub.f32 \t%f105, %f99, %f104;\n\tmul.f32 \t%f106, %f105, 0f4060E0E1;\n\tdiv.rn.f32 \t%f107, %f106, 0f3FED844D;\n\tadd.f32 \t%f108, %f107, 0f44000000;\n\tmov.f32 \t%f109, 0f00000000;\n\tmax.f32 \t%f110, %f108, %f109;\n\tmov.f32 \t%f111, 0f447FC000;\n\tmin.f32 \t%f112, %f110, %f111;\n\tcvt.rni.s32.f32 \t%r51, %f112;\n\tcvt.u16.u32 \t%rs55, %r51;\n\tshl.b16 \t%rs56, %rs55, 6;\n\tsub.f32 \t%f113, %f101, %f104;\n\tmul.f32 \t%f114, %f113, 0f4060E0E1;\n\tdiv.rn.f32 \t%f115, %f114, 0f3FC9930C;\n\tadd.f32 \t%f116, %f115, 0f44000000;\n\tmax.f32 \t%f117, %f116, %f109;\n\tmin.f32 \t%f118, %f117, %f111;\n\tcvt.rni.s32.f32 \t%r52, %f118;\n\tcvt.u16.u32 \t%rs57, %r52;\n\tshl.b16 \t%rs58, %rs57, 6;\n\tfma.rn.f32 \t%f119, %f104, 0f405BDBDC, 0f42800000;\n\tmax.f32 \t%f120, %f119, %f109;\n\tmin.f32 \t%f121, %f120, %f111;\n\tcvt.rni.s32.f32 \t%r53, %f121;\n\tcvt.u16.u32 \t%rs59, %r53;\n\tshl.b16 \t%rs60, %rs59, 6;\n\tcvt.s64.s32 \t%rd8, %r1;\n\tmul.wide.s32 \t%rd35, %r1, 2;\n\tadd.s64 \t%rd9, %rd4, %rd35;\n\tst.global.u16 \t[%rd9], %rs60;\n\tmul.lo.s64 \t%rd36, %rd3, %rd2;\n\tadd.s64 \t%rd37, %rd9, %rd36;\n\tst.global.u16 \t[%rd37], %rs56;\n\tadd.s64 \t%rd38, %rd37, %rd36;\n\tst.global.u16 \t[%rd38], %rs58;\n\tadd.s32 \t%r5, %r1, 1;\n\tsetp.ge.s32 \t%p8, %r5, %r8;\n\t@%p8 bra \t$L__BB5_9;\n\n\tcvt.rn.f32.s32 \t%f122, %r5;\n\ttex.2d.v4.u32.f32 \t{%r54, %r55, %r56, %r57}, [%rd11, {%f122, %f1}];\n\tcvt.u16.u32 \t%rs61, %r54;\n\tand.b16 \t%rs62, %rs61, 255;\n\tcvt.u16.u32 \t%rs63, %r55;\n\tand.b16 \t%rs64, %rs63, 255;\n\tcvt.u16.u32 \t%rs65, %r56;\n\tand.b16 \t%rs66, %rs65, 255;\n\tcvt.rn.f32.u16 \t%f123, %rs62;\n\tcvt.rn.f32.u16 \t%f124, %rs64;\n\tcvt.rn.f32.u16 \t%f125, %rs66;\n\tmul.f32 \t%f126, %f125, 0f3E59B3D0;\n\tfma.rn.f32 \t%f127, %f124, 0f3F371759, %f126;\n\tfma.rn.f32 \t%f128, %f123, 0f3D93DD98, %f127;\n\tsub.f32 \t%f129, %f123, %f128;\n\tmul.f32 \t%f130, %f129, 0f4060E0E1;\n\tdiv.rn.f32 \t%f131, %f130, 0f3FED844D;\n\tadd.f32 \t%f132, %f131, 0f44000000;\n\tmax.f32 \t%f134, %f132, %f109;\n\tmin.f32 \t%f136, %f134, %f111;\n\tcvt.rni.s32.f32 \t%r58, %f136;\n\tcvt.u16.u32 \t%rs67, %r58;\n\tshl.b16 \t%rs68, %rs67, 6;\n\tsub.f32 \t%f137, %f125, %f128;\n\tmul.f32 \t%f138, %f137, 0f4060E0E1;\n\tdiv.rn.f32 \t%f139, %f138, 0f3FC9930C;\n\tadd.f32 \t%f140, %f139, 0f44000000;\n\tmax.f32 \t%f141, %f140, %f109;\n\tmin.f32 \t%f142, %f141, %f111;\n\tcvt.rni.s32.f32 \t%r59, %f142;\n\tcvt.u16.u32 \t%rs69, %r59;\n\tshl.b16 \t%rs70, %rs69, 6;\n\tfma.rn.f32 \t%f143, %f128, 0f405BDBDC, 0f42800000;\n\tmax.f32 \t%f144, %f143, %f109;\n\tmin.f32 \t%f145, %f144, %f111;\n\tcvt.rni.s32.f32 \t%r60, %f145;\n\tcvt.u16.u32 \t%rs71, %r60;\n\tshl.b16 \t%rs72, %rs71, 6;\n\tadd.s64 \t%rd39, %rd9, 2;\n\tst.global.u16 \t[%rd9+2], %rs72;\n\tadd.s64 \t%rd41, %rd39, %rd36;\n\tst.global.u16 \t[%rd37+2], %rs68;\n\tadd.s64 \t%rd43, %rd41, %rd36;\n\tst.global.u16 \t[%rd43], %rs70;\n\n$L__BB5_9:\n\tadd.s32 \t%r6, %r2, 1;\n\tsetp.ge.s32 \t%p9, %r6, %r9;\n\t@%p9 bra \t$L__BB5_12;\n\n\tcvt.s64.s32 \t%rd44, %r6;\n\tmul.lo.s64 \t%rd45, %rd44, %rd2;\n\tadd.s64 \t%rd46, %rd1, %rd45;\n\tcvt.rn.f32.s32 \t%f4, %r6;\n\ttex.2d.v4.u32.f32 \t{%r61, %r62, %r63, %r64}, [%rd11, {%f2, %f4}];\n\tcvt.u16.u32 \t%rs73, %r61;\n\tand.b16 \t%rs74, %rs73, 255;\n\tcvt.u16.u32 \t%rs75, %r62;\n\tand.b16 \t%rs76, %rs75, 255;\n\tcvt.u16.u32 \t%rs77, %r63;\n\tand.b16 \t%rs78, %rs77, 255;\n\tcvt.rn.f32.u16 \t%f146, %rs74;\n\tcvt.rn.f32.u16 \t%f147, %rs76;\n\tcvt.rn.f32.u16 \t%f148, %rs78;\n\tmul.f32 \t%f149, %f148, 0f3E59B3D0;\n\tfma.rn.f32 \t%f150, %f147, 0f3F371759, %f149;\n\tfma.rn.f32 \t%f151, %f146, 0f3D93DD98, %f150;\n\tsub.f32 \t%f152, %f146, %f151;\n\tmul.f32 \t%f153, %f152, 0f4060E0E1;\n\tdiv.rn.f32 \t%f154, %f153, 0f3FED844D;\n\tadd.f32 \t%f155, %f154, 0f44000000;\n\tmov.f32 \t%f156, 0f00000000;\n\tmax.f32 \t%f157, %f155, %f156;\n\tmov.f32 \t%f158, 0f447FC000;\n\tmin.f32 \t%f159, %f157, %f158;\n\tcvt.rni.s32.f32 \t%r65, %f159;\n\tcvt.u16.u32 \t%rs79, %r65;\n\tshl.b16 \t%rs80, %rs79, 6;\n\tsub.f32 \t%f160, %f148, %f151;\n\tmul.f32 \t%f161, %f160, 0f4060E0E1;\n\tdiv.rn.f32 \t%f162, %f161, 0f3FC9930C;\n\tadd.f32 \t%f163, %f162, 0f44000000;\n\tmax.f32 \t%f164, %f163, %f156;\n\tmin.f32 \t%f165, %f164, %f158;\n\tcvt.rni.s32.f32 \t%r66, %f165;\n\tcvt.u16.u32 \t%rs81, %r66;\n\tshl.b16 \t%rs82, %rs81, 6;\n\tfma.rn.f32 \t%f166, %f151, 0f405BDBDC, 0f42800000;\n\tmax.f32 \t%f167, %f166, %f156;\n\tmin.f32 \t%f168, %f167, %f158;\n\tcvt.rni.s32.f32 \t%r67, %f168;\n\tcvt.u16.u32 \t%rs83, %r67;\n\tshl.b16 \t%rs84, %rs83, 6;\n\tshl.b64 \t%rd47, %rd8, 1;\n\tadd.s64 \t%rd10, %rd46, %rd47;\n\tst.global.u16 \t[%rd10], %rs84;\n\tadd.s64 \t%rd49, %rd10, %rd36;\n\tst.global.u16 \t[%rd49], %rs80;\n\tadd.s64 \t%rd50, %rd49, %rd36;\n\tst.global.u16 \t[%rd50], %rs82;\n\t@%p8 bra \t$L__BB5_12;\n\n\tcvt.rn.f32.s32 \t%f169, %r5;\n\ttex.2d.v4.u32.f32 \t{%r68, %r69, %r70, %r71}, [%rd11, {%f169, %f4}];\n\tcvt.u16.u32 \t%rs85, %r68;\n\tand.b16 \t%rs86, %rs85, 255;\n\tcvt.u16.u32 \t%rs87, %r69;\n\tand.b16 \t%rs88, %rs87, 255;\n\tcvt.u16.u32 \t%rs89, %r70;\n\tand.b16 \t%rs90, %rs89, 255;\n\tcvt.rn.f32.u16 \t%f170, %rs86;\n\tcvt.rn.f32.u16 \t%f171, %rs88;\n\tcvt.rn.f32.u16 \t%f172, %rs90;\n\tmul.f32 \t%f173, %f172, 0f3E59B3D0;\n\tfma.rn.f32 \t%f174, %f171, 0f3F371759, %f173;\n\tfma.rn.f32 \t%f175, %f170, 0f3D93DD98, %f174;\n\tsub.f32 \t%f176, %f170, %f175;\n\tmul.f32 \t%f177, %f176, 0f4060E0E1;\n\tdiv.rn.f32 \t%f178, %f177, 0f3FED844D;\n\tadd.f32 \t%f179, %f178, 0f44000000;\n\tmax.f32 \t%f181, %f179, %f156;\n\tmin.f32 \t%f183, %f181, %f158;\n\tcvt.rni.s32.f32 \t%r72, %f183;\n\tcvt.u16.u32 \t%rs91, %r72;\n\tshl.b16 \t%rs92, %rs91, 6;\n\tsub.f32 \t%f184, %f172, %f175;\n\tmul.f32 \t%f185, %f184, 0f4060E0E1;\n\tdiv.rn.f32 \t%f186, %f185, 0f3FC9930C;\n\tadd.f32 \t%f187, %f186, 0f44000000;\n\tmax.f32 \t%f188, %f187, %f156;\n\tmin.f32 \t%f189, %f188, %f158;\n\tcvt.rni.s32.f32 \t%r73, %f189;\n\tcvt.u16.u32 \t%rs93, %r73;\n\tshl.b16 \t%rs94, %rs93, 6;\n\tfma.rn.f32 \t%f190, %f175, 0f405BDBDC, 0f42800000;\n\tmax.f32 \t%f191, %f190, %f156;\n\tmin.f32 \t%f192, %f191, %f158;\n\tcvt.rni.s32.f32 \t%r74, %f192;\n\tcvt.u16.u32 \t%rs95, %r74;\n\tshl.b16 \t%rs96, %rs95, 6;\n\tadd.s64 \t%rd51, %rd10, 2;\n\tst.global.u16 \t[%rd10+2], %rs96;\n\tadd.s64 \t%rd53, %rd51, %rd36;\n\tst.global.u16 \t[%rd49+2], %rs92;\n\tadd.s64 \t%rd55, %rd53, %rd36;\n\tst.global.u16 \t[%rd55], %rs94;\n\n$L__BB5_12:\n\tret;\n\n}\n\t// .globl\tband_hash\n.visible .entry band_hash(\n\t.param .u64 band_hash_param_0,\n\t.param .u32 band_hash_param_1,\n\t.param .u32 band_hash_param_2,\n\t.param .u32 band_hash_param_3,\n\t.param .u32 band_hash_param_4,\n\t.param .u64 band_hash_param_5\n)\n{\n\t.reg .pred \t%p<18>;\n\t.reg .b32 \t%r<42>;\n\t.reg .b64 \t%rd<172>;\n\t// demoted variable\n\t.shared .align 8 .b8 _ZZ9band_hashE4sums[2048];\n\n\tld.param.u64 \t%rd23, [band_hash_param_0];\n\tld.param.u32 \t%r19, [band_hash_param_1];\n\tld.param.u32 \t%r20, [band_hash_param_2];\n\tld.param.u32 \t%r21, [band_hash_param_3];\n\tld.param.u32 \t%r22, [band_hash_param_4];\n\tld.param.u64 \t%rd24, [band_hash_param_5];\n\tmov.u32 \t%r1, %ctaid.x;\n\tmul.lo.s32 \t%r2, %r1, %r22;\n\tadd.s32 \t%r23, %r2, %r22;\n\tmin.s32 \t%r3, %r23, %r21;\n\tmov.u32 \t%r4, %nctaid.y;\n\tshl.b32 \t%r5, %r4, 8;\n\tsetp.lt.s32 \t%p1, %r2, %r3;\n\t@%p1 bra \t$L__BB6_2;\n\tbra.uni \t$L__BB6_1;\n\n$L__BB6_2:\n\tcvt.s64.s32 \t%rd1, %r19;\n\tcvt.s64.s32 \t%rd2, %r20;\n\tmov.u32 \t%r24, %ctaid.y;\n\tshl.b32 \t%r25, %r24, 8;\n\tmov.u32 \t%r26, %tid.x;\n\tadd.s32 \t%r6, %r25, %r26;\n\tnot.b32 \t%r27, %r26;\n\tadd.s32 \t%r28, %r27, %r20;\n\tsub.s32 \t%r29, %r28, %r25;\n\tdiv.u32 \t%r7, %r29, %r5;\n\tadd.s32 \t%r30, %r7, 1;\n\tand.b32 \t%r8, %r30, 3;\n\tcvt.s64.s32 \t%rd3, %r6;\n\tadd.s32 \t%r9, %r6, %r5;\n\tcvt.s64.s32 \t%rd4, %r9;\n\tadd.s32 \t%r10, %r9, %r5;\n\tcvt.s64.s32 \t%rd5, %r10;\n\tadd.s32 \t%r11, %r10, %r5;\n\tmul.wide.s32 \t%rd6, %r5, 4;\n\tcvta.to.global.u64 \t%rd7, %rd23;\n\tmov.u64 \t%rd170, 0;\n\tmov.u32 \t%r39, %r2;\n\n$L__BB6_3:\n\tsub.s32 \t%r31, %r39, %r2;\n\tcvt.s64.s32 \t%rd27, %r31;\n\tmul.lo.s64 \t%rd9, %rd27, %rd2;\n\tsetp.ge.s32 \t%p2, %r6, %r20;\n\t@%p2 bra \t$L__BB6_10;\n\n\tcvt.s64.s32 \t%rd29, %r39;\n\tmul.lo.s64 \t%rd30, %rd29, %rd1;\n\tadd.s64 \t%rd10, %rd7, %rd30;\n\tsetp.eq.s32 \t%p3, %r8, 0;\n\tmov.u32 \t%r40, %r6;\n\t@%p3 bra \t$L__BB6_8;\n\n\tadd.s64 \t%rd31, %rd9, %rd3;\n\tshl.b64 \t%rd32, %rd3, 2;\n\tadd.s64 \t%rd11, %rd10, %rd32;\n\tld.global.u32 \t%rd33, [%rd11];\n\tand.b64 \t%rd34, %rd33, 16777215;\n\tbfi.b64 \t%rd35, %rd31, %rd34, 24, 40;\n\tshr.u64 \t%rd36, %rd31, 6;\n\tand.b64 \t%rd37, %rd36, 17179869183;\n\txor.b64 \t%rd38, %rd35, %rd37;\n\tmul.lo.s64 \t%rd39, %rd38, -4658895280553007687;\n\tshr.u64 \t%rd40, %rd39, 27;\n\txor.b64 \t%rd41, %rd40, %rd39;\n\tmul.lo.s64 \t%rd42, %rd41, -7723592293110705685;\n\tshr.u64 \t%rd43, %rd42, 31;\n\txor.b64 \t%rd44, %rd43, %rd42;\n\tadd.s64 \t%rd170, %rd44, %rd170;\n\tsetp.eq.s32 \t%p4, %r8, 1;\n\tmov.u32 \t%r40, %r9;\n\t@%p4 bra \t$L__BB6_8;\n\n\tsetp.eq.s32 \t%p5, %r8, 2;\n\tadd.s64 \t%rd45, %rd9, %rd4;\n\tcvt.s64.s32 \t%rd13, %r5;\n\tadd.s64 \t%rd14, %rd11, %rd6;\n\tld.global.u32 \t%rd47, [%rd14];\n\tand.b64 \t%rd48, %rd47, 16777215;\n\tbfi.b64 \t%rd49, %rd45, %rd48, 24, 40;\n\tshr.u64 \t%rd50, %rd45, 6;\n\tand.b64 \t%rd51, %rd50, 17179869183;\n\txor.b64 \t%rd52, %rd49, %rd51;\n\tmul.lo.s64 \t%rd53, %rd52, -4658895280553007687;\n\tshr.u64 \t%rd54, %rd53, 27;\n\txor.b64 \t%rd55, %rd54, %rd53;\n\tmul.lo.s64 \t%rd56, %rd55, -7723592293110705685;\n\tshr.u64 \t%rd57, %rd56, 31;\n\txor.b64 \t%rd58, %rd57, %rd56;\n\tadd.s64 \t%rd170, %rd58, %rd170;\n\tmov.u32 \t%r40, %r10;\n\t@%p5 bra \t$L__BB6_8;\n\n\tadd.s64 \t%rd59, %rd9, %rd5;\n\tshl.b64 \t%rd60, %rd13, 2;\n\tadd.s64 \t%rd61, %rd14, %rd60;\n\tld.global.u32 \t%rd62, [%rd61];\n\tand.b64 \t%rd63, %rd62, 16777215;\n\tbfi.b64 \t%rd64, %rd59, %rd63, 24, 40;\n\tshr.u64 \t%rd65, %rd59, 6;\n\tand.b64 \t%rd66, %rd65, 17179869183;\n\txor.b64 \t%rd67, %rd64, %rd66;\n\tmul.lo.s64 \t%rd68, %rd67, -4658895280553007687;\n\tshr.u64 \t%rd69, %rd68, 27;\n\txor.b64 \t%rd70, %rd69, %rd68;\n\tmul.lo.s64 \t%rd71, %rd70, -7723592293110705685;\n\tshr.u64 \t%rd72, %rd71, 31;\n\txor.b64 \t%rd73, %rd72, %rd71;\n\tadd.s64 \t%rd170, %rd73, %rd170;\n\tmov.u32 \t%r40, %r11;\n\n$L__BB6_8:\n\tsetp.lt.u32 \t%p6, %r7, 3;\n\t@%p6 bra \t$L__BB6_10;\n\n$L__BB6_9:\n\tcvt.s64.s32 \t%rd74, %r40;\n\tadd.s64 \t%rd75, %rd9, %rd74;\n\tmul.wide.s32 \t%rd76, %r40, 4;\n\tadd.s64 \t%rd77, %rd10, %rd76;\n\tld.global.u32 \t%rd78, [%rd77];\n\tand.b64 \t%rd79, %rd78, 16777215;\n\tbfi.b64 \t%rd80, %rd75, %rd79, 24, 40;\n\tshr.u64 \t%rd81, %rd75, 6;\n\tand.b64 \t%rd82, %rd81, 17179869183;\n\txor.b64 \t%rd83, %rd80, %rd82;\n\tmul.lo.s64 \t%rd84, %rd83, -4658895280553007687;\n\tshr.u64 \t%rd85, %rd84, 27;\n\txor.b64 \t%rd86, %rd85, %rd84;\n\tmul.lo.s64 \t%rd87, %rd86, -7723592293110705685;\n\tshr.u64 \t%rd88, %rd87, 31;\n\txor.b64 \t%rd89, %rd88, %rd87;\n\tadd.s64 \t%rd90, %rd89, %rd170;\n\tadd.s32 \t%r32, %r40, %r5;\n\tcvt.s64.s32 \t%rd91, %r32;\n\tadd.s64 \t%rd92, %rd9, %rd91;\n\tadd.s64 \t%rd93, %rd77, %rd6;\n\tld.global.u32 \t%rd94, [%rd93];\n\tand.b64 \t%rd95, %rd94, 16777215;\n\tbfi.b64 \t%rd96, %rd92, %rd95, 24, 40;\n\tshr.u64 \t%rd97, %rd92, 6;\n\tand.b64 \t%rd98, %rd97, 17179869183;\n\txor.b64 \t%rd99, %rd96, %rd98;\n\tmul.lo.s64 \t%rd100, %rd99, -4658895280553007687;\n\tshr.u64 \t%rd101, %rd100, 27;\n\txor.b64 \t%rd102, %rd101, %rd100;\n\tmul.lo.s64 \t%rd103, %rd102, -7723592293110705685;\n\tshr.u64 \t%rd104, %rd103, 31;\n\txor.b64 \t%rd105, %rd104, %rd103;\n\tadd.s64 \t%rd106, %rd105, %rd90;\n\tadd.s32 \t%r33, %r32, %r5;\n\tcvt.s64.s32 \t%rd107, %r33;\n\tadd.s64 \t%rd108, %rd9, %rd107;\n\tadd.s64 \t%rd109, %rd93, %rd6;\n\tld.global.u32 \t%rd110, [%rd109];\n\tand.b64 \t%rd111, %rd110, 16777215;\n\tbfi.b64 \t%rd112, %rd108, %rd111, 24, 40;\n\tshr.u64 \t%rd113, %rd108, 6;\n\tand.b64 \t%rd114, %rd113, 17179869183;\n\txor.b64 \t%rd115, %rd112, %rd114;\n\tmul.lo.s64 \t%rd116, %rd115, -4658895280553007687;\n\tshr.u64 \t%rd117, %rd116, 27;\n\txor.b64 \t%rd118, %rd117, %rd116;\n\tmul.lo.s64 \t%rd119, %rd118, -7723592293110705685;\n\tshr.u64 \t%rd120, %rd119, 31;\n\txor.b64 \t%rd121, %rd120, %rd119;\n\tadd.s64 \t%rd122, %rd121, %rd106;\n\tadd.s32 \t%r34, %r33, %r5;\n\tcvt.s64.s32 \t%rd123, %r34;\n\tadd.s64 \t%rd124, %rd9, %rd123;\n\tadd.s64 \t%rd125, %rd109, %rd6;\n\tld.global.u32 \t%rd126, [%rd125];\n\tand.b64 \t%rd127, %rd126, 16777215;\n\tbfi.b64 \t%rd128, %rd124, %rd127, 24, 40;\n\tshr.u64 \t%rd129, %rd124, 6;\n\tand.b64 \t%rd130, %rd129, 17179869183;\n\txor.b64 \t%rd131, %rd128, %rd130;\n\tmul.lo.s64 \t%rd132, %rd131, -4658895280553007687;\n\tshr.u64 \t%rd133, %rd132, 27;\n\txor.b64 \t%rd134, %rd133, %rd132;\n\tmul.lo.s64 \t%rd135, %rd134, -7723592293110705685;\n\tshr.u64 \t%rd136, %rd135, 31;\n\txor.b64 \t%rd137, %rd136, %rd135;\n\tadd.s64 \t%rd170, %rd137, %rd122;\n\tadd.s32 \t%r40, %r34, %r5;\n\tsetp.lt.s32 \t%p7, %r40, %r20;\n\t@%p7 bra \t$L__BB6_9;\n\n$L__BB6_10:\n\tadd.s32 \t%r39, %r39, 1;\n\tsetp.lt.s32 \t%p8, %r39, %r3;\n\t@%p8 bra \t$L__BB6_3;\n\tbra.uni \t$L__BB6_11;\n\n$L__BB6_1:\n\tmov.u64 \t%rd170, 0;\n\n$L__BB6_11:\n\tmov.u32 \t%r17, %tid.x;\n\tshl.b32 \t%r35, %r17, 3;\n\tmov.u32 \t%r36, _ZZ9band_hashE4sums;\n\tadd.s32 \t%r18, %r36, %r35;\n\tst.shared.u64 \t[%r18], %rd170;\n\tbar.sync \t0;\n\tsetp.gt.u32 \t%p9, %r17, 127;\n\t@%p9 bra \t$L__BB6_13;\n\n\tld.shared.u64 \t%rd138, [%r18];\n\tld.shared.u64 \t%rd139, [%r18+1024];\n\tadd.s64 \t%rd140, %rd138, %rd139;\n\tst.shared.u64 \t[%r18], %rd140;\n\n$L__BB6_13:\n\tbar.sync \t0;\n\tsetp.gt.u32 \t%p10, %r17, 63;\n\t@%p10 bra \t$L__BB6_15;\n\n\tld.shared.u64 \t%rd141, [%r18];\n\tld.shared.u64 \t%rd142, [%r18+512];\n\tadd.s64 \t%rd143, %rd141, %rd142;\n\tst.shared.u64 \t[%r18], %rd143;\n\n$L__BB6_15:\n\tbar.sync \t0;\n\tsetp.gt.u32 \t%p11, %r17, 31;\n\t@%p11 bra \t$L__BB6_17;\n\n\tld.shared.u64 \t%rd144, [%r18];\n\tld.shared.u64 \t%rd145, [%r18+256];\n\tadd.s64 \t%rd146, %rd144, %rd145;\n\tst.shared.u64 \t[%r18], %rd146;\n\n$L__BB6_17:\n\tbar.sync \t0;\n\tsetp.gt.u32 \t%p12, %r17, 15;\n\t@%p12 bra \t$L__BB6_19;\n\n\tld.shared.u64 \t%rd147, [%r18];\n\tld.shared.u64 \t%rd148, [%r18+128];\n\tadd.s64 \t%rd149, %rd147, %rd148;\n\tst.shared.u64 \t[%r18], %rd149;\n\n$L__BB6_19:\n\tbar.sync \t0;\n\tsetp.gt.u32 \t%p13, %r17, 7;\n\t@%p13 bra \t$L__BB6_21;\n\n\tld.shared.u64 \t%rd150, [%r18];\n\tld.shared.u64 \t%rd151, [%r18+64];\n\tadd.s64 \t%rd152, %rd150, %rd151;\n\tst.shared.u64 \t[%r18], %rd152;\n\n$L__BB6_21:\n\tbar.sync \t0;\n\tsetp.gt.u32 \t%p14, %r17, 3;\n\t@%p14 bra \t$L__BB6_23;\n\n\tld.shared.u64 \t%rd153, [%r18];\n\tld.shared.u64 \t%rd154, [%r18+32];\n\tadd.s64 \t%rd155, %rd153, %rd154;\n\tst.shared.u64 \t[%r18], %rd155;\n\n$L__BB6_23:\n\tbar.sync \t0;\n\tsetp.gt.u32 \t%p15, %r17, 1;\n\t@%p15 bra \t$L__BB6_25;\n\n\tld.shared.u64 \t%rd156, [%r18];\n\tld.shared.u64 \t%rd157, [%r18+16];\n\tadd.s64 \t%rd158, %rd156, %rd157;\n\tst.shared.u64 \t[%r18], %rd158;\n\n$L__BB6_25:\n\tbar.sync \t0;\n\tsetp.ne.s32 \t%p16, %r17, 0;\n\t@%p16 bra \t$L__BB6_27;\n\n\tld.shared.u64 \t%rd159, [_ZZ9band_hashE4sums+8];\n\tld.shared.u64 \t%rd160, [%r18];\n\tadd.s64 \t%rd161, %rd160, %rd159;\n\tst.shared.u64 \t[%r18], %rd161;\n\n$L__BB6_27:\n\tbar.sync \t0;\n\t@%p16 bra \t$L__BB6_29;\n\n\tld.shared.u64 \t%rd162, [_ZZ9band_hashE4sums];\n\tmov.u32 \t%r37, %ctaid.y;\n\tmad.lo.s32 \t%r38, %r1, %r4, %r37;\n\tcvta.to.global.u64 \t%rd163, %rd24;\n\tmul.wide.u32 \t%rd164, %r38, 8;\n\tadd.s64 \t%rd165, %rd163, %rd164;\n\tst.global.u64 \t[%rd165], %rd162;\n\n$L__BB6_29:\n\tret;\n\n}\n\n";Expand description
The ARGB/ABGR → NV12 convert and the band hash, as PTX the driver JIT-compiles at session
open. PTX is the portable form: libcuda compiles it for whatever GPU is present, so nothing
beyond the driver NVENC already needs has to be installed, and .version 3.1/.target sm_30
keeps every NVENC-capable GPU in range.