diff options
Diffstat (limited to 'zluda_rt/src/tests/buffer_id_callable.ptx')
| -rw-r--r-- | zluda_rt/src/tests/buffer_id_callable.ptx | 110 |
1 files changed, 110 insertions, 0 deletions
diff --git a/zluda_rt/src/tests/buffer_id_callable.ptx b/zluda_rt/src/tests/buffer_id_callable.ptx new file mode 100644 index 0000000..4826996 --- /dev/null +++ b/zluda_rt/src/tests/buffer_id_callable.ptx @@ -0,0 +1,110 @@ +// +// Generated by NVIDIA NVVM Compiler +// +// Compiler Build ID: CL-31833905 +// Cuda compilation tools, release 11.8, V11.8.89 +// Based on NVVM 7.0.1 +// + +.version 7.8 +.target sm_52 +.address_size 64 + + // .globl _Z8callablev +.visible .global .align 1 .b8 buffers[1]; +.visible .global .align 4 .b8 program[4]; +.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo7programE[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_typename7programE[32] = {114, 116, 67, 97, 108, 108, 97, 98, 108, 101, 80, 114, 111, 103, 114, 97, 109, 73, 100, 60, 118, 111, 105, 100, 40, 118, 111, 105, 100, 41, 62, 0}; +.visible .global .align 4 .u32 _ZN21rti_internal_typeenum7programE = 4920; +.visible .global .align 1 .b8 _ZN21rti_internal_semantic7programE[1]; +.visible .global .align 1 .b8 _ZN23rti_internal_annotation7programE[1]; + +.visible .func _Z8callablev() +{ + .reg .b32 %r<17>; + .reg .b64 %rd<34>; + + + mov.u64 %rd33, buffers; + cvta.global.u64 %rd2, %rd33; + mov.u32 %r14, 1; + mov.u32 %r15, 4; + mov.u64 %rd32, 0; + // begin inline asm + call (%rd1), _rt_buffer_get_64, (%rd2, %r14, %r15, %rd32, %rd32, %rd32, %rd32); + // end inline asm + ld.u32 %r3, [%rd1]; + mov.u64 %rd8, 2; + // begin inline asm + call (%rd7), _rt_buffer_get_id_64, (%r3, %r14, %r15, %rd8, %rd32, %rd32, %rd32); + // end inline asm + mov.u32 %r16, 18364300; + st.u32 [%rd7], %r16; + // begin inline asm + call (%rd12), _rt_buffer_get_64, (%rd2, %r14, %r15, %rd32, %rd32, %rd32, %rd32); + // end inline asm + ld.u32 %r8, [%rd12]; + // begin inline asm + call (%rd18, %rd19, %rd20, %rd21), _rt_buffer_get_id_size_64, (%r8, %r14, %r15); + // end inline asm + // begin inline asm + call (%rd22), _rt_buffer_get_64, (%rd2, %r14, %r15, %rd32, %rd32, %rd32, %rd32); + // end inline asm + ld.u32 %r13, [%rd22]; + mov.u64 %rd29, 1; + // begin inline asm + call (%rd28), _rt_buffer_get_id_64, (%r13, %r14, %r15, %rd29, %rd32, %rd32, %rd32); + // end inline asm + st.u32 [%rd28], %rd18; + ret; + +} + // .globl _Z5startv +.visible .entry _Z5startv() +{ + .reg .b32 %r<2>; + .reg .b64 %rd<4>; + + + ld.global.u32 %r1, [program]; + mov.u64 %rd2, 0; + // begin inline asm + call (%rd1), _rt_callable_program_from_id_v2_64, (%r1, %rd2); + // end inline asm + { // callseq 0, 0 + .reg .b32 temp_param_reg; + prototype_0 : .callprototype ()_ (); + call + %rd1, + ( + ) + , prototype_0; + } // callseq 0 + ret; + +} + |