
64-bit ELF: type=ET_EXEC, ABI=8, sm=90, toolkit=13.3, flags=0x9005a04
Sections:
Index Offset   Size ES Align                      Type    Flags Link     Info Name
    1     40    1ab  0  1                       STRTAB        0    0        0 .shstrtab
    2    21a    1f5  0  1                       STRTAB        0    0        0 .strtab
    3    410    138 18  8                       SYMTAB        0    2        d .symtab
    4    548     68  0  1                     PROGBITS        0    0        0 .debug_frame
    5    5b0     67  0  1                     PROGBITS        0    0        0 .debug_line
    6    617     51  0  1                     PROGBITS        0    0        0 .nv_debug_line_sass
    7    668    3ac  0  1                     PROGBITS        0    0        0 .nv_debug_ptx_txt
    8    a14     b4  0  4                         NOTE  2000000    0        0 .note.nv.tkinfo
    9    ac8     20  0  4                         NOTE  1000040    8        b .note.nv.cuinfo
    a    ae8     24  0  4                    CUDA_INFO        0    3        0 .nv.info
    b    b0c     1c  0  4             CUDA_COMPAT_INFO        0    0        0 .nv.compat
    c    b28     84  0  4                    CUDA_INFO       40    3       11 .nv.info._Z5saxpyifPKfPf
    d    bac     20  8  4               CUDA_CALLGRAPH        0    3        0 .nv.callgraph
    e    bd0     18 18  8                         RELA       40    3        5 .rela.debug_line
    f    be8     18 18  8                         RELA       40    3        6 .rela.nv_debug_line_sass
   10    c00     18 18  8                         RELA       40    3        4 .rela.debug_frame
   11    c80    200  0 80                     PROGBITS        6    3        b .text._Z5saxpyifPKfPf
   12    e80      0  0  1                       NOBITS        3    0        0 .nv.shared.reserved.0
   13    e80    228  0  4                     PROGBITS       42    0       11 .nv.constant0._Z5saxpyifPKfPf

.section .strtab

.section .shstrtab

.section .symtab
 index           value           size      info    other  shndx    name  
   0               0               0        0        0      0     (null)
 0x1               0               0      0x3        0    0x8     .note.nv.tkinfo
 0x2               0               0      0x3        0    0x9     .note.nv.cuinfo
 0x3               0               0      0x3        0   0x11     .text._Z5saxpyifPKfPf
 0x4               0             0x4     0x21        0      0     .nv.reservedSmem.offset0
 0x5               0               0     0x20     0xa0   0x12     __nv_reservedSMEM_offset_0_alias
 0x6               0               0      0x3        0    0x4     .debug_frame
 0x7               0               0      0x3        0    0x5     .debug_line
 0x8               0               0      0x3        0    0x6     .nv_debug_line_sass
 0x9               0               0      0x3        0    0x7     .nv_debug_ptx_txt
 0xa               0               0      0x3        0    0xd     .nv.callgraph
 0xb               0           0x200     0x12     0x10   0x11     _Z5saxpyifPKfPf
 0xc               0               0      0x3        0   0x13     .nv.constant0._Z5saxpyifPKfPf


.nv.shared.reserved.0
   No data to be dumped. 


.nv.constant0._Z5saxpyifPKfPf
0x00000000 0x00000000 0x00000000 0x00000000
0x00000000 0x00000000 0x00000000 0x00000000
0x00000000 0x00000000 0x00000000 0x00000000
0x00000000 0x00000000 0x00000000 0x00000000
0x00000000 0x00000000 0x00000000 0x00000000
0x00000000 0x00000000 0x00000000 0x00000000
0x00000000 0x00000000 0x00000000 0x00000000
0x00000000 0x00000000 0x00000000 0x00000000
0x00000000 0x00000000 0x00000000 0x00000000
0x00000000 0x00000000 0x00000000 0x00000000
0x00000000 0x00000000 0x00000000 0x00000000
0x00000000 0x00000000 0x00000000 0x00000000
0x00000000 0x00000000 0x00000000 0x00000000
0x00000000 0x00000000 0x00000000 0x00000000
0x00000000 0x00000000 0x00000000 0x00000000
0x00000000 0x00000000 0x00000000 0x00000000
0x00000000 0x00000000 0x00000000 0x00000000
0x00000000 0x00000000 0x00000000 0x00000000
0x00000000 0x00000000 0x00000000 0x00000000
0x00000000 0x00000000 0x00000000 0x00000000
0x00000000 0x00000000 0x00000000 0x00000000
0x00000000 0x00000000 0x00000000 0x00000000
0x00000000 0x00000000 0x00000000 0x00000000
0x00000000 0x00000000 0x00000000 0x00000000
0x00000000 0x00000000 0x00000000 0x00000000
0x00000000 0x00000000 0x00000000 0x00000000
0x00000000 0x00000000 0x00000000 0x00000000
0x00000000 0x00000000 0x00000000 0x00000000
0x00000000 0x00000000 0x00000000 0x00000000
0x00000000 0x00000000 0x00000000 0x00000000
0x00000000 0x00000000 0x00000000 0x00000000
0x00000000 0x00000000 0x00000000 0x00000000
0x00000000 0x00000000 0x00000000 0x00000000
0x00000000 0x00000000 0x00000000 0x00000000
0x00000000 0x00000000 


.section .note.nv.tkinfo
Entry: 1
  Owner                  Data size        Description
  NVIDIA Corp            156              NVIDIA CUDA Toolkit Information
	Note Version: 2
	Tool Name: ptxas
	Tool Version: Cuda compilation tools, release 13.3, V13.3.73
	Tool Branch: Build cuda_13.3.r13.3/compiler.38244171_0
	Tool Command Line Arguments: -v  -lineinfo  -arch sm_90 -m 64 


.section .note.nv.cuinfo
Entry: 1
  Owner                  Data size        Description
  NVIDIA Corp            8                NVIDIA CUDA Information
	Note Version: 2
	CUDA Virtual SM: sm_90
	CUDA Tool Kit Version: 13.3


.nv.info
	<0x1>
	Attribute:	EIATTR_REGCOUNT
	Format:	EIFMT_SVAL
	Value:	function: _Z5saxpyifPKfPf(0xb)	register count: 10
	<0x2>
	Attribute:	EIATTR_FRAME_SIZE
	Format:	EIFMT_SVAL
	Value:	function: _Z5saxpyifPKfPf(0xb)	frame size: 0x0
	<0x3>
	Attribute:	EIATTR_MIN_STACK_SIZE
	Format:	EIFMT_SVAL
	Value:	function: _Z5saxpyifPKfPf(0xb)	min stack size: 0x0


.nv.compat
	<0x1>
	Attribute:	EICOMPAT_ATTR_CUDA_ACCELERATOR_TARGET
	Format:	EIFMT_BVAL
	Value:	0x0
	<0x2>
	Attribute:	EICOMPAT_ATTR_ISA_CLASS
	Format:	EIFMT_BVAL
	Value:	0x1
	<0x3>
	Attribute:	EICOMPAT_ATTR_MERCURY_ISA_MAJOR_MINOR_VERSION_V2
	Format:	EIFMT_HVAL
	Value:	1.1
	<0x4>
	Attribute:	EICOMPAT_ATTR_INST_TENSORMAP_V1
	Format:	EIFMT_BVAL
	Value:	0x0
	<0x5>
	Attribute:	EICOMPAT_ATTR_CAN_FASTPATH_FINALIZE
	Format:	EIFMT_SVAL
	Value:	0x0 0x0 


.nv.info._Z5saxpyifPKfPf
	<0x1>
	Attribute:	EIATTR_LANGUAGE
	Format:	EIFMT_SVAL
	Value:	Language:	 PTX

	<0x2>
	Attribute:	EIATTR_CUDA_API_VERSION
	Format:	EIFMT_SVAL
	Value:	0x85 
	<0x3>
	Attribute:	EIATTR_KPARAM_INFO
	Format:	EIFMT_SVAL
	Value:	Index : 0x0	Ordinal : 0x3	Offset  : 0x10	Size    : 0x8
		Pointee's logAlignment : 0x0	Space : 0x0	cbank : 0x1f	Parameter Space : CBANK	
	<0x4>
	Attribute:	EIATTR_KPARAM_INFO
	Format:	EIFMT_SVAL
	Value:	Index : 0x0	Ordinal : 0x2	Offset  : 0x8	Size    : 0x8
		Pointee's logAlignment : 0x0	Space : 0x0	cbank : 0x1f	Parameter Space : CBANK	
	<0x5>
	Attribute:	EIATTR_KPARAM_INFO
	Format:	EIFMT_SVAL
	Value:	Index : 0x0	Ordinal : 0x1	Offset  : 0x4	Size    : 0x4
		Pointee's logAlignment : 0x0	Space : 0x0	cbank : 0x1f	Parameter Space : CBANK	
	<0x6>
	Attribute:	EIATTR_KPARAM_INFO
	Format:	EIFMT_SVAL
	Value:	Index : 0x0	Ordinal : 0x0	Offset  : 0x0	Size    : 0x4
		Pointee's logAlignment : 0x0	Space : 0x0	cbank : 0x1f	Parameter Space : CBANK	
	<0x7>
	Attribute:	EIATTR_SPARSE_MMA_MASK
	Format:	EIFMT_HVAL
	Value:	0x0
	<0x8>
	Attribute:	EIATTR_MAXREG_COUNT
	Format:	EIFMT_HVAL
	Value:	0xff
	<0x9>
	Attribute:	EIATTR_MERCURY_ISA_VERSION
	Format:	EIFMT_HVAL
	Value:	1.1
	<0x10>
	Attribute:	EIATTR_EXIT_INSTR_OFFSETS
	Format:	EIFMT_SVAL
	Value:	0x70 0x120 
	<0x11>
	Attribute:	EIATTR_CBANK_PARAM_SIZE
	Format:	EIFMT_HVAL
	Value:	0x18
	<0x12>
	Attribute:	EIATTR_PARAM_CBANK
	Format:	EIFMT_SVAL
	Value:	0xc 0x180210 
	<0x13>
	Attribute:	EIATTR_SW_WAR
	Format:	EIFMT_SVAL
	Value:	0x8 
	<0x14>
	Attribute:	EIATTR_NVSAL_SW_WAR
	Format:	EIFMT_HVAL
	Value:	 0x1


.nv.callgraph
 <0,-1>
 <0,-2>
 <0,-3>
 <0,-4>


.text._Z5saxpyifPKfPf
lmem=0	smem=0
0xff017b82 0x00000a00 0x00000800 0x000fe200
0x00007919 0x00000000 0x00002100 0x000e2e00
0x000479c3 0x00000000 0x00002500 0x000e3000
0xff077b82 0x00000000 0x00000800 0x000e2400
0x07077c24 0x00000004 0x0f8e0200 0x001fe200
0x00047ab9 0x00008400 0x00000800 0x000fc800
0x07007c0c 0x00000004 0x0bf06270 0x000fda00
0x0000094d 0x00000000 0x03800000 0x000fea00
0xff027b82 0x00008600 0x00000a00 0x000e2200
0x00047ab9 0x00008200 0x00000a00 0x000fe200
0x00067ab9 0x00008500 0x00000800 0x000fcc00
0xff047b82 0x00008800 0x00000a00 0x000e6200
0x07027825 0x00000004 0x078e0202 0x001fcc00
0x02027981 0x00000004 0x0c1e1900 0x000ea200
0x07047825 0x00000004 0x078e0204 0x002fca00
0x04077981 0x00000004 0x0c1e1900 0x000ea400
0x02077c23 0x00000006 0x08000007 0x004fca00
0x04007986 0x00000007 0x0c101904 0x000fe200
0x0000794d 0x00000000 0x03800000 0x000fea00
0x00fc7947 0xfffffffc 0x0383ffff 0x000fc000
0x00007918 0x00000000 0x00000000 0x000fc000
0x00007918 0x00000000 0x00000000 0x000fc000
0x00007918 0x00000000 0x00000000 0x000fc000
0x00007918 0x00000000 0x00000000 0x000fc000
0x00007918 0x00000000 0x00000000 0x000fc000
0x00007918 0x00000000 0x00000000 0x000fc000
0x00007918 0x00000000 0x00000000 0x000fc000
0x00007918 0x00000000 0x00000000 0x000fc000
0x00007918 0x00000000 0x00000000 0x000fc000
0x00007918 0x00000000 0x00000000 0x000fc000
0x00007918 0x00000000 0x00000000 0x000fc000
0x00007918 0x00000000 0x00000000 0x000fc000


.section .debug_line
  Offset:                         0
  Length:                         99
  DWARF Version:                   2
  Prologue Length:                58
  Minimum instruction length:      1
  Initial value of is_stmt:        1
  Line base:                       -5
  Line range:                      14
  Opcode base:                     10

Opcodes:
  Opcode 1 has 0 args
  Opcode 2 has 1 args
  Opcode 3 has 1 args
  Opcode 4 has 1 args
  Opcode 5 has 1 args
  Opcode 6 has 0 args
  Opcode 7 has 0 args
  Opcode 8 has 0 args
  Opcode 9 has 1 args

The Directory Table:
 /work/tests/fixtures/src

The File Name Table:
 Entry   Dir         Time    Size     Name
     1     1   1783195392     235     saxpy.cu

Line Number Statements:
  Extended opcode 2: set Address to 0x0
  Set File Name to entry 1 in the File Name Table
  Advance Line by 1 to 2
  Copy
  Special opcode 230: advance Address by 16 to 0x10 and Line by 1 to 3
  Advance Line by 1 to 4
  Advance PC by 64 to 0x50
  Copy
  Advance Line by 1 to 5
  Advance PC by 208 to 0x120
  Copy
  Advance PC by 224 to 0x200
  Extended opcode 1: End of Sequence



.section .rela.debug_line	RELA
0x47    _Z5saxpyifPKfPf    R_CUDA_64    0x0

.section .nv_debug_line_sass
  Offset:                         0
  Length:                         77
  DWARF Version:                   2
  Prologue Length:                16
  Minimum instruction length:      1
  Initial value of is_stmt:        1
  Line base:                       -5
  Line range:                      14
  Opcode base:                     10

Opcodes:
  Opcode 1 has 0 args
  Opcode 2 has 1 args
  Opcode 3 has 1 args
  Opcode 4 has 1 args
  Opcode 5 has 1 args
  Opcode 6 has 0 args
  Opcode 7 has 0 args
  Opcode 8 has 0 args
  Opcode 9 has 1 args

The Directory Table:

The File Name Table:
 Entry   Dir         Time    Size     Name

Line Number Statements:
  Extended opcode 2: set Address to 0x0
  Set File Name to entry 0 in the File Name Table
  Advance Line by 20 to 21
  Copy
  Advance Line by 15 to 36
  Advance PC by 16 to 0x10
  Copy
  Special opcode 227: advance Address by 16 to 0x20 and Line by -2 to 34
  Special opcode 230: advance Address by 16 to 0x30 and Line by 1 to 35
  Special opcode 231: advance Address by 16 to 0x40 and Line by 2 to 37
  Special opcode 231: advance Address by 16 to 0x50 and Line by 2 to 39
  Advance Line by 1 to 40
  Advance PC by 32 to 0x70
  Copy
  Special opcode 236: advance Address by 16 to 0x80 and Line by 7 to 47
  Advance Line by -26 to 21
  Advance PC by 16 to 0x90
  Copy
  Advance Line by 30 to 51
  Advance PC by 16 to 0xa0
  Copy
  Special opcode 227: advance Address by 16 to 0xb0 and Line by -2 to 49
  Special opcode 227: advance Address by 16 to 0xc0 and Line by -2 to 47
  Special opcode 230: advance Address by 16 to 0xd0 and Line by 1 to 48
  Special opcode 230: advance Address by 16 to 0xe0 and Line by 1 to 49
  Special opcode 230: advance Address by 16 to 0xf0 and Line by 1 to 50
  Special opcode 230: advance Address by 16 to 0x100 and Line by 1 to 51
  Special opcode 230: advance Address by 16 to 0x110 and Line by 1 to 52
  Special opcode 233: advance Address by 16 to 0x120 and Line by 4 to 56
  Advance PC by 224 to 0x200
  Extended opcode 1: End of Sequence



.section .rela.nv_debug_line_sass	RELA
0x1d    _Z5saxpyifPKfPf    R_CUDA_64    0x0

.section .nv_debug_ptx_txt








.version 9.3
.target sm_90
.address_size 64



.visible .entry _Z5saxpyifPKfPf(
.param .u32 _Z5saxpyifPKfPf_param_0,
.param .f32 _Z5saxpyifPKfPf_param_1,
.param .u64 _Z5saxpyifPKfPf_param_2,
.param .u64 _Z5saxpyifPKfPf_param_3
)
{
.reg .pred 	%p<2>;
.reg .f32 	%f<5>;
.reg .b32 	%r<6>;
.reg .b64 	%rd<8>;



ld.param.u32 	%r2, [_Z5saxpyifPKfPf_param_0];
ld.param.f32 	%f1, [_Z5saxpyifPKfPf_param_1];
ld.param.u64 	%rd1, [_Z5saxpyifPKfPf_param_2];
ld.param.u64 	%rd2, [_Z5saxpyifPKfPf_param_3];

mov.u32 	%r3, %ctaid.x;
mov.u32 	%r4, %ntid.x;
mov.u32 	%r5, %tid.x;
mad.lo.s32 	%r1, %r3, %r4, %r5;

setp.ge.s32 	%p1, %r1, %r2;
@%p1 bra 	$L__BB0_2;


cvta.to.global.u64 	%rd3, %rd2;
cvta.to.global.u64 	%rd4, %rd1;

mul.wide.s32 	%rd5, %r1, 4;
add.s64 	%rd6, %rd4, %rd5;
ld.global.f32 	%f2, [%rd6];
add.s64 	%rd7, %rd3, %rd5;
ld.global.f32 	%f3, [%rd7];
fma.rn.f32 	%f4, %f2, %f1, %f3;
st.global.f32 	[%rd7], %f4;

$L__BB0_2:

ret;

}

.section .debug_frame
decodeDebugFrame, frameBuf 0xffffffff, total_length 104
CIE length 36, cie_id -1
version 3
augmentation slen 1
augmentation 
code_align_factor slen 1
data_align_factor slen 1
 Debug Frame Common Information Entry
  length:                 36
  CIE_id :                -1
  version:                3
  augmentation:           
  code align factor:      4
  data align factor:      -4
  return address register 0xffffffff
  initial instructions: 19 bytes, ptr = 0x8080810c, frameBuf = 0xffffffff
  DW_CFA_def_cfa register R1, offset 0
  DW_CFA_same_value R255
  DW_CFA_same_value R1
  DW_CFA_nop
  DW_CFA_nop
  DW_CFA_nop
 Debug Frame Description Entry
  length:                 44
  CIE_pointer:            0
  initial_location:       0x0
  address_range:          0x200
  function:               _Z5saxpyifPKfPf
  instructions: 20 bytes
  DW_CFA_advance_loc4 delta 32
  DW_CFA_def_cfa register R1, offset 0
  DW_CFA_advance_loc4 delta 40
  DW_CFA_nop
  DW_CFA_nop
  DW_CFA_nop
  DW_CFA_nop

.section .rela.debug_frame	RELA
0x44    _Z5saxpyifPKfPf    R_CUDA_64    0x0
