diff options
Diffstat (limited to 'zluda_rt/src/tests/exception.ptx')
| -rw-r--r-- | zluda_rt/src/tests/exception.ptx | 294 |
1 files changed, 294 insertions, 0 deletions
diff --git a/zluda_rt/src/tests/exception.ptx b/zluda_rt/src/tests/exception.ptx new file mode 100644 index 0000000..b0304e0 --- /dev/null +++ b/zluda_rt/src/tests/exception.ptx @@ -0,0 +1,294 @@ +// +// Generated by NVIDIA NVVM Compiler +// +// Compiler Build ID: CL-27506705 +// Cuda compilation tools, release 10.2, V10.2.89 +// Based on LLVM 3.4svn +// + +.version 6.5 +.target sm_30 +.address_size 64 + + // .globl _Z14call_callable2j +.visible .global .align 1 .b8 var_buffer[1]; +.visible .global .align 4 .b8 bvh[4]; +.visible .global .align 8 .b8 launch_index[8]; +.visible .global .align 4 .b8 callable1[4]; +.visible .global .align 4 .b8 callable2[4]; +.visible .global .align 4 .b8 callable3[4]; +.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo3bvhE[8] = {82, 97, 121, 0, 4, 0, 0, 0}; +.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo12launch_indexE[8] = {82, 97, 121, 0, 8, 0, 0, 0}; +.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo9callable1E[8] = {82, 97, 121, 0, 4, 0, 0, 0}; +.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo9callable2E[8] = {82, 97, 121, 0, 4, 0, 0, 0}; +.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo9callable3E[8] = {82, 97, 121, 0, 4, 0, 0, 0}; +.visible .global .align 8 .u64 _ZN21rti_internal_register20reg_bitness_detectorE; +.visible .global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail0E; +.visible .global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail1E; +.visible .global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail2E; +.visible .global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail3E; +.visible .global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail4E; +.visible .global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail5E; +.visible .global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail6E; +.visible .global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail7E; +.visible .global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail8E; +.visible .global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail9E; +.visible .global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail0E; +.visible .global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail1E; +.visible .global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail2E; +.visible .global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail3E; +.visible .global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail4E; +.visible .global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail5E; +.visible .global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail6E; +.visible .global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail7E; +.visible .global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail8E; +.visible .global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail9E; +.visible .global .align 4 .u32 _ZN21rti_internal_register14reg_rayIndex_xE; +.visible .global .align 4 .u32 _ZN21rti_internal_register14reg_rayIndex_yE; +.visible .global .align 4 .u32 _ZN21rti_internal_register14reg_rayIndex_zE; +.visible .global .align 1 .b8 _ZN21rti_internal_typename3bvhE[9] = {114, 116, 79, 98, 106, 101, 99, 116, 0}; +.visible .global .align 1 .b8 _ZN21rti_internal_typename12launch_indexE[6] = {117, 105, 110, 116, 50, 0}; +.visible .global .align 1 .b8 _ZN21rti_internal_typename9callable1E[13] = {105, 110, 116, 95, 111, 112, 101, 114, 97, 116, 111, 114, 0}; +.visible .global .align 1 .b8 _ZN21rti_internal_typename9callable2E[13] = {105, 110, 116, 95, 111, 112, 101, 114, 97, 116, 111, 114, 0}; +.visible .global .align 1 .b8 _ZN21rti_internal_typename9callable3E[13] = {105, 110, 116, 95, 111, 112, 101, 114, 97, 116, 111, 114, 0}; +.visible .global .align 4 .u32 _ZN21rti_internal_typeenum3bvhE = 4919; +.visible .global .align 4 .u32 _ZN21rti_internal_typeenum12launch_indexE = 4919; +.visible .global .align 4 .u32 _ZN21rti_internal_typeenum9callable1E = 4920; +.visible .global .align 4 .u32 _ZN21rti_internal_typeenum9callable2E = 4920; +.visible .global .align 4 .u32 _ZN21rti_internal_typeenum9callable3E = 4920; +.visible .global .align 1 .b8 _ZN21rti_internal_semantic3bvhE[1]; +.visible .global .align 1 .b8 _ZN21rti_internal_semantic12launch_indexE[14] = {114, 116, 76, 97, 117, 110, 99, 104, 73, 110, 100, 101, 120, 0}; +.visible .global .align 1 .b8 _ZN21rti_internal_semantic9callable1E[1]; +.visible .global .align 1 .b8 _ZN21rti_internal_semantic9callable2E[1]; +.visible .global .align 1 .b8 _ZN21rti_internal_semantic9callable3E[1]; +.visible .global .align 1 .b8 _ZN23rti_internal_annotation3bvhE[1]; +.visible .global .align 1 .b8 _ZN23rti_internal_annotation12launch_indexE[1]; +.visible .global .align 1 .b8 _ZN23rti_internal_annotation9callable1E[1]; +.visible .global .align 1 .b8 _ZN23rti_internal_annotation9callable2E[1]; +.visible .global .align 1 .b8 _ZN23rti_internal_annotation9callable3E[1]; + +.visible .func (.param .b32 func_retval0) _Z14call_callable2j( + .param .b32 _Z14call_callable2j_param_0 +) +{ + .reg .b32 %r<4>; + .reg .b64 %rd<4>; + + + ld.param.u32 %r2, [_Z14call_callable2j_param_0]; + ld.global.u32 %r1, [callable2]; + mov.u64 %rd2, 0; + // inline asm + call (%rd1), _rt_callable_program_from_id_v2_64, (%r1, %rd2); + // inline asm + // Callseq Start 0 + { + .reg .b32 temp_param_reg; + // <end>} + .param .b32 param0; + st.param.b32 [param0+0], %r2; + .param .b32 retval0; + prototype_0 : .callprototype (.param .b32 _) _ (.param .b32 _) ; + call (retval0), + %rd1, + ( + param0 + ) + , prototype_0; + ld.param.b32 %r3, [retval0+0]; + + //{ + }// Callseq End 0 + st.param.b32 [func_retval0+0], %r3; + ret; +} + + // .globl _Z14throw_callablej +.visible .func (.param .b32 func_retval0) _Z14throw_callablej( + .param .b32 _Z14throw_callablej_param_0 +) +{ + .reg .b32 %r<3>; + + + ld.param.u32 %r2, [_Z14throw_callablej_param_0]; + add.s32 %r1, %r2, 1024; + // inline asm + call _rt_throw, (%r1); + // inline asm + st.param.b32 [func_retval0+0], %r2; + ret; +} + + // .globl _Z18throw_callable_subv +.visible .func _Z18throw_callable_subv( + +) +{ + .reg .b32 %r<4>; + .reg .b64 %rd<4>; + + + ld.global.u32 %r1, [callable3]; + mov.u64 %rd2, 0; + // inline asm + call (%rd1), _rt_callable_program_from_id_v2_64, (%r1, %rd2); + // inline asm + mov.u32 %r2, 1; + // Callseq Start 1 + { + .reg .b32 temp_param_reg; + // <end>} + .param .b32 param0; + st.param.b32 [param0+0], %r2; + .param .b32 retval0; + prototype_1 : .callprototype (.param .b32 _) _ (.param .b32 _) ; + call (retval0), + %rd1, + ( + param0 + ) + , prototype_1; + ld.param.b32 %r3, [retval0+0]; + + //{ + }// Callseq End 1 + ret; +} + + // .globl _Z19throw_callable_mainj +.visible .func (.param .b32 func_retval0) _Z19throw_callable_mainj( + .param .b32 _Z19throw_callable_mainj_param_0 +) +{ + .reg .b32 %r<2>; + + + ld.param.u32 %r1, [_Z19throw_callable_mainj_param_0]; + // Callseq Start 2 + { + .reg .b32 temp_param_reg; + // <end>} + call.uni + _Z18throw_callable_subv, + ( + ); + + //{ + }// Callseq End 2 + st.param.b32 [func_retval0+0], %r1; + ret; +} + + // .globl _Z5tracev +.visible .entry _Z5tracev( + +) +{ + .local .align 1 .b8 __local_depot4[1]; + .reg .b64 %SP; + .reg .b64 %SPL; + .reg .b16 %rs<2>; + .reg .f32 %f<9>; + .reg .b32 %r<7>; + .reg .b64 %rd<3>; + + + mov.u64 %SPL, __local_depot4; + cvta.local.u64 %SP, %SPL; + add.u64 %rd1, %SP, 0; + add.u64 %rd2, %SPL, 0; + ld.global.u32 %r6, [launch_index]; + cvt.rn.f32.u32 %f1, %r6; + mov.u16 %rs1, 0; + st.local.u8 [%rd2], %rs1; + ld.global.u32 %r1, [bvh]; + mov.u32 %r3, 255; + mov.u32 %r4, 0; + mov.u32 %r5, 1; + mov.f32 %f3, 0fBF800000; + mov.f32 %f6, 0f3F800000; + mov.f32 %f7, 0f00000000; + mov.f32 %f8, 0f6C4ECB8F; + // inline asm + call _rt_trace_mask_flags_64, (%r1, %f1, %f7, %f3, %f7, %f7, %f6, %r4, %f7, %f8, %r3, %r4, %rd1, %r5); + // inline asm + ret; +} + + // .globl _Z6throw_v +.visible .entry _Z6throw_v( + +) +{ + .reg .b32 %r<2>; + + + mov.u32 %r1, 1024; + // inline asm + call _rt_throw, (%r1); + // inline asm + ret; +} + + // .globl _Z9exceptionv +.visible .entry _Z9exceptionv( + +) +{ + .reg .b32 %r<4>; + .reg .b64 %rd<8>; + + + // inline asm + call (%r1), _rt_get_exception_code, (); + // inline asm + mov.u64 %rd7, var_buffer; + cvta.global.u64 %rd2, %rd7; + mov.u32 %r2, 1; + mov.u32 %r3, 4; + mov.u64 %rd6, 0; + // inline asm + call (%rd1), _rt_buffer_get_64, (%rd2, %r2, %r3, %rd6, %rd6, %rd6, %rd6); + // inline asm + st.u32 [%rd1], %r1; + ret; +} + + // .globl _Z14call_callable1v +.visible .entry _Z14call_callable1v( + +) +{ + .reg .b32 %r<4>; + .reg .b64 %rd<4>; + + + ld.global.u32 %r1, [callable1]; + mov.u64 %rd2, 0; + // inline asm + call (%rd1), _rt_callable_program_from_id_v2_64, (%r1, %rd2); + // inline asm + mov.u32 %r2, 1; + // Callseq Start 3 + { + .reg .b32 temp_param_reg; + // <end>} + .param .b32 param0; + st.param.b32 [param0+0], %r2; + .param .b32 retval0; + prototype_3 : .callprototype (.param .b32 _) _ (.param .b32 _) ; + call (retval0), + %rd1, + ( + param0 + ) + , prototype_3; + ld.param.b32 %r3, [retval0+0]; + + //{ + }// Callseq End 3 + ret; +} + + |