.amdgcn_target "amdgcn-unknown-amdhsa--gfx803"
	.text
.Ltext0:
	.file 1 "tst.c"
	.section	.text.startup,"ax",@progbits
	.align	4
	.globl	main


	.rodata
	.p2align	6
	.amdhsa_kernel	main
	  .amdhsa_user_sgpr_private_segment_buffer	1 ; (s0, s1, s2, s3)
	  .amdhsa_user_sgpr_dispatch_ptr	1 ; (s4, s5)
	  .amdhsa_user_sgpr_queue_ptr	1 ; (s6, s7)
	  .amdhsa_user_sgpr_kernarg_segment_ptr	1 ; (s8, s9)
	  ;.amdhsa_user_sgpr_dispatch_id	0
	  ;.amdhsa_user_sgpr_flat_scratch_init	0
	  ;.amdhsa_user_sgpr_private_segment_size	0
	  .amdhsa_system_sgpr_workgroup_id_x	1 ; (s10)
	  ;.amdhsa_system_sgpr_workgroup_id_y	0
	  ;.amdhsa_system_sgpr_workgroup_id_z	0
	  ;.amdhsa_system_sgpr_workgroup_info	0
	  .amdhsa_system_sgpr_private_segment_wavefront_offset	1 ; (s11)
	  ; work_item_id_X	1 (v0)
	  ; work_item_id_Y	1 (v1)
	  ; work_item_id_Z	1 (v2)
	  .amdhsa_system_vgpr_workitem_id	2
	  .amdhsa_next_free_vgpr	0
	  .amdhsa_next_free_sgpr	25
	  .amdhsa_reserve_vcc	1
	  .amdhsa_reserve_flat_scratch	0
	  .amdhsa_reserve_xnack_mask	0
	  .amdhsa_private_segment_fixed_size	16384
	  .amdhsa_group_segment_fixed_size	65536
	  .amdhsa_float_denorm_mode_32	3
	  .amdhsa_float_denorm_mode_16_64	3
	.end_amdhsa_kernel
        .amdgpu_metadata
        amdhsa.version:
          - 1
          - 0
        amdhsa.kernels:
          - .name: main
            .symbol: main.kd
            .kernarg_segment_size: 0
            .kernarg_segment_align: 0
            .group_segment_fixed_size: 65536
            .private_segment_fixed_size: 16384
            .wavefront_size: 64
            .sgpr_count: 25
            .vgpr_count: 0
            .max_flat_workgroup_size: 1024
        .end_amdgpu_metadata
	.text
	.align	256
	.type	main,@function
main:
.LFB0:
	.loc 1 1 12
	; using flat addressing in function
	; HSA kernel entry point
	; local vars size: 0
	; outgoing args size: 0
	s_mov_b32	s14, s0
	s_and_b32	s15, s1, 65535
	s_add_u32	s14, s14, s11
	s_addc_u32	s15, s15, 0
	s_mov_b64	s[16:17], s[14:15]
.LCFI0:
	s_mov_b32	m0, 65536
	.loc 1 1 14
	.loc 1 1 24 is_stmt 0
	s_mov_b32	s24, 0
	s_load_dwordx2	s[8:9], s[8:9], 16
	s_waitcnt	lgkmcnt(0)
	s_store_dword	s24, s[8:9], 0
	s_waitcnt	lgkmcnt(0)
	s_dcache_wb
	s_endpgm
.LFE0:
	.size	main, .-main
	.section	.debug_frame,"",@progbits
.Lframe0:
	.4byte	.LECIE0-.LSCIE0
.LSCIE0:
	.4byte	0xffffffff
	.byte	0x3
	.string	""
	.byte	0x1
	.byte	0x4
	.byte	0x10
	.byte	0xc
	.byte	0x30
	.byte	0
	.byte	0x9
	.byte	0x10
	.byte	0x32
	.align	8
.LECIE0:
.LSFDE0:
	.4byte	.LEFDE0-.LASFDE0
.LASFDE0:
	.4byte	.Lframe0
	.8byte	.LFB0
	.8byte	.LFE0-.LFB0
	.byte	0x4
	.4byte	.LCFI0-.LFB0
	.byte	0x12
	.byte	0x2e
	.byte	0x7e
	.align	8
.LEFDE0:
	.text
.Letext0:
	.section	.debug_info,"",@progbits
.Ldebug_info0:
	.4byte	0x4c
	.2byte	0x5
	.byte	0x1
	.byte	0x8
	.4byte	.Ldebug_abbrev0
	.byte	0x1
	.4byte	.LASF0
	.byte	0x1d
	.4byte	.LASF1
	.4byte	.LASF2
	.4byte	.LLRL0
	.8byte	0
	.4byte	.Ldebug_line0
	.byte	0x2
	.4byte	.LASF3
	.byte	0x1
	.byte	0x1
	.byte	0x5
	.4byte	0x48
	.8byte	.LFB0
	.8byte	.LFE0-.LFB0
	.byte	0x1
	.byte	0x9c
	.byte	0x3
	.byte	0x4
	.byte	0x5
	.string	"int"
	.byte	0
	.section	.debug_abbrev,"",@progbits
.Ldebug_abbrev0:
	.byte	0x1
	.byte	0x11
	.byte	0x1
	.byte	0x25
	.byte	0xe
	.byte	0x13
	.byte	0xb
	.byte	0x3
	.byte	0xe
	.byte	0x1b
	.byte	0xe
	.byte	0x55
	.byte	0x17
	.byte	0x11
	.byte	0x1
	.byte	0x10
	.byte	0x17
	.byte	0
	.byte	0
	.byte	0x2
	.byte	0x2e
	.byte	0
	.byte	0x3f
	.byte	0x19
	.byte	0x3
	.byte	0xe
	.byte	0x3a
	.byte	0xb
	.byte	0x3b
	.byte	0xb
	.byte	0x39
	.byte	0xb
	.byte	0x49
	.byte	0x13
	.byte	0x11
	.byte	0x1
	.byte	0x12
	.byte	0x7
	.byte	0x40
	.byte	0x18
	.byte	0x7a
	.byte	0x19
	.byte	0
	.byte	0
	.byte	0x3
	.byte	0x24
	.byte	0
	.byte	0xb
	.byte	0xb
	.byte	0x3e
	.byte	0xb
	.byte	0x3
	.byte	0x8
	.byte	0
	.byte	0
	.byte	0
	.section	.debug_aranges,"",@progbits
	.4byte	0x2c
	.2byte	0x2
	.4byte	.Ldebug_info0
	.byte	0x8
	.byte	0
	.2byte	0
	.2byte	0
	.8byte	.LFB0
	.8byte	.LFE0-.LFB0
	.8byte	0
	.8byte	0
	.section	.debug_rnglists,"",@progbits
.Ldebug_ranges0:
	.4byte	.Ldebug_ranges3-.Ldebug_ranges2
.Ldebug_ranges2:
	.2byte	0x5
	.byte	0x8
	.byte	0
	.4byte	0
.LLRL0:
	.byte	0x6
	.8byte	.LFB0
	.8byte	.LFE0
	.byte	0
.Ldebug_ranges3:
	.section	.debug_line,"",@progbits
.Ldebug_line0:
	.section	.debug_str,"MS",@progbits,1
.LASF0:
	.string	"GNU C17 12.0.0 20210918 (experimental) [master r12-3644-g7afcb534239] -march=fiji -msram-ecc=off -g -O2"
.LASF2:
	.string	"/packages/gcc/snap/gcc-snapshot-20210918"
.LASF1:
	.string	"tst.c"
.LASF3:
	.string	"main"
	.ident	"GCC: (GNU) 12.0.0 20210918 (experimental) [master r12-3644-g7afcb534239]"