diff options
Diffstat (limited to 'VRCSDK3Worlds/Assets/Editor/x64/Bakery/exception.ptx')
| -rw-r--r-- | VRCSDK3Worlds/Assets/Editor/x64/Bakery/exception.ptx | 934 |
1 files changed, 934 insertions, 0 deletions
diff --git a/VRCSDK3Worlds/Assets/Editor/x64/Bakery/exception.ptx b/VRCSDK3Worlds/Assets/Editor/x64/Bakery/exception.ptx new file mode 100644 index 00000000..6bf2bf19 --- /dev/null +++ b/VRCSDK3Worlds/Assets/Editor/x64/Bakery/exception.ptx @@ -0,0 +1,934 @@ +// +// Generated by NVIDIA NVVM Compiler +// +// Compiler Build ID: CL-23083092 +// Cuda compilation tools, release 9.1, V9.1.85 +// Based on LLVM 3.4svn +// + +.version 6.1 +.target sm_30 +.address_size 64 + + // .globl _Z6oxMainv +.extern .func (.param .b32 func_retval0) vprintf +( + .param .b64 vprintf_param_0, + .param .b64 vprintf_param_1 +) +; +.global .align 8 .b8 pixelID[8]; +.global .align 8 .b8 resolution[8]; +.global .align 4 .b8 normal[12]; +.global .align 4 .b8 camPos[12]; +.global .align 4 .b8 root[4]; +.global .align 4 .u32 imageEnabled; +.global .texref lightmap; +.global .align 16 .b8 tileInfo[16]; +.global .align 4 .u32 additive; +.global .align 4 .b8 _ZN21rti_internal_typeinfo7pixelIDE[8] = {82, 97, 121, 0, 8, 0, 0, 0}; +.global .align 4 .b8 _ZN21rti_internal_typeinfo10resolutionE[8] = {82, 97, 121, 0, 8, 0, 0, 0}; +.global .align 4 .b8 _ZN21rti_internal_typeinfo6normalE[8] = {82, 97, 121, 0, 12, 0, 0, 0}; +.global .align 4 .b8 _ZN21rti_internal_typeinfo6camPosE[8] = {82, 97, 121, 0, 12, 0, 0, 0}; +.global .align 4 .b8 _ZN21rti_internal_typeinfo4rootE[8] = {82, 97, 121, 0, 4, 0, 0, 0}; +.global .align 4 .b8 _ZN21rti_internal_typeinfo12imageEnabledE[8] = {82, 97, 121, 0, 4, 0, 0, 0}; +.global .align 4 .b8 _ZN21rti_internal_typeinfo8tileInfoE[8] = {82, 97, 121, 0, 16, 0, 0, 0}; +.global .align 4 .b8 _ZN21rti_internal_typeinfo8additiveE[8] = {82, 97, 121, 0, 4, 0, 0, 0}; +.global .align 8 .u64 _ZN21rti_internal_register20reg_bitness_detectorE; +.global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail0E; +.global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail1E; +.global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail2E; +.global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail3E; +.global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail4E; +.global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail5E; +.global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail6E; +.global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail7E; +.global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail8E; +.global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail9E; +.global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail0E; +.global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail1E; +.global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail2E; +.global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail3E; +.global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail4E; +.global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail5E; +.global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail6E; +.global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail7E; +.global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail8E; +.global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail9E; +.global .align 4 .u32 _ZN21rti_internal_register14reg_rayIndex_xE; +.global .align 4 .u32 _ZN21rti_internal_register14reg_rayIndex_yE; +.global .align 4 .u32 _ZN21rti_internal_register14reg_rayIndex_zE; +.global .align 8 .b8 _ZN21rti_internal_typename7pixelIDE[6] = {117, 105, 110, 116, 50, 0}; +.global .align 8 .b8 _ZN21rti_internal_typename10resolutionE[6] = {117, 105, 110, 116, 50, 0}; +.global .align 8 .b8 _ZN21rti_internal_typename6normalE[7] = {102, 108, 111, 97, 116, 51, 0}; +.global .align 8 .b8 _ZN21rti_internal_typename6camPosE[7] = {102, 108, 111, 97, 116, 51, 0}; +.global .align 16 .b8 _ZN21rti_internal_typename4rootE[9] = {114, 116, 79, 98, 106, 101, 99, 116, 0}; +.global .align 4 .b8 _ZN21rti_internal_typename12imageEnabledE[4] = {105, 110, 116, 0}; +.global .align 8 .b8 _ZN21rti_internal_typename8tileInfoE[6] = {117, 105, 110, 116, 52, 0}; +.global .align 4 .b8 _ZN21rti_internal_typename8additiveE[4] = {105, 110, 116, 0}; +.global .align 4 .u32 _ZN21rti_internal_typeenum7pixelIDE = 4919; +.global .align 4 .u32 _ZN21rti_internal_typeenum10resolutionE = 4919; +.global .align 4 .u32 _ZN21rti_internal_typeenum6normalE = 4919; +.global .align 4 .u32 _ZN21rti_internal_typeenum6camPosE = 4919; +.global .align 4 .u32 _ZN21rti_internal_typeenum4rootE = 4919; +.global .align 4 .u32 _ZN21rti_internal_typeenum12imageEnabledE = 4919; +.global .align 4 .u32 _ZN21rti_internal_typeenum8tileInfoE = 4919; +.global .align 4 .u32 _ZN21rti_internal_typeenum8additiveE = 4919; +.global .align 16 .b8 _ZN21rti_internal_semantic7pixelIDE[14] = {114, 116, 76, 97, 117, 110, 99, 104, 73, 110, 100, 101, 120, 0}; +.global .align 16 .b8 _ZN21rti_internal_semantic10resolutionE[12] = {114, 116, 76, 97, 117, 110, 99, 104, 68, 105, 109, 0}; +.global .align 16 .b8 _ZN21rti_internal_semantic6normalE[17] = {97, 116, 116, 114, 105, 98, 117, 116, 101, 32, 110, 111, 114, 109, 97, 108, 0}; +.global .align 1 .b8 _ZN21rti_internal_semantic6camPosE[1]; +.global .align 1 .b8 _ZN21rti_internal_semantic4rootE[1]; +.global .align 1 .b8 _ZN21rti_internal_semantic12imageEnabledE[1]; +.global .align 1 .b8 _ZN21rti_internal_semantic8tileInfoE[1]; +.global .align 1 .b8 _ZN21rti_internal_semantic8additiveE[1]; +.global .align 1 .b8 _ZN23rti_internal_annotation7pixelIDE[1]; +.global .align 1 .b8 _ZN23rti_internal_annotation10resolutionE[1]; +.global .align 1 .b8 _ZN23rti_internal_annotation6normalE[1]; +.global .align 1 .b8 _ZN23rti_internal_annotation6camPosE[1]; +.global .align 1 .b8 _ZN23rti_internal_annotation4rootE[1]; +.global .align 1 .b8 _ZN23rti_internal_annotation12imageEnabledE[1]; +.global .align 1 .b8 _ZN23rti_internal_annotation8tileInfoE[1]; +.global .align 1 .b8 _ZN23rti_internal_annotation8additiveE[1]; +.global .align 16 .b8 $str[64] = {67, 97, 117, 103, 104, 116, 32, 82, 84, 95, 69, 88, 67, 69, 80, 84, 73, 79, 78, 95, 83, 84, 65, 67, 75, 95, 79, 86, 69, 82, 70, 76, 79, 87, 10, 32, 32, 108, 97, 117, 110, 99, 104, 32, 105, 110, 100, 101, 120, 32, 58, 32, 37, 100, 44, 32, 37, 100, 44, 32, 37, 100, 10, 0}; +.global .align 16 .b8 $str1[218] = {67, 97, 117, 103, 104, 116, 32, 82, 84, 95, 69, 88, 67, 69, 80, 84, 73, 79, 78, 95, 66, 85, 70, 70, 69, 82, 95, 73, 78, 68, 69, 88, 95, 79, 85, 84, 95, 79, 70, 95, 66, 79, 85, 78, 68, 83, 10, 32, 32, 108, 97, 117, 110, 99, 104, 32, 105, 110, 100, 101, 120, 32, 32, 32, 58, 32, 37, 100, 44, 32, 37, 100, 44, 32, 37, 100, 10, 32, 32, 98, 117, 102, 102, 101, 114, 32, 97, 100, 100, 114, 101, 115, 115, 32, 58, 32, 48, 120, 37, 108, 108, 88, 10, 32, 32, 100, 105, 109, 101, 110, 115, 105, 111, 110, 97, 108, 105, 116, 121, 32, 58, 32, 37, 100, 10, 32, 32, 115, 105, 122, 101, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 58, 32, 37, 108, 108, 100, 120, 37, 108, 108, 100, 120, 37, 108, 108, 100, 10, 32, 32, 101, 108, 101, 109, 101, 110, 116, 32, 115, 105, 122, 101, 32, 32, 32, 58, 32, 37, 100, 10, 32, 32, 97, 99, 99, 101, 115, 115, 101, 100, 32, 105, 110, 100, 101, 120, 32, 58, 32, 37, 108, 108, 100, 44, 32, 37, 108, 108, 100, 44, 32, 37, 108, 108, 100, 10, 0}; +.global .align 16 .b8 $str2[40] = {67, 97, 117, 103, 104, 116, 32, 82, 84, 95, 69, 88, 67, 69, 80, 84, 73, 79, 78, 95, 80, 82, 79, 71, 82, 65, 77, 95, 73, 68, 95, 73, 78, 86, 65, 76, 73, 68, 10, 0}; +.global .align 16 .b8 $str3[46] = {9, 112, 114, 111, 103, 114, 97, 109, 32, 73, 68, 32, 101, 113, 117, 97, 108, 32, 116, 111, 32, 82, 84, 95, 80, 82, 79, 71, 82, 65, 77, 95, 73, 68, 95, 78, 85, 76, 76, 32, 117, 115, 101, 100, 10, 0}; +.global .align 16 .b8 $str4[56] = {9, 112, 114, 111, 103, 114, 97, 109, 32, 73, 68, 32, 40, 37, 100, 41, 32, 105, 115, 32, 110, 111, 116, 32, 105, 110, 32, 116, 104, 101, 32, 118, 97, 108, 105, 100, 32, 114, 97, 110, 103, 101, 32, 111, 102, 32, 91, 49, 44, 115, 105, 122, 101, 41, 10, 0}; +.global .align 16 .b8 $str5[39] = {9, 112, 114, 111, 103, 114, 97, 109, 32, 73, 68, 32, 111, 102, 32, 97, 32, 100, 101, 108, 101, 116, 101, 100, 32, 112, 114, 111, 103, 114, 97, 109, 32, 117, 115, 101, 100, 10, 0}; +.global .align 16 .b8 $str6[40] = {67, 97, 117, 103, 104, 116, 32, 82, 84, 95, 69, 88, 67, 69, 80, 84, 73, 79, 78, 95, 84, 69, 88, 84, 85, 82, 69, 95, 73, 68, 95, 73, 78, 86, 65, 76, 73, 68, 10, 0}; +.global .align 16 .b8 $str7[33] = {9, 116, 101, 120, 116, 117, 114, 101, 32, 73, 68, 32, 40, 37, 100, 41, 32, 105, 115, 32, 105, 110, 118, 97, 108, 105, 100, 32, 40, 48, 41, 10, 0}; +.global .align 16 .b8 $str8[56] = {9, 116, 101, 120, 116, 117, 114, 101, 32, 73, 68, 32, 40, 37, 100, 41, 32, 105, 115, 32, 110, 111, 116, 32, 105, 110, 32, 116, 104, 101, 32, 118, 97, 108, 105, 100, 32, 114, 97, 110, 103, 101, 32, 111, 102, 32, 91, 49, 44, 115, 105, 122, 101, 41, 10, 0}; +.global .align 16 .b8 $str9[34] = {9, 116, 101, 120, 116, 117, 114, 101, 32, 73, 68, 32, 40, 37, 100, 41, 32, 105, 115, 32, 105, 110, 118, 97, 108, 105, 100, 32, 40, 45, 49, 41, 10, 0}; +.global .align 16 .b8 $str10[39] = {67, 97, 117, 103, 104, 116, 32, 82, 84, 95, 69, 88, 67, 69, 80, 84, 73, 79, 78, 95, 66, 85, 70, 70, 69, 82, 95, 73, 68, 95, 73, 78, 86, 65, 76, 73, 68, 10, 0}; +.global .align 16 .b8 $str11[44] = {9, 98, 117, 102, 102, 101, 114, 32, 73, 68, 32, 101, 113, 117, 97, 108, 32, 116, 111, 32, 82, 84, 95, 66, 85, 70, 70, 69, 82, 95, 73, 68, 95, 78, 85, 76, 76, 32, 117, 115, 101, 100, 10, 0}; +.global .align 16 .b8 $str12[55] = {9, 98, 117, 102, 102, 101, 114, 32, 73, 68, 32, 40, 37, 100, 41, 32, 105, 115, 32, 110, 111, 116, 32, 105, 110, 32, 116, 104, 101, 32, 118, 97, 108, 105, 100, 32, 114, 97, 110, 103, 101, 32, 111, 102, 32, 91, 49, 44, 115, 105, 122, 101, 41, 10, 0}; +.global .align 16 .b8 $str13[37] = {9, 66, 117, 102, 102, 101, 114, 32, 73, 68, 32, 111, 102, 32, 97, 32, 100, 101, 108, 101, 116, 101, 100, 32, 98, 117, 102, 102, 101, 114, 32, 117, 115, 101, 100, 10, 0}; +.global .align 16 .b8 $str14[145] = {67, 97, 117, 103, 104, 116, 32, 82, 84, 95, 69, 88, 67, 69, 80, 84, 73, 79, 78, 95, 73, 78, 68, 69, 88, 95, 79, 85, 84, 95, 79, 70, 95, 66, 79, 85, 78, 68, 83, 10, 32, 32, 108, 97, 117, 110, 99, 104, 32, 105, 110, 100, 101, 120, 32, 32, 32, 58, 32, 37, 100, 44, 32, 37, 100, 44, 32, 37, 100, 10, 32, 32, 98, 117, 102, 102, 101, 114, 32, 97, 100, 100, 114, 101, 115, 115, 32, 58, 32, 48, 120, 37, 108, 108, 88, 10, 32, 32, 115, 105, 122, 101, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 58, 32, 37, 108, 108, 100, 10, 32, 32, 97, 99, 99, 101, 115, 115, 101, 100, 32, 105, 110, 100, 101, 120, 32, 58, 32, 37, 108, 108, 100, 10, 0}; +.global .align 16 .b8 $str15[179] = {67, 97, 117, 103, 104, 116, 32, 82, 84, 95, 69, 88, 67, 69, 80, 84, 73, 79, 78, 95, 73, 78, 86, 65, 76, 73, 68, 95, 82, 65, 89, 10, 32, 32, 108, 97, 117, 110, 99, 104, 32, 105, 110, 100, 101, 120, 32, 32, 58, 32, 37, 100, 44, 32, 37, 100, 44, 32, 37, 100, 10, 32, 32, 114, 97, 121, 32, 111, 114, 105, 103, 105, 110, 32, 32, 32, 32, 58, 32, 37, 102, 32, 37, 102, 32, 37, 102, 10, 32, 32, 114, 97, 121, 32, 100, 105, 114, 101, 99, 116, 105, 111, 110, 32, 58, 32, 37, 102, 32, 37, 102, 32, 37, 102, 10, 32, 32, 114, 97, 121, 32, 116, 121, 112, 101, 32, 32, 32, 32, 32, 32, 58, 32, 37, 100, 10, 32, 32, 114, 97, 121, 32, 116, 109, 105, 110, 32, 32, 32, 32, 32, 32, 58, 32, 37, 102, 10, 32, 32, 114, 97, 121, 32, 116, 109, 97, 120, 32, 32, 32, 32, 32, 32, 58, 32, 37, 102, 10, 0}; +.global .align 16 .b8 $str16[84] = {67, 97, 117, 103, 104, 116, 32, 82, 84, 95, 69, 88, 67, 69, 80, 84, 73, 79, 78, 95, 73, 78, 84, 69, 82, 78, 65, 76, 95, 69, 82, 82, 79, 82, 10, 32, 32, 108, 97, 117, 110, 99, 104, 32, 105, 110, 100, 101, 120, 32, 58, 32, 37, 100, 44, 32, 37, 100, 44, 32, 37, 100, 10, 32, 32, 101, 114, 114, 111, 114, 32, 105, 100, 32, 32, 32, 32, 32, 58, 32, 37, 100, 10, 0}; +.global .align 16 .b8 $str17[57] = {67, 97, 117, 103, 104, 116, 32, 82, 84, 95, 69, 88, 67, 69, 80, 84, 73, 79, 78, 95, 85, 83, 69, 82, 43, 37, 100, 10, 32, 32, 108, 97, 117, 110, 99, 104, 32, 105, 110, 100, 101, 120, 32, 58, 32, 37, 100, 44, 32, 37, 100, 44, 32, 37, 100, 10, 0}; +.global .align 16 .b8 $str18[54] = {67, 97, 117, 103, 104, 116, 32, 117, 110, 107, 110, 111, 119, 110, 32, 101, 120, 99, 101, 112, 116, 105, 111, 110, 10, 32, 32, 108, 97, 117, 110, 99, 104, 32, 105, 110, 100, 101, 120, 32, 58, 32, 37, 100, 44, 32, 37, 100, 44, 32, 37, 100, 10, 0}; + +.visible .entry _Z6oxMainv( + +) +{ + .local .align 16 .b8 __local_depot0[208]; + .reg .b64 %SP; + .reg .b64 %SPL; + .reg .pred %p<40>; + .reg .f32 %f<9>; + .reg .b32 %r<84>; + .reg .f64 %fd<9>; + .reg .b64 %rd<90>; + + + mov.u64 %rd89, __local_depot0; + cvta.local.u64 %SP, %rd89; + // inline asm + call (%r39), _rt_get_exception_code, (); + // inline asm + // inline asm + call (%r40), _rt_get_exception_code, (); + // inline asm + setp.eq.s32 %p1, %r40, 1020; + @%p1 bra BB0_58; + bra.uni BB0_1; + +BB0_58: + ld.volatile.global.u32 %r36, [_ZN21rti_internal_register14reg_rayIndex_xE]; + ld.volatile.global.u32 %r37, [_ZN21rti_internal_register14reg_rayIndex_yE]; + ld.volatile.global.u32 %r38, [_ZN21rti_internal_register14reg_rayIndex_zE]; + // inline asm + call (%r81), _rt_print_active, (); + // inline asm + setp.eq.s32 %p39, %r81, 0; + @%p39 bra BB0_60; + + add.u64 %rd83, %SP, 184; + cvta.to.local.u64 %rd84, %rd83; + st.local.v2.u32 [%rd84], {%r36, %r37}; + st.local.u32 [%rd84+8], %r38; + mov.u64 %rd85, $str; + cvta.global.u64 %rd86, %rd85; + // Callseq Start 18 + { + .reg .b32 temp_param_reg; + // <end>} + .param .b64 param0; + st.param.b64 [param0+0], %rd86; + .param .b64 param1; + st.param.b64 [param1+0], %rd83; + .param .b32 retval0; + call.uni (retval0), + vprintf, + ( + param0, + param1 + ); + ld.param.b32 %r82, [retval0+0]; + + //{ + }// Callseq End 18 + bra.uni BB0_60; + +BB0_1: + setp.eq.s32 %p2, %r40, 1021; + @%p2 bra BB0_52; + bra.uni BB0_2; + +BB0_52: + ld.volatile.global.u32 %r30, [_ZN21rti_internal_register21reg_exception_detail0E]; + ld.volatile.global.u32 %r31, [_ZN21rti_internal_register14reg_rayIndex_xE]; + ld.volatile.global.u32 %r32, [_ZN21rti_internal_register14reg_rayIndex_yE]; + ld.volatile.global.u32 %r33, [_ZN21rti_internal_register14reg_rayIndex_zE]; + ld.volatile.global.u64 %rd5, [_ZN21rti_internal_register24reg_exception_64_detail0E]; + ld.volatile.global.u32 %r34, [_ZN21rti_internal_register21reg_exception_detail0E]; + ld.volatile.global.u64 %rd6, [_ZN21rti_internal_register24reg_exception_64_detail1E]; + mov.u64 %rd88, 1; + setp.lt.u32 %p36, %r30, 2; + mov.u64 %rd87, %rd88; + @%p36 bra BB0_54; + + ld.volatile.global.u64 %rd87, [_ZN21rti_internal_register24reg_exception_64_detail2E]; + +BB0_54: + setp.lt.u32 %p37, %r30, 3; + @%p37 bra BB0_56; + + ld.volatile.global.u64 %rd88, [_ZN21rti_internal_register24reg_exception_64_detail3E]; + +BB0_56: + ld.volatile.global.u32 %r35, [_ZN21rti_internal_register21reg_exception_detail1E]; + ld.volatile.global.u64 %rd11, [_ZN21rti_internal_register24reg_exception_64_detail4E]; + ld.volatile.global.u64 %rd12, [_ZN21rti_internal_register24reg_exception_64_detail5E]; + ld.volatile.global.u64 %rd13, [_ZN21rti_internal_register24reg_exception_64_detail6E]; + // inline asm + call (%r79), _rt_print_active, (); + // inline asm + setp.eq.s32 %p38, %r79, 0; + @%p38 bra BB0_60; + + add.u64 %rd79, %SP, 96; + cvta.to.local.u64 %rd80, %rd79; + st.local.v2.u32 [%rd80], {%r31, %r32}; + st.local.u32 [%rd80+8], %r33; + st.local.u32 [%rd80+24], %r34; + st.local.u32 [%rd80+56], %r35; + st.local.u64 [%rd80+16], %rd5; + st.local.u64 [%rd80+32], %rd6; + st.local.u64 [%rd80+40], %rd87; + st.local.u64 [%rd80+48], %rd88; + st.local.u64 [%rd80+64], %rd11; + st.local.u64 [%rd80+72], %rd12; + st.local.u64 [%rd80+80], %rd13; + mov.u64 %rd81, $str1; + cvta.global.u64 %rd82, %rd81; + // Callseq Start 17 + { + .reg .b32 temp_param_reg; + // <end>} + .param .b64 param0; + st.param.b64 [param0+0], %rd82; + .param .b64 param1; + st.param.b64 [param1+0], %rd79; + .param .b32 retval0; + call.uni (retval0), + vprintf, + ( + param0, + param1 + ); + ld.param.b32 %r80, [retval0+0]; + + //{ + }// Callseq End 17 + +BB0_60: + ret; + +BB0_2: + setp.eq.s32 %p3, %r40, 1006; + @%p3 bra BB0_41; + bra.uni BB0_3; + +BB0_41: + // inline asm + call (%r70), _rt_print_active, (); + // inline asm + setp.eq.s32 %p29, %r70, 0; + @%p29 bra BB0_43; + + mov.u64 %rd64, $str2; + cvta.global.u64 %rd65, %rd64; + mov.u64 %rd66, 0; + // Callseq Start 13 + { + .reg .b32 temp_param_reg; + // <end>} + .param .b64 param0; + st.param.b64 [param0+0], %rd65; + .param .b64 param1; + st.param.b64 [param1+0], %rd66; + .param .b32 retval0; + call.uni (retval0), + vprintf, + ( + param0, + param1 + ); + ld.param.b32 %r71, [retval0+0]; + + //{ + }// Callseq End 13 + +BB0_43: + ld.volatile.global.u32 %r72, [_ZN21rti_internal_register21reg_exception_detail1E]; + setp.eq.s32 %p30, %r72, 0; + @%p30 bra BB0_50; + + setp.eq.s32 %p31, %r72, 1; + @%p31 bra BB0_48; + bra.uni BB0_45; + +BB0_48: + ld.volatile.global.u32 %r29, [_ZN21rti_internal_register21reg_exception_detail0E]; + // inline asm + call (%r75), _rt_print_active, (); + // inline asm + setp.eq.s32 %p34, %r75, 0; + @%p34 bra BB0_60; + + add.u64 %rd70, %SP, 88; + cvta.to.local.u64 %rd71, %rd70; + st.local.u32 [%rd71], %r29; + mov.u64 %rd72, $str4; + cvta.global.u64 %rd73, %rd72; + // Callseq Start 15 + { + .reg .b32 temp_param_reg; + // <end>} + .param .b64 param0; + st.param.b64 [param0+0], %rd73; + .param .b64 param1; + st.param.b64 [param1+0], %rd70; + .param .b32 retval0; + call.uni (retval0), + vprintf, + ( + param0, + param1 + ); + ld.param.b32 %r76, [retval0+0]; + + //{ + }// Callseq End 15 + bra.uni BB0_60; + +BB0_3: + setp.eq.s32 %p4, %r40, 1007; + @%p4 bra BB0_30; + bra.uni BB0_4; + +BB0_30: + // inline asm + call (%r61), _rt_print_active, (); + // inline asm + setp.eq.s32 %p22, %r61, 0; + @%p22 bra BB0_32; + + mov.u64 %rd49, $str6; + cvta.global.u64 %rd50, %rd49; + mov.u64 %rd51, 0; + // Callseq Start 9 + { + .reg .b32 temp_param_reg; + // <end>} + .param .b64 param0; + st.param.b64 [param0+0], %rd50; + .param .b64 param1; + st.param.b64 [param1+0], %rd51; + .param .b32 retval0; + call.uni (retval0), + vprintf, + ( + param0, + param1 + ); + ld.param.b32 %r62, [retval0+0]; + + //{ + }// Callseq End 9 + +BB0_32: + ld.volatile.global.u32 %r63, [_ZN21rti_internal_register21reg_exception_detail1E]; + setp.eq.s32 %p23, %r63, 0; + @%p23 bra BB0_39; + + setp.eq.s32 %p24, %r63, 1; + @%p24 bra BB0_37; + bra.uni BB0_34; + +BB0_37: + ld.volatile.global.u32 %r27, [_ZN21rti_internal_register21reg_exception_detail0E]; + // inline asm + call (%r66), _rt_print_active, (); + // inline asm + setp.eq.s32 %p27, %r66, 0; + @%p27 bra BB0_60; + + add.u64 %rd56, %SP, 72; + cvta.to.local.u64 %rd57, %rd56; + st.local.u32 [%rd57], %r27; + mov.u64 %rd58, $str8; + cvta.global.u64 %rd59, %rd58; + // Callseq Start 11 + { + .reg .b32 temp_param_reg; + // <end>} + .param .b64 param0; + st.param.b64 [param0+0], %rd59; + .param .b64 param1; + st.param.b64 [param1+0], %rd56; + .param .b32 retval0; + call.uni (retval0), + vprintf, + ( + param0, + param1 + ); + ld.param.b32 %r67, [retval0+0]; + + //{ + }// Callseq End 11 + bra.uni BB0_60; + +BB0_4: + setp.eq.s32 %p5, %r40, 1018; + @%p5 bra BB0_19; + bra.uni BB0_5; + +BB0_19: + // inline asm + call (%r52), _rt_print_active, (); + // inline asm + setp.eq.s32 %p15, %r52, 0; + @%p15 bra BB0_21; + + mov.u64 %rd36, $str10; + cvta.global.u64 %rd37, %rd36; + mov.u64 %rd38, 0; + // Callseq Start 5 + { + .reg .b32 temp_param_reg; + // <end>} + .param .b64 param0; + st.param.b64 [param0+0], %rd37; + .param .b64 param1; + st.param.b64 [param1+0], %rd38; + .param .b32 retval0; + call.uni (retval0), + vprintf, + ( + param0, + param1 + ); + ld.param.b32 %r53, [retval0+0]; + + //{ + }// Callseq End 5 + +BB0_21: + ld.volatile.global.u32 %r54, [_ZN21rti_internal_register21reg_exception_detail1E]; + setp.eq.s32 %p16, %r54, 0; + @%p16 bra BB0_28; + + setp.eq.s32 %p17, %r54, 1; + @%p17 bra BB0_26; + bra.uni BB0_23; + +BB0_26: + ld.volatile.global.u32 %r25, [_ZN21rti_internal_register21reg_exception_detail0E]; + // inline asm + call (%r57), _rt_print_active, (); + // inline asm + setp.eq.s32 %p20, %r57, 0; + @%p20 bra BB0_60; + + add.u64 %rd42, %SP, 56; + cvta.to.local.u64 %rd43, %rd42; + st.local.u32 [%rd43], %r25; + mov.u64 %rd44, $str12; + cvta.global.u64 %rd45, %rd44; + // Callseq Start 7 + { + .reg .b32 temp_param_reg; + // <end>} + .param .b64 param0; + st.param.b64 [param0+0], %rd45; + .param .b64 param1; + st.param.b64 [param1+0], %rd42; + .param .b32 retval0; + call.uni (retval0), + vprintf, + ( + param0, + param1 + ); + ld.param.b32 %r58, [retval0+0]; + + //{ + }// Callseq End 7 + bra.uni BB0_60; + +BB0_50: + // inline asm + call (%r77), _rt_print_active, (); + // inline asm + setp.eq.s32 %p35, %r77, 0; + @%p35 bra BB0_60; + + mov.u64 %rd74, $str3; + cvta.global.u64 %rd75, %rd74; + mov.u64 %rd76, 0; + // Callseq Start 16 + { + .reg .b32 temp_param_reg; + // <end>} + .param .b64 param0; + st.param.b64 [param0+0], %rd75; + .param .b64 param1; + st.param.b64 [param1+0], %rd76; + .param .b32 retval0; + call.uni (retval0), + vprintf, + ( + param0, + param1 + ); + ld.param.b32 %r78, [retval0+0]; + + //{ + }// Callseq End 16 + bra.uni BB0_60; + +BB0_45: + setp.ne.s32 %p32, %r72, 2; + @%p32 bra BB0_60; + + // inline asm + call (%r73), _rt_print_active, (); + // inline asm + setp.eq.s32 %p33, %r73, 0; + @%p33 bra BB0_60; + + mov.u64 %rd67, $str5; + cvta.global.u64 %rd68, %rd67; + mov.u64 %rd69, 0; + // Callseq Start 14 + { + .reg .b32 temp_param_reg; + // <end>} + .param .b64 param0; + st.param.b64 [param0+0], %rd68; + .param .b64 param1; + st.param.b64 [param1+0], %rd69; + .param .b32 retval0; + call.uni (retval0), + vprintf, + ( + param0, + param1 + ); + ld.param.b32 %r74, [retval0+0]; + + //{ + }// Callseq End 14 + bra.uni BB0_60; + +BB0_5: + setp.eq.s32 %p6, %r40, 1019; + @%p6 bra BB0_17; + bra.uni BB0_6; + +BB0_17: + ld.volatile.global.u32 %r50, [_ZN21rti_internal_register21reg_exception_detail0E]; + ld.volatile.global.u32 %r22, [_ZN21rti_internal_register14reg_rayIndex_xE]; + ld.volatile.global.u32 %r23, [_ZN21rti_internal_register14reg_rayIndex_yE]; + ld.volatile.global.u32 %r24, [_ZN21rti_internal_register14reg_rayIndex_zE]; + ld.volatile.global.u64 %rd2, [_ZN21rti_internal_register24reg_exception_64_detail0E]; + ld.volatile.global.u64 %rd3, [_ZN21rti_internal_register24reg_exception_64_detail1E]; + ld.volatile.global.u64 %rd4, [_ZN21rti_internal_register24reg_exception_64_detail2E]; + // inline asm + call (%r49), _rt_print_active, (); + // inline asm + setp.eq.s32 %p14, %r49, 0; + @%p14 bra BB0_60; + + add.u64 %rd32, %SP, 16; + cvta.to.local.u64 %rd33, %rd32; + st.local.v2.u32 [%rd33], {%r22, %r23}; + st.local.u32 [%rd33+8], %r24; + st.local.u64 [%rd33+16], %rd2; + st.local.u64 [%rd33+24], %rd3; + st.local.u64 [%rd33+32], %rd4; + mov.u64 %rd34, $str14; + cvta.global.u64 %rd35, %rd34; + // Callseq Start 4 + { + .reg .b32 temp_param_reg; + // <end>} + .param .b64 param0; + st.param.b64 [param0+0], %rd35; + .param .b64 param1; + st.param.b64 [param1+0], %rd32; + .param .b32 retval0; + call.uni (retval0), + vprintf, + ( + param0, + param1 + ); + ld.param.b32 %r51, [retval0+0]; + + //{ + }// Callseq End 4 + bra.uni BB0_60; + +BB0_39: + ld.volatile.global.u32 %r28, [_ZN21rti_internal_register21reg_exception_detail0E]; + // inline asm + call (%r68), _rt_print_active, (); + // inline asm + setp.eq.s32 %p28, %r68, 0; + @%p28 bra BB0_60; + + add.u64 %rd60, %SP, 80; + cvta.to.local.u64 %rd61, %rd60; + st.local.u32 [%rd61], %r28; + mov.u64 %rd62, $str7; + cvta.global.u64 %rd63, %rd62; + // Callseq Start 12 + { + .reg .b32 temp_param_reg; + // <end>} + .param .b64 param0; + st.param.b64 [param0+0], %rd63; + .param .b64 param1; + st.param.b64 [param1+0], %rd60; + .param .b32 retval0; + call.uni (retval0), + vprintf, + ( + param0, + param1 + ); + ld.param.b32 %r69, [retval0+0]; + + //{ + }// Callseq End 12 + bra.uni BB0_60; + +BB0_34: + setp.ne.s32 %p25, %r63, 2; + @%p25 bra BB0_60; + + ld.volatile.global.u32 %r26, [_ZN21rti_internal_register21reg_exception_detail0E]; + // inline asm + call (%r64), _rt_print_active, (); + // inline asm + setp.eq.s32 %p26, %r64, 0; + @%p26 bra BB0_60; + + add.u64 %rd52, %SP, 64; + cvta.to.local.u64 %rd53, %rd52; + st.local.u32 [%rd53], %r26; + mov.u64 %rd54, $str9; + cvta.global.u64 %rd55, %rd54; + // Callseq Start 10 + { + .reg .b32 temp_param_reg; + // <end>} + .param .b64 param0; + st.param.b64 [param0+0], %rd55; + .param .b64 param1; + st.param.b64 [param1+0], %rd52; + .param .b32 retval0; + call.uni (retval0), + vprintf, + ( + param0, + param1 + ); + ld.param.b32 %r65, [retval0+0]; + + //{ + }// Callseq End 10 + bra.uni BB0_60; + +BB0_6: + setp.eq.s32 %p7, %r40, 1022; + @%p7 bra BB0_15; + bra.uni BB0_7; + +BB0_15: + ld.volatile.global.u32 %r10, [_ZN21rti_internal_register14reg_rayIndex_xE]; + ld.volatile.global.u32 %r11, [_ZN21rti_internal_register14reg_rayIndex_yE]; + ld.volatile.global.u32 %r12, [_ZN21rti_internal_register14reg_rayIndex_zE]; + ld.volatile.global.u32 %r13, [_ZN21rti_internal_register21reg_exception_detail0E]; + ld.volatile.global.u32 %r14, [_ZN21rti_internal_register21reg_exception_detail1E]; + ld.volatile.global.u32 %r15, [_ZN21rti_internal_register21reg_exception_detail2E]; + ld.volatile.global.u32 %r16, [_ZN21rti_internal_register21reg_exception_detail3E]; + ld.volatile.global.u32 %r17, [_ZN21rti_internal_register21reg_exception_detail4E]; + ld.volatile.global.u32 %r18, [_ZN21rti_internal_register21reg_exception_detail5E]; + ld.volatile.global.u32 %r19, [_ZN21rti_internal_register21reg_exception_detail6E]; + ld.volatile.global.u32 %r20, [_ZN21rti_internal_register21reg_exception_detail7E]; + ld.volatile.global.u32 %r21, [_ZN21rti_internal_register21reg_exception_detail8E]; + // inline asm + call (%r47), _rt_print_active, (); + // inline asm + setp.eq.s32 %p13, %r47, 0; + @%p13 bra BB0_60; + + mov.b32 %f1, %r13; + cvt.f64.f32 %fd1, %f1; + mov.b32 %f2, %r14; + cvt.f64.f32 %fd2, %f2; + mov.b32 %f3, %r15; + cvt.f64.f32 %fd3, %f3; + mov.b32 %f4, %r16; + cvt.f64.f32 %fd4, %f4; + mov.b32 %f5, %r17; + cvt.f64.f32 %fd5, %f5; + mov.b32 %f6, %r18; + cvt.f64.f32 %fd6, %f6; + mov.b32 %f7, %r20; + cvt.f64.f32 %fd7, %f7; + mov.b32 %f8, %r21; + cvt.f64.f32 %fd8, %f8; + add.u64 %rd28, %SP, 96; + cvta.to.local.u64 %rd29, %rd28; + st.local.v2.u32 [%rd29], {%r10, %r11}; + st.local.u32 [%rd29+8], %r12; + st.local.u32 [%rd29+64], %r19; + st.local.f64 [%rd29+16], %fd1; + st.local.f64 [%rd29+24], %fd2; + st.local.f64 [%rd29+32], %fd3; + st.local.f64 [%rd29+40], %fd4; + st.local.f64 [%rd29+48], %fd5; + st.local.f64 [%rd29+56], %fd6; + st.local.f64 [%rd29+72], %fd7; + st.local.f64 [%rd29+80], %fd8; + mov.u64 %rd30, $str15; + cvta.global.u64 %rd31, %rd30; + // Callseq Start 3 + { + .reg .b32 temp_param_reg; + // <end>} + .param .b64 param0; + st.param.b64 [param0+0], %rd31; + .param .b64 param1; + st.param.b64 [param1+0], %rd28; + .param .b32 retval0; + call.uni (retval0), + vprintf, + ( + param0, + param1 + ); + ld.param.b32 %r48, [retval0+0]; + + //{ + }// Callseq End 3 + bra.uni BB0_60; + +BB0_28: + // inline asm + call (%r59), _rt_print_active, (); + // inline asm + setp.eq.s32 %p21, %r59, 0; + @%p21 bra BB0_60; + + mov.u64 %rd46, $str11; + cvta.global.u64 %rd47, %rd46; + mov.u64 %rd48, 0; + // Callseq Start 8 + { + .reg .b32 temp_param_reg; + // <end>} + .param .b64 param0; + st.param.b64 [param0+0], %rd47; + .param .b64 param1; + st.param.b64 [param1+0], %rd48; + .param .b32 retval0; + call.uni (retval0), + vprintf, + ( + param0, + param1 + ); + ld.param.b32 %r60, [retval0+0]; + + //{ + }// Callseq End 8 + bra.uni BB0_60; + +BB0_23: + setp.ne.s32 %p18, %r54, 2; + @%p18 bra BB0_60; + + // inline asm + call (%r55), _rt_print_active, (); + // inline asm + setp.eq.s32 %p19, %r55, 0; + @%p19 bra BB0_60; + + mov.u64 %rd39, $str13; + cvta.global.u64 %rd40, %rd39; + mov.u64 %rd41, 0; + // Callseq Start 6 + { + .reg .b32 temp_param_reg; + // <end>} + .param .b64 param0; + st.param.b64 [param0+0], %rd40; + .param .b64 param1; + st.param.b64 [param1+0], %rd41; + .param .b32 retval0; + call.uni (retval0), + vprintf, + ( + param0, + param1 + ); + ld.param.b32 %r56, [retval0+0]; + + //{ + }// Callseq End 6 + bra.uni BB0_60; + +BB0_7: + setp.eq.s32 %p8, %r40, 1023; + add.u64 %rd14, %SP, 0; + cvta.to.local.u64 %rd15, %rd14; + add.s64 %rd1, %rd15, 4; + @%p8 bra BB0_13; + bra.uni BB0_8; + +BB0_13: + ld.volatile.global.u32 %r6, [_ZN21rti_internal_register14reg_rayIndex_xE]; + ld.volatile.global.u32 %r7, [_ZN21rti_internal_register14reg_rayIndex_yE]; + ld.volatile.global.u32 %r8, [_ZN21rti_internal_register14reg_rayIndex_zE]; + ld.volatile.global.u32 %r9, [_ZN21rti_internal_register21reg_exception_detail0E]; + // inline asm + call (%r45), _rt_print_active, (); + // inline asm + setp.eq.s32 %p12, %r45, 0; + @%p12 bra BB0_60; + + st.local.u32 [%rd15], %r6; + st.local.u32 [%rd1], %r7; + st.local.v2.u32 [%rd1+4], {%r8, %r9}; + mov.u64 %rd26, $str16; + cvta.global.u64 %rd27, %rd26; + // Callseq Start 2 + { + .reg .b32 temp_param_reg; + // <end>} + .param .b64 param0; + st.param.b64 [param0+0], %rd27; + .param .b64 param1; + st.param.b64 [param1+0], %rd14; + .param .b32 retval0; + call.uni (retval0), + vprintf, + ( + param0, + param1 + ); + ld.param.b32 %r46, [retval0+0]; + + //{ + }// Callseq End 2 + bra.uni BB0_60; + +BB0_8: + add.s32 %r2, %r40, -1024; + setp.lt.u32 %p9, %r2, 64512; + ld.volatile.global.u32 %r3, [_ZN21rti_internal_register14reg_rayIndex_xE]; + ld.volatile.global.u32 %r4, [_ZN21rti_internal_register14reg_rayIndex_yE]; + ld.volatile.global.u32 %r5, [_ZN21rti_internal_register14reg_rayIndex_zE]; + @%p9 bra BB0_11; + bra.uni BB0_9; + +BB0_11: + // inline asm + call (%r43), _rt_print_active, (); + // inline asm + setp.eq.s32 %p11, %r43, 0; + @%p11 bra BB0_60; + + add.s32 %r83, %r40, -1024; + st.local.u32 [%rd15], %r83; + st.local.u32 [%rd1], %r3; + st.local.v2.u32 [%rd1+4], {%r4, %r5}; + mov.u64 %rd22, $str17; + cvta.global.u64 %rd23, %rd22; + // Callseq Start 1 + { + .reg .b32 temp_param_reg; + // <end>} + .param .b64 param0; + st.param.b64 [param0+0], %rd23; + .param .b64 param1; + st.param.b64 [param1+0], %rd14; + .param .b32 retval0; + call.uni (retval0), + vprintf, + ( + param0, + param1 + ); + ld.param.b32 %r44, [retval0+0]; + + //{ + }// Callseq End 1 + bra.uni BB0_60; + +BB0_9: + // inline asm + call (%r41), _rt_print_active, (); + // inline asm + setp.eq.s32 %p10, %r41, 0; + @%p10 bra BB0_60; + + add.u64 %rd16, %SP, 184; + cvta.to.local.u64 %rd17, %rd16; + st.local.v2.u32 [%rd17], {%r3, %r4}; + st.local.u32 [%rd17+8], %r5; + mov.u64 %rd18, $str18; + cvta.global.u64 %rd19, %rd18; + // Callseq Start 0 + { + .reg .b32 temp_param_reg; + // <end>} + .param .b64 param0; + st.param.b64 [param0+0], %rd19; + .param .b64 param1; + st.param.b64 [param1+0], %rd16; + .param .b32 retval0; + call.uni (retval0), + vprintf, + ( + param0, + param1 + ); + ld.param.b32 %r42, [retval0+0]; + + //{ + }// Callseq End 0 + bra.uni BB0_60; +} + + |