935 lines
30 KiB
Plaintext
935 lines
30 KiB
Plaintext
|
//
|
||
|
|
// 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;
|
||
|
|
}
|
||
|
|
|
||
|
|
|