/*
   Copyright (C) 2021-2022 Free Software Foundation, Inc.
   Copyright (C) 2021-2022 Advanced Micro Devices, Inc. All rights reserved.

   This program is free software; you can redistribute it and/or modify
   it under the terms of the GNU General Public License as published by
   the Free Software Foundation; either version 3 of the License, or
   (at your option) any later version.

   This program is distributed in the hope that it will be useful,
   but WITHOUT ANY WARRANTY; without even the implied warranty of
   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
   GNU General Public License for more details.

   You should have received a copy of the GNU General Public License
   along with this program.  If not, see <http://www.gnu.org/licenses/>.

   This file has been generated from the lane-pc-vega20-kernel.cl
   OpenCL kernel source file, by using a set of compilation steps:

	bash$ clang -cc1 -x cl -O0 -debug-info-kind=line-tables-only
		    -disable-O0-optnone -triple amdgcn-amd-amdhsa
		    -target-cpu "gfx906" -Dcl_khr_fp64 -Dcl_khr_fp16
		    -finclude-default-header -emit-llvm-bc
		    -I /opt/rocm/opencl/include/CL/
		    -o lane-pc-vega20-kernel.raw.bc
		    lane-pc-vega20-kernel.cl

	bash$ opt -f -O0 -o lane-pc-vega20-kernel.orig.bc
		  lane-pc-vega20-kernel.raw.bc

	bash$ llvm-link -f -o lane-pc-vega20-kernel.linked.bc
			lane-pc-vega20-kernel.orig.bc
			/opt/rocm/amdgcn/bitcode/opencl.bc
			/opt/rocm/amdgcn/bitcode/ockl.bc

	bash$ llc -amdgpu-spill-cfi-saved-regs -relocation-model=pic
		  -march=amdgcn -mcpu="gfx906" -filetype=asm -o
		  -o lane-pc-vega20-kernel.S
		  lane-pc-vega20-kernel.linked.bc

   In addition, the file has been hand modified to expose labels
   that mark the end of each function so that they can be referenced
   by the hand written DWARF information found in lane-pc-vega20.exp
   test file.

   To support a new DW_AT_LLVM_lane_pc attribute, set of VPC_LINE_
   labels, that mark points of lane execution divergence / convergence,
   have also been added.

   Also, the .debug_info section had to be removed, so that it
   would not mask the same section generated by the test file.  */

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx906"
	.protected	__ockl_get_global_id    ; -- Begin function __ockl_get_global_id
	.weak	__ockl_get_global_id
	.p2align	2
	.type	__ockl_get_global_id,@function
__ockl_get_global_id:                   ; @__ockl_get_global_id
.Lfunc_begin0:
	.cfi_sections .debug_frame
	.cfi_startproc
; %bb.0:
	.cfi_llvm_def_aspace_cfa 64, 0, 6
	.cfi_escape 0x10, 0x10, 0x08, 0x90, 0x3e, 0x93, 0x04, 0x90, 0x3f, 0x93, 0x04 ;
	.cfi_undefined 2560
	.cfi_undefined 2561
	.cfi_undefined 2562
	.cfi_undefined 2563
	.cfi_undefined 2564
	.cfi_undefined 2565
	.cfi_undefined 2566
	.cfi_undefined 2567
	.cfi_undefined 36
	.cfi_undefined 37
	.cfi_undefined 38
	.cfi_undefined 39
	.cfi_undefined 42
	.cfi_undefined 43
	.cfi_undefined 44
	.cfi_undefined 45
	.cfi_undefined 46
	.cfi_undefined 47
	.cfi_undefined 48
	.cfi_undefined 49
	.cfi_undefined 50
	.cfi_undefined 51
	.cfi_undefined 52
	.cfi_undefined 53
	s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
	s_or_saveexec_b64 s[6:7], -1
	buffer_store_dword v8, off, s[0:3], s32 ; 4-byte Folded Spill
	.cfi_offset 2568, 0
	s_mov_b64 exec, s[6:7]
	v_writelane_b32 v8, s30, 0
	v_writelane_b32 v8, s31, 1
	.cfi_escape 0x10, 0x10, 0x0c, 0x90, 0x88, 0x14, 0x9d, 0x20, 0x00, 0x90, 0x88, 0x14, 0x9d, 0x20, 0x20 ;
	v_writelane_b32 v8, exec_lo, 2
	v_writelane_b32 v8, exec_hi, 3
	.cfi_escape 0x10, 0x11, 0x0c, 0x90, 0x88, 0x14, 0x9d, 0x20, 0x40, 0x90, 0x88, 0x14, 0x9d, 0x20, 0x60 ;
	v_cmp_lt_i32_e32 vcc, 0, v0
	s_mov_b64 s[6:7], 0
	s_mov_b64 s[16:17], 0
                                        ; implicit-def: $vgpr2
                                        ; implicit-def: $vgpr3
                                        ; implicit-def: $vgpr1
	s_and_saveexec_b64 s[10:11], vcc
	s_xor_b64 s[10:11], exec, s[10:11]
	s_cbranch_execz BB0_8
; %bb.1:                                ; %NodeBlock
	v_cmp_lt_i32_e32 vcc, 1, v0
	s_mov_b64 s[16:17], 0
                                        ; implicit-def: $sgpr15
                                        ; implicit-def: $vgpr1
	s_and_saveexec_b64 s[18:19], vcc
	s_xor_b64 s[18:19], exec, s[18:19]
	s_cbranch_execz BB0_5
; %bb.2:                                ; %LeafBlock1
	v_cmp_eq_u32_e32 vcc, 2, v0
	s_mov_b64 s[16:17], -1
                                        ; implicit-def: $sgpr15
                                        ; implicit-def: $vgpr1
	s_and_saveexec_b64 s[20:21], vcc
	s_cbranch_execz BB0_4
; %bb.3:
	s_load_dword s15, s[4:5], 0x8
	v_bfe_u32 v1, v31, 20, 10
	s_xor_b64 s[16:17], exec, -1
BB0_4:                                  ; %Flow22
	s_or_b64 exec, exec, s[20:21]
	s_and_b64 s[16:17], s[16:17], exec
BB0_5:                                  ; %Flow21
	s_or_saveexec_b64 s[18:19], s[18:19]
	s_waitcnt lgkmcnt(0)
	v_mov_b32_e32 v2, s15
	v_mov_b32_e32 v3, s14
	s_xor_b64 exec, exec, s[18:19]
	s_cbranch_execz BB0_7
; %bb.6:
	s_load_dword s14, s[4:5], 0x4
	v_bfe_u32 v1, v31, 10, 10
	v_mov_b32_e32 v3, s13
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s14, s14, 16
	v_mov_b32_e32 v2, s14
BB0_7:                                  ; %Flow23
	s_or_b64 exec, exec, s[18:19]
	s_and_b64 s[16:17], s[16:17], exec
BB0_8:                                  ; %Flow20
	s_or_saveexec_b64 s[10:11], s[10:11]
	s_xor_b64 exec, exec, s[10:11]
; %bb.9:                                ; %LeafBlock
	v_cmp_ne_u32_e32 vcc, 0, v0
	s_andn2_b64 s[14:15], s[16:17], exec
	s_and_b64 s[16:17], vcc, exec
	s_mov_b64 s[6:7], exec
	s_or_b64 s[16:17], s[14:15], s[16:17]
; %bb.10:                               ; %Flow24
	s_or_b64 exec, exec, s[10:11]
	s_and_saveexec_b64 s[10:11], s[16:17]
	s_xor_b64 s[10:11], exec, s[10:11]
; %bb.11:                               ; %NewDefault
	v_mov_b32_e32 v3, 0
	v_mov_b32_e32 v2, 1
	s_andn2_b64 s[6:7], s[6:7], exec
	v_mov_b32_e32 v1, v3
; %bb.12:                               ; %Flow25
	s_or_b64 exec, exec, s[10:11]
	s_and_saveexec_b64 s[10:11], s[6:7]
	s_cbranch_execz BB0_14
; %bb.13:
	s_load_dword s4, s[4:5], 0x4
	v_and_b32_e32 v1, 0x3ff, v31
	v_mov_b32_e32 v3, s12
	s_waitcnt lgkmcnt(0)
	v_mov_b32_e32 v2, s4
BB0_14:
	s_or_b64 exec, exec, s[10:11]
	v_cmp_lt_i32_e32 vcc, 0, v0
	s_mov_b64 s[4:5], 0
	s_mov_b64 s[12:13], 0
	s_mov_b64 s[10:11], 0
                                        ; implicit-def: $vgpr4_vgpr5
	s_and_saveexec_b64 s[6:7], vcc
	s_xor_b64 s[6:7], exec, s[6:7]
	s_cbranch_execz BB0_22
; %bb.15:                               ; %NodeBlock10
	v_cmp_lt_i32_e32 vcc, 1, v0
	s_mov_b64 s[12:13], 0
	s_mov_b64 s[14:15], 0
                                        ; implicit-def: $sgpr10_sgpr11
	s_and_saveexec_b64 s[16:17], vcc
	s_xor_b64 s[16:17], exec, s[16:17]
	s_cbranch_execz BB0_19
; %bb.16:                               ; %LeafBlock8
	v_cmp_eq_u32_e32 vcc, 2, v0
	s_mov_b64 s[18:19], -1
	s_mov_b64 s[12:13], 0
                                        ; implicit-def: $sgpr10_sgpr11
	s_and_saveexec_b64 s[14:15], vcc
; %bb.17:
	s_add_u32 s10, s8, 16
	s_addc_u32 s11, s9, 0
	s_mov_b64 s[12:13], exec
	s_xor_b64 s[18:19], exec, -1
; %bb.18:                               ; %Flow15
	s_or_b64 exec, exec, s[14:15]
	s_and_b64 s[14:15], s[12:13], exec
	s_and_b64 s[12:13], s[18:19], exec
BB0_19:                                 ; %Flow14
	s_or_saveexec_b64 s[16:17], s[16:17]
	v_mov_b32_e32 v4, s10
	v_mov_b32_e32 v5, s11
	s_xor_b64 exec, exec, s[16:17]
; %bb.20:
	s_add_u32 s10, s8, 8
	s_addc_u32 s11, s9, 0
	v_mov_b32_e32 v4, s10
	v_mov_b32_e32 v5, s11
	s_or_b64 s[14:15], s[14:15], exec
; %bb.21:                               ; %Flow16
	s_or_b64 exec, exec, s[16:17]
	s_and_b64 s[10:11], s[14:15], exec
	s_and_b64 s[12:13], s[12:13], exec
BB0_22:                                 ; %Flow
	s_or_saveexec_b64 s[6:7], s[6:7]
	s_xor_b64 exec, exec, s[6:7]
; %bb.23:                               ; %LeafBlock6
	v_cmp_ne_u32_e32 vcc, 0, v0
	s_andn2_b64 s[12:13], s[12:13], exec
	s_and_b64 s[14:15], vcc, exec
	s_mov_b64 s[4:5], exec
	s_or_b64 s[12:13], s[12:13], s[14:15]
; %bb.24:                               ; %Flow17
	s_or_b64 exec, exec, s[6:7]
	s_and_saveexec_b64 s[6:7], s[12:13]
	s_xor_b64 s[6:7], exec, s[6:7]
; %bb.25:                               ; %NewDefault5
	s_andn2_b64 s[4:5], s[4:5], exec
; %bb.26:                               ; %Flow18
	s_or_b64 exec, exec, s[6:7]
	s_and_saveexec_b64 s[6:7], s[4:5]
; %bb.27:
	v_mov_b32_e32 v4, s8
	v_mov_b32_e32 v5, s9
	s_or_b64 s[10:11], s[10:11], exec
; %bb.28:                               ; %Flow19
	s_or_b64 exec, exec, s[6:7]
	v_mov_b32_e32 v6, 0
	v_mov_b32_e32 v7, 0
	s_and_saveexec_b64 s[4:5], s[10:11]
	s_cbranch_execz BB0_30
; %bb.29:
	global_load_dwordx2 v[6:7], v[4:5], off
BB0_30:
	s_or_b64 exec, exec, s[4:5]
	v_and_b32_e32 v0, 0xffff, v2
	v_mul_lo_u32 v0, v3, v0
	v_add_u32_e32 v0, v0, v1
	s_waitcnt vmcnt(0)
	v_add_co_u32_e32 v0, vcc, v6, v0
	v_addc_co_u32_e32 v1, vcc, 0, v7, vcc
	s_or_saveexec_b64 s[4:5], -1
	buffer_load_dword v8, off, s[0:3], s32  ; 4-byte Folded Reload
	s_mov_b64 exec, s[4:5]
	s_waitcnt vmcnt(0)
	s_setpc_b64 s[30:31]
.Lfunc_end0:
	.size	__ockl_get_global_id, .Lfunc_end0-__ockl_get_global_id
	.cfi_endproc
                                        ; -- End function
	.section	.AMDGPU.csdata
; Function info:
; codeLenInByte = 600
; NumSgprs: 37
; NumVgprs: 32
; ScratchSize: 8
; MemoryBound: 0
	.text
	.protected	_Z13get_global_idj      ; -- Begin function _Z13get_global_idj
	.weak	_Z13get_global_idj
	.p2align	2
	.type	_Z13get_global_idj,@function
_Z13get_global_idj:                     ; @_Z13get_global_idj
.Lfunc_begin1:
	.cfi_startproc
; %bb.0:
	.cfi_llvm_def_aspace_cfa 64, 0, 6
	.cfi_escape 0x10, 0x10, 0x08, 0x90, 0x3e, 0x93, 0x04, 0x90, 0x3f, 0x93, 0x04 ;
	.cfi_undefined 2560
	.cfi_undefined 2561
	.cfi_undefined 2562
	.cfi_undefined 2563
	.cfi_undefined 2564
	.cfi_undefined 2565
	.cfi_undefined 2566
	.cfi_undefined 2567
	.cfi_undefined 2568
	.cfi_undefined 2569
	.cfi_undefined 2570
	.cfi_undefined 2571
	.cfi_undefined 2572
	.cfi_undefined 2573
	.cfi_undefined 2574
	.cfi_undefined 2575
	.cfi_undefined 2576
	.cfi_undefined 2577
	.cfi_undefined 2578
	.cfi_undefined 2579
	.cfi_undefined 2580
	.cfi_undefined 2581
	.cfi_undefined 2582
	.cfi_undefined 2583
	.cfi_undefined 2584
	.cfi_undefined 2585
	.cfi_undefined 2586
	.cfi_undefined 2587
	.cfi_undefined 2588
	.cfi_undefined 2589
	.cfi_undefined 2590
	.cfi_undefined 2591
	.cfi_undefined 2592
	.cfi_undefined 2593
	.cfi_undefined 2594
	.cfi_undefined 2595
	.cfi_undefined 2596
	.cfi_undefined 2597
	.cfi_undefined 2598
	.cfi_undefined 2599
	.cfi_undefined 2608
	.cfi_undefined 2609
	.cfi_undefined 2610
	.cfi_undefined 2611
	.cfi_undefined 2612
	.cfi_undefined 2613
	.cfi_undefined 2614
	.cfi_undefined 2615
	.cfi_undefined 2624
	.cfi_undefined 2625
	.cfi_undefined 2626
	.cfi_undefined 2627
	.cfi_undefined 2628
	.cfi_undefined 2629
	.cfi_undefined 2630
	.cfi_undefined 2631
	.cfi_undefined 2640
	.cfi_undefined 2641
	.cfi_undefined 2642
	.cfi_undefined 2643
	.cfi_undefined 2644
	.cfi_undefined 2645
	.cfi_undefined 2646
	.cfi_undefined 2647
	.cfi_undefined 2656
	.cfi_undefined 2657
	.cfi_undefined 2658
	.cfi_undefined 2659
	.cfi_undefined 2660
	.cfi_undefined 2661
	.cfi_undefined 2662
	.cfi_undefined 2663
	.cfi_undefined 2672
	.cfi_undefined 2673
	.cfi_undefined 2674
	.cfi_undefined 2675
	.cfi_undefined 2676
	.cfi_undefined 2677
	.cfi_undefined 2678
	.cfi_undefined 2679
	.cfi_undefined 2688
	.cfi_undefined 2689
	.cfi_undefined 2690
	.cfi_undefined 2691
	.cfi_undefined 2692
	.cfi_undefined 2693
	.cfi_undefined 2694
	.cfi_undefined 2695
	.cfi_undefined 2704
	.cfi_undefined 2705
	.cfi_undefined 2706
	.cfi_undefined 2707
	.cfi_undefined 2708
	.cfi_undefined 2709
	.cfi_undefined 2710
	.cfi_undefined 2711
	.cfi_undefined 2720
	.cfi_undefined 2721
	.cfi_undefined 2722
	.cfi_undefined 2723
	.cfi_undefined 2724
	.cfi_undefined 2725
	.cfi_undefined 2726
	.cfi_undefined 2727
	.cfi_undefined 2736
	.cfi_undefined 2737
	.cfi_undefined 2738
	.cfi_undefined 2739
	.cfi_undefined 2740
	.cfi_undefined 2741
	.cfi_undefined 2742
	.cfi_undefined 2743
	.cfi_undefined 2752
	.cfi_undefined 2753
	.cfi_undefined 2754
	.cfi_undefined 2755
	.cfi_undefined 2756
	.cfi_undefined 2757
	.cfi_undefined 2758
	.cfi_undefined 2759
	.cfi_undefined 2768
	.cfi_undefined 2769
	.cfi_undefined 2770
	.cfi_undefined 2771
	.cfi_undefined 2772
	.cfi_undefined 2773
	.cfi_undefined 2774
	.cfi_undefined 2775
	.cfi_undefined 2784
	.cfi_undefined 2785
	.cfi_undefined 2786
	.cfi_undefined 2787
	.cfi_undefined 2788
	.cfi_undefined 2789
	.cfi_undefined 2790
	.cfi_undefined 2791
	.cfi_undefined 2800
	.cfi_undefined 2801
	.cfi_undefined 2802
	.cfi_undefined 2803
	.cfi_undefined 2804
	.cfi_undefined 2805
	.cfi_undefined 2806
	.cfi_undefined 2807
	.cfi_undefined 32
	.cfi_undefined 33
	.cfi_undefined 34
	.cfi_undefined 35
	.cfi_undefined 36
	.cfi_undefined 37
	.cfi_undefined 38
	.cfi_undefined 39
	.cfi_undefined 40
	.cfi_undefined 41
	.cfi_undefined 42
	.cfi_undefined 43
	.cfi_undefined 44
	.cfi_undefined 45
	.cfi_undefined 46
	.cfi_undefined 47
	.cfi_undefined 48
	.cfi_undefined 49
	.cfi_undefined 50
	.cfi_undefined 51
	.cfi_undefined 52
	.cfi_undefined 53
	.cfi_undefined 54
	.cfi_undefined 55
	.cfi_undefined 56
	.cfi_undefined 57
	.cfi_undefined 58
	.cfi_undefined 59
	.cfi_undefined 60
	.cfi_undefined 61
	.cfi_undefined 62
	.cfi_undefined 63
	s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
	s_or_saveexec_b64 s[16:17], -1
	buffer_store_dword v40, off, s[0:3], s32 ; 4-byte Folded Spill
	.cfi_offset 2600, 0
	s_mov_b64 exec, s[16:17]
	v_writelane_b32 v40, s30, 0
	v_writelane_b32 v40, s31, 1
	.cfi_escape 0x10, 0x10, 0x0c, 0x90, 0xa8, 0x14, 0x9d, 0x20, 0x00, 0x90, 0xa8, 0x14, 0x9d, 0x20, 0x20 ;
	v_writelane_b32 v40, exec_lo, 2
	v_writelane_b32 v40, exec_hi, 3
	.cfi_escape 0x10, 0x11, 0x0c, 0x90, 0xa8, 0x14, 0x9d, 0x20, 0x40, 0x90, 0xa8, 0x14, 0x9d, 0x20, 0x60 ;
	s_getpc_b64 s[16:17]
	s_add_u32 s16, s16, __ockl_get_global_id@rel32@lo+4
	s_addc_u32 s17, s17, __ockl_get_global_id@rel32@hi+12
	s_or_saveexec_b64 s[18:19], -1
	buffer_load_dword v40, off, s[0:3], s32 ; 4-byte Folded Reload
	s_mov_b64 exec, s[18:19]
	s_setpc_b64 s[16:17]
.Lfunc_end1:
	.size	_Z13get_global_idj, .Lfunc_end1-_Z13get_global_idj
	.cfi_endproc
                                        ; -- End function
	.section	.AMDGPU.csdata
; Function info:
; codeLenInByte = 92
; NumSgprs: 37
; NumVgprs: 41
; ScratchSize: 16
; MemoryBound: 0
	.text
	.hidden GenValue
	.globl	GenValue                        ; -- Begin function GenValue
	.hidden Lfunc_end2
	.globl	Lfunc_end2
	.hidden VPC_LINE_42
	.globl	VPC_LINE_42
	.hidden VPC_LINE_47
	.globl	VPC_LINE_47
	.p2align	2
	.type	GenValue,@function
GenValue:                               ; @GenValue
GenValue$local:
.Lfunc_begin2:
	.file	1 "./lane-pc-vega20-kernel.cl"
	.loc	1 36 0                          ; lane-pc-vega20-kernel.cl:36:0
	.cfi_startproc
; %bb.0:                                ; %entry
	.cfi_llvm_def_aspace_cfa 64, 0, 6
	.cfi_escape 0x10, 0x10, 0x08, 0x90, 0x3e, 0x93, 0x04, 0x90, 0x3f, 0x93, 0x04 ;
	.cfi_undefined 2560
	.cfi_undefined 2561
	.cfi_undefined 2562
	.cfi_undefined 2563
	.cfi_undefined 2564
	.cfi_undefined 2565
	.cfi_undefined 2566
	.cfi_undefined 2567
	.cfi_undefined 2568
	.cfi_undefined 2569
	.cfi_undefined 2570
	.cfi_undefined 2571
	.cfi_undefined 2572
	.cfi_undefined 2573
	.cfi_undefined 2574
	.cfi_undefined 2575
	.cfi_undefined 2576
	.cfi_undefined 2577
	.cfi_undefined 2578
	.cfi_undefined 2579
	.cfi_undefined 2580
	.cfi_undefined 2581
	.cfi_undefined 2582
	.cfi_undefined 2583
	.cfi_undefined 2584
	.cfi_undefined 2585
	.cfi_undefined 2586
	.cfi_undefined 2587
	.cfi_undefined 2588
	.cfi_undefined 2589
	.cfi_undefined 2590
	.cfi_undefined 2591
	.cfi_undefined 2592
	.cfi_undefined 2593
	.cfi_undefined 2594
	.cfi_undefined 2595
	.cfi_undefined 2596
	.cfi_undefined 2597
	.cfi_undefined 2598
	.cfi_undefined 2599
	.cfi_undefined 2608
	.cfi_undefined 2609
	.cfi_undefined 2610
	.cfi_undefined 2611
	.cfi_undefined 2612
	.cfi_undefined 2613
	.cfi_undefined 2614
	.cfi_undefined 2615
	.cfi_undefined 2624
	.cfi_undefined 2625
	.cfi_undefined 2626
	.cfi_undefined 2627
	.cfi_undefined 2628
	.cfi_undefined 2629
	.cfi_undefined 2630
	.cfi_undefined 2631
	.cfi_undefined 2640
	.cfi_undefined 2641
	.cfi_undefined 2642
	.cfi_undefined 2643
	.cfi_undefined 2644
	.cfi_undefined 2645
	.cfi_undefined 2646
	.cfi_undefined 2647
	.cfi_undefined 2656
	.cfi_undefined 2657
	.cfi_undefined 2658
	.cfi_undefined 2659
	.cfi_undefined 2660
	.cfi_undefined 2661
	.cfi_undefined 2662
	.cfi_undefined 2663
	.cfi_undefined 2672
	.cfi_undefined 2673
	.cfi_undefined 2674
	.cfi_undefined 2675
	.cfi_undefined 2676
	.cfi_undefined 2677
	.cfi_undefined 2678
	.cfi_undefined 2679
	.cfi_undefined 2688
	.cfi_undefined 2689
	.cfi_undefined 2690
	.cfi_undefined 2691
	.cfi_undefined 2692
	.cfi_undefined 2693
	.cfi_undefined 2694
	.cfi_undefined 2695
	.cfi_undefined 2704
	.cfi_undefined 2705
	.cfi_undefined 2706
	.cfi_undefined 2707
	.cfi_undefined 2708
	.cfi_undefined 2709
	.cfi_undefined 2710
	.cfi_undefined 2711
	.cfi_undefined 2720
	.cfi_undefined 2721
	.cfi_undefined 2722
	.cfi_undefined 2723
	.cfi_undefined 2724
	.cfi_undefined 2725
	.cfi_undefined 2726
	.cfi_undefined 2727
	.cfi_undefined 2736
	.cfi_undefined 2737
	.cfi_undefined 2738
	.cfi_undefined 2739
	.cfi_undefined 2740
	.cfi_undefined 2741
	.cfi_undefined 2742
	.cfi_undefined 2743
	.cfi_undefined 2752
	.cfi_undefined 2753
	.cfi_undefined 2754
	.cfi_undefined 2755
	.cfi_undefined 2756
	.cfi_undefined 2757
	.cfi_undefined 2758
	.cfi_undefined 2759
	.cfi_undefined 2768
	.cfi_undefined 2769
	.cfi_undefined 2770
	.cfi_undefined 2771
	.cfi_undefined 2772
	.cfi_undefined 2773
	.cfi_undefined 2774
	.cfi_undefined 2775
	.cfi_undefined 2784
	.cfi_undefined 2785
	.cfi_undefined 2786
	.cfi_undefined 2787
	.cfi_undefined 2788
	.cfi_undefined 2789
	.cfi_undefined 2790
	.cfi_undefined 2791
	.cfi_undefined 2800
	.cfi_undefined 2801
	.cfi_undefined 2802
	.cfi_undefined 2803
	.cfi_undefined 2804
	.cfi_undefined 2805
	.cfi_undefined 2806
	.cfi_undefined 2807
	.cfi_undefined 32
	.cfi_undefined 33
	.cfi_undefined 34
	.cfi_undefined 35
	.cfi_undefined 36
	.cfi_undefined 37
	.cfi_undefined 38
	.cfi_undefined 39
	.cfi_undefined 40
	.cfi_undefined 41
	.cfi_undefined 42
	.cfi_undefined 43
	.cfi_undefined 44
	.cfi_undefined 45
	.cfi_undefined 46
	.cfi_undefined 47
	.cfi_undefined 48
	.cfi_undefined 49
	.cfi_undefined 50
	.cfi_undefined 51
	.cfi_undefined 52
	.cfi_undefined 53
	.cfi_undefined 54
	.cfi_undefined 55
	.cfi_undefined 56
	.cfi_undefined 57
	.cfi_undefined 58
	.cfi_undefined 59
	.cfi_undefined 60
	.cfi_undefined 61
	.cfi_undefined 62
	.cfi_undefined 63
	s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
	s_or_saveexec_b64 s[16:17], -1
	buffer_store_dword v40, off, s[0:3], s32 offset:128 ; 4-byte Folded Spill
	.cfi_offset 2600, 8192
	s_mov_b64 exec, s[16:17]
	v_writelane_b32 v40, s30, 2
	v_writelane_b32 v40, s31, 3
	.cfi_escape 0x10, 0x10, 0x0c, 0x90, 0xa8, 0x14, 0x9d, 0x20, 0x40, 0x90, 0xa8, 0x14, 0x9d, 0x20, 0x60 ;
	v_writelane_b32 v40, exec_lo, 4
	v_writelane_b32 v40, exec_hi, 5
	.cfi_escape 0x10, 0x11, 0x0e, 0x90, 0xa8, 0x14, 0x9d, 0x20, 0x80, 0x01, 0x90, 0xa8, 0x14, 0x9d, 0x20, 0xa0, 0x01 ;
	v_writelane_b32 v40, s33, 6
	.cfi_escape 0x10, 0x41, 0x05, 0x90, 0xa8, 0x14, 0xe4, 0x18 ;
	s_mov_b32 s33, s32
	.cfi_def_cfa_register 65
	s_add_u32 s32, s32, 0x2400
	s_getpc_b64 s[24:25]
	s_add_u32 s24, s24, __const.GenValue.array@rel32@lo+116
	s_addc_u32 s25, s25, __const.GenValue.array@rel32@hi+124
	s_getpc_b64 s[26:27]
	s_add_u32 s26, s26, __const.GenValue.array@rel32@lo+100
	s_addc_u32 s27, s27, __const.GenValue.array@rel32@hi+108
.Ltmp0:
	.loc	1 37 7 prologue_end             ; lane-pc-vega20-kernel.cl:37:7
	s_load_dwordx4 s[16:19], s[24:25], 0x0
	s_load_dwordx4 s[20:23], s[26:27], 0x0
	s_getpc_b64 s[24:25]
	s_add_u32 s24, s24, __const.GenValue.array@rel32@lo+84
	s_addc_u32 s25, s25, __const.GenValue.array@rel32@hi+92
	s_getpc_b64 s[26:27]
	s_add_u32 s26, s26, __const.GenValue.array@rel32@lo+68
	s_addc_u32 s27, s27, __const.GenValue.array@rel32@hi+76
	v_writelane_b32 v40, s30, 0
	s_waitcnt lgkmcnt(0)
	v_mov_b32_e32 v0, s19
	buffer_store_dword v0, off, s[0:3], s33 offset:124
	v_mov_b32_e32 v0, s18
	buffer_store_dword v0, off, s[0:3], s33 offset:120
	v_mov_b32_e32 v0, s17
	buffer_store_dword v0, off, s[0:3], s33 offset:116
	v_mov_b32_e32 v0, s16
	buffer_store_dword v0, off, s[0:3], s33 offset:112
	v_mov_b32_e32 v0, s23
	buffer_store_dword v0, off, s[0:3], s33 offset:108
	v_mov_b32_e32 v0, s22
	buffer_store_dword v0, off, s[0:3], s33 offset:104
	v_mov_b32_e32 v0, s21
	buffer_store_dword v0, off, s[0:3], s33 offset:100
	v_mov_b32_e32 v0, s20
	buffer_store_dword v0, off, s[0:3], s33 offset:96
	s_load_dwordx4 s[16:19], s[24:25], 0x0
	s_load_dwordx4 s[20:23], s[26:27], 0x0
	s_getpc_b64 s[24:25]
	s_add_u32 s24, s24, __const.GenValue.array@rel32@lo+52
	s_addc_u32 s25, s25, __const.GenValue.array@rel32@hi+60
	s_getpc_b64 s[26:27]
	s_add_u32 s26, s26, __const.GenValue.array@rel32@lo+36
	s_addc_u32 s27, s27, __const.GenValue.array@rel32@hi+44
	v_writelane_b32 v40, s31, 1
	s_waitcnt lgkmcnt(0)
	v_mov_b32_e32 v0, s19
	buffer_store_dword v0, off, s[0:3], s33 offset:92
	v_mov_b32_e32 v0, s18
	buffer_store_dword v0, off, s[0:3], s33 offset:88
	v_mov_b32_e32 v0, s17
	buffer_store_dword v0, off, s[0:3], s33 offset:84
	v_mov_b32_e32 v0, s16
	buffer_store_dword v0, off, s[0:3], s33 offset:80
	v_mov_b32_e32 v0, s23
	buffer_store_dword v0, off, s[0:3], s33 offset:76
	v_mov_b32_e32 v0, s22
	buffer_store_dword v0, off, s[0:3], s33 offset:72
	v_mov_b32_e32 v0, s21
	buffer_store_dword v0, off, s[0:3], s33 offset:68
	v_mov_b32_e32 v0, s20
	buffer_store_dword v0, off, s[0:3], s33 offset:64
	s_load_dwordx4 s[16:19], s[24:25], 0x0
	s_load_dwordx4 s[20:23], s[26:27], 0x0
	s_getpc_b64 s[24:25]
	s_add_u32 s24, s24, __const.GenValue.array@rel32@lo+20
	s_addc_u32 s25, s25, __const.GenValue.array@rel32@hi+28
	s_getpc_b64 s[26:27]
	s_add_u32 s26, s26, __const.GenValue.array@rel32@lo+4
	s_addc_u32 s27, s27, __const.GenValue.array@rel32@hi+12
	s_waitcnt lgkmcnt(0)
	v_mov_b32_e32 v0, s19
	buffer_store_dword v0, off, s[0:3], s33 offset:60
	v_mov_b32_e32 v0, s18
	buffer_store_dword v0, off, s[0:3], s33 offset:56
	v_mov_b32_e32 v0, s17
	buffer_store_dword v0, off, s[0:3], s33 offset:52
	v_mov_b32_e32 v0, s16
	buffer_store_dword v0, off, s[0:3], s33 offset:48
	v_mov_b32_e32 v0, s23
	buffer_store_dword v0, off, s[0:3], s33 offset:44
	v_mov_b32_e32 v0, s22
	buffer_store_dword v0, off, s[0:3], s33 offset:40
	v_mov_b32_e32 v0, s21
	buffer_store_dword v0, off, s[0:3], s33 offset:36
	v_mov_b32_e32 v0, s20
	buffer_store_dword v0, off, s[0:3], s33 offset:32
	s_load_dwordx4 s[16:19], s[24:25], 0x0
	s_load_dwordx4 s[20:23], s[26:27], 0x0
	s_waitcnt lgkmcnt(0)
	v_mov_b32_e32 v0, s19
	buffer_store_dword v0, off, s[0:3], s33 offset:28
	v_mov_b32_e32 v0, s18
	buffer_store_dword v0, off, s[0:3], s33 offset:24
	v_mov_b32_e32 v0, s17
	buffer_store_dword v0, off, s[0:3], s33 offset:20
	v_mov_b32_e32 v0, s16
	buffer_store_dword v0, off, s[0:3], s33 offset:16
	v_mov_b32_e32 v0, s23
	buffer_store_dword v0, off, s[0:3], s33 offset:12
	v_mov_b32_e32 v0, s22
	buffer_store_dword v0, off, s[0:3], s33 offset:8
	v_mov_b32_e32 v0, s21
	buffer_store_dword v0, off, s[0:3], s33 offset:4
	v_mov_b32_e32 v0, s20
	buffer_store_dword v0, off, s[0:3], s33
	.loc	1 39 18                         ; lane-pc-vega20-kernel.cl:39:18
	v_mov_b32_e32 v0, 0
	s_getpc_b64 s[16:17]
	s_add_u32 s16, s16, _Z13get_global_idj@rel32@lo+4
	s_addc_u32 s17, s17, _Z13get_global_idj@rel32@hi+12
	s_swappc_b64 s[30:31], s[16:17]
	.loc	1 42 15                         ; lane-pc-vega20-kernel.cl:42:15
	v_and_b32_e32 v1, 1, v0
	v_cmp_eq_u32_e32 vcc, 1, v1
	v_mov_b32_e32 v1, 0
	s_xor_b64 s[6:7], vcc, -1
	.loc	1 45 27                         ; lane-pc-vega20-kernel.cl:45:27
	v_lshlrev_b64 v[1:2], 2, v[0:1]
	s_getpc_b64 s[4:5]
	s_add_u32 s4, s4, const_array@rel32@lo+4
	s_addc_u32 s5, s5, const_array@rel32@hi+12
	v_mov_b32_e32 v3, s5
	v_add_co_u32_e32 v1, vcc, s4, v1
	s_getpc_b64 s[4:5]
	s_add_u32 s4, s4, const_array@rel32@lo+44
	s_addc_u32 s5, s5, const_array@rel32@hi+52
	v_addc_co_u32_e32 v2, vcc, v2, v3, vcc
	.loc	1 0 0 is_stmt 0                 ; lane-pc-vega20-kernel.cl:0:0
	v_mov_b32_e32 v3, s4
	v_mov_b32_e32 v4, s5
	s_and_saveexec_b64 s[4:5], s[6:7]
; %bb.1:                                ; %if.else
	v_mov_b32_e32 v4, v2
	v_mov_b32_e32 v3, v1
; %bb.2:                                ; %if.end
VPC_LINE_42:
	s_or_b64 exec, exec, s[4:5]
	s_mov_b32 s4, 0xb6db6db7
	.loc	1 47 15 is_stmt 1               ; lane-pc-vega20-kernel.cl:47:15
	v_mul_lo_u32 v5, v0, s4
	s_mov_b32 s4, 0x24924924
	v_add_u32_e32 v5, 0x49249249, v5
	v_cmp_lt_u32_e32 vcc, s4, v5
	s_and_saveexec_b64 s[4:5], vcc
; %bb.3:                                ; %if.else8
	.loc	1 48 15 is_stmt 1                ; lane-pc-vega20-kernel.cl:48:15
	s_getpc_b64 s[6:7]
	s_add_u32 s6, s6, const_array@rel32@lo+52
	s_addc_u32 s7, s7, const_array@rel32@hi+60
	v_mov_b32_e32 v1, s6
	v_mov_b32_e32 v2, s7
; %bb.4:                                ; %if.end9
VPC_LINE_47:
	s_or_b64 exec, exec, s[4:5]
	.loc	1 52 23 is_stmt 1               ; lane-pc-vega20-kernel.cl:52:23
	global_load_dword v5, v[3:4], off
	.loc	1 52 45 is_stmt 0               ; lane-pc-vega20-kernel.cl:52:45
	global_load_dword v6, v[1:2], off
	v_lshrrev_b32_e64 v1, 6, s33
	.loc	1 52 10                         ; lane-pc-vega20-kernel.cl:52:10
	v_lshl_add_u32 v0, v0, 2, v1
	buffer_load_dword v0, v0, s[0:3], 0 offen
	.loc	1 52 3                          ; lane-pc-vega20-kernel.cl:52:3
	v_readlane_b32 s4, v40, 0
	v_readlane_b32 s5, v40, 1
	s_sub_u32 s32, s32, 0x2400
	v_readlane_b32 s33, v40, 6
	.loc	1 52 43                         ; lane-pc-vega20-kernel.cl:52:43
	s_waitcnt vmcnt(1)
	v_mul_lo_u32 v1, v5, v6
	.loc	1 52 21                         ; lane-pc-vega20-kernel.cl:52:21
	s_waitcnt vmcnt(0)
	v_add_u32_e32 v0, v0, v1
	.cfi_def_cfa_register 64
	s_or_saveexec_b64 s[6:7], -1
	buffer_load_dword v40, off, s[0:3], s32 offset:128 ; 4-byte Folded Reload
	s_mov_b64 exec, s[6:7]
	.loc	1 52 3                          ; lane-pc-vega20-kernel.cl:52:3
	s_waitcnt vmcnt(0)
	s_setpc_b64 s[4:5]
.Ltmp1:
Lfunc_end2:
	.size	GenValue, Lfunc_end2-GenValue
	.cfi_endproc
                                        ; -- End function
	.section	.AMDGPU.csdata
; Function info:
; codeLenInByte = 1028
; NumSgprs: 38
; NumVgprs: 41
; ScratchSize: 160
; MemoryBound: 0
	.text
	.hidden ChangeLocalContent
	.globl	ChangeLocalContent              ; -- Begin function ChangeLocalContent
	.hidden Lfunc_end3
	.globl	Lfunc_end3
	.hidden Llist_start1
	.globl	Llist_start1
	.hidden VPC_LINE_61
	.globl	VPC_LINE_61
	.hidden VPC_LINE_66
	.globl	VPC_LINE_66
	.hidden VPC_LINE_71
	.globl	VPC_LINE_71
	.hidden VPC_LINE_71_2
	.globl	VPC_LINE_71_2
	.p2align	2
	.type	ChangeLocalContent,@function
ChangeLocalContent:                     ; @ChangeLocalContent
ChangeLocalContent$local:
.Lfunc_begin3:
	.loc	1 56 0 is_stmt 1                ; lane-pc-vega20-kernel.cl:56:0
	.cfi_startproc
; %bb.0:                                ; %entry
	.cfi_llvm_def_aspace_cfa 64, 0, 6
	.cfi_escape 0x10, 0x10, 0x08, 0x90, 0x3e, 0x93, 0x04, 0x90, 0x3f, 0x93, 0x04 ;
	.cfi_undefined 2560
	.cfi_undefined 2561
	.cfi_undefined 2562
	.cfi_undefined 2563
	.cfi_undefined 2564
	.cfi_undefined 2565
	.cfi_undefined 2566
	.cfi_undefined 2567
	.cfi_undefined 2568
	.cfi_undefined 2569
	.cfi_undefined 2570
	.cfi_undefined 2571
	.cfi_undefined 2572
	.cfi_undefined 2573
	.cfi_undefined 2574
	.cfi_undefined 2575
	.cfi_undefined 2576
	.cfi_undefined 2577
	.cfi_undefined 2578
	.cfi_undefined 2579
	.cfi_undefined 2580
	.cfi_undefined 2581
	.cfi_undefined 2582
	.cfi_undefined 2583
	.cfi_undefined 2584
	.cfi_undefined 2585
	.cfi_undefined 2586
	.cfi_undefined 2587
	.cfi_undefined 2588
	.cfi_undefined 2589
	.cfi_undefined 2590
	.cfi_undefined 2591
	.cfi_undefined 2592
	.cfi_undefined 2593
	.cfi_undefined 2594
	.cfi_undefined 2595
	.cfi_undefined 2596
	.cfi_undefined 2597
	.cfi_undefined 2598
	.cfi_undefined 2599
	.cfi_undefined 2608
	.cfi_undefined 2609
	.cfi_undefined 2610
	.cfi_undefined 2611
	.cfi_undefined 2612
	.cfi_undefined 2613
	.cfi_undefined 2614
	.cfi_undefined 2615
	.cfi_undefined 2624
	.cfi_undefined 2625
	.cfi_undefined 2626
	.cfi_undefined 2627
	.cfi_undefined 2628
	.cfi_undefined 2629
	.cfi_undefined 2630
	.cfi_undefined 2631
	.cfi_undefined 2640
	.cfi_undefined 2641
	.cfi_undefined 2642
	.cfi_undefined 2643
	.cfi_undefined 2644
	.cfi_undefined 2645
	.cfi_undefined 2646
	.cfi_undefined 2647
	.cfi_undefined 2656
	.cfi_undefined 2657
	.cfi_undefined 2658
	.cfi_undefined 2659
	.cfi_undefined 2660
	.cfi_undefined 2661
	.cfi_undefined 2662
	.cfi_undefined 2663
	.cfi_undefined 2672
	.cfi_undefined 2673
	.cfi_undefined 2674
	.cfi_undefined 2675
	.cfi_undefined 2676
	.cfi_undefined 2677
	.cfi_undefined 2678
	.cfi_undefined 2679
	.cfi_undefined 2688
	.cfi_undefined 2689
	.cfi_undefined 2690
	.cfi_undefined 2691
	.cfi_undefined 2692
	.cfi_undefined 2693
	.cfi_undefined 2694
	.cfi_undefined 2695
	.cfi_undefined 2704
	.cfi_undefined 2705
	.cfi_undefined 2706
	.cfi_undefined 2707
	.cfi_undefined 2708
	.cfi_undefined 2709
	.cfi_undefined 2710
	.cfi_undefined 2711
	.cfi_undefined 2720
	.cfi_undefined 2721
	.cfi_undefined 2722
	.cfi_undefined 2723
	.cfi_undefined 2724
	.cfi_undefined 2725
	.cfi_undefined 2726
	.cfi_undefined 2727
	.cfi_undefined 2736
	.cfi_undefined 2737
	.cfi_undefined 2738
	.cfi_undefined 2739
	.cfi_undefined 2740
	.cfi_undefined 2741
	.cfi_undefined 2742
	.cfi_undefined 2743
	.cfi_undefined 2752
	.cfi_undefined 2753
	.cfi_undefined 2754
	.cfi_undefined 2755
	.cfi_undefined 2756
	.cfi_undefined 2757
	.cfi_undefined 2758
	.cfi_undefined 2759
	.cfi_undefined 2768
	.cfi_undefined 2769
	.cfi_undefined 2770
	.cfi_undefined 2771
	.cfi_undefined 2772
	.cfi_undefined 2773
	.cfi_undefined 2774
	.cfi_undefined 2775
	.cfi_undefined 2784
	.cfi_undefined 2785
	.cfi_undefined 2786
	.cfi_undefined 2787
	.cfi_undefined 2788
	.cfi_undefined 2789
	.cfi_undefined 2790
	.cfi_undefined 2791
	.cfi_undefined 2800
	.cfi_undefined 2801
	.cfi_undefined 2802
	.cfi_undefined 2803
	.cfi_undefined 2804
	.cfi_undefined 2805
	.cfi_undefined 2806
	.cfi_undefined 2807
	.cfi_undefined 32
	.cfi_undefined 33
	.cfi_undefined 34
	.cfi_undefined 35
	.cfi_undefined 36
	.cfi_undefined 37
	.cfi_undefined 38
	.cfi_undefined 39
	.cfi_undefined 40
	.cfi_undefined 41
	.cfi_undefined 42
	.cfi_undefined 43
	.cfi_undefined 44
	.cfi_undefined 45
	.cfi_undefined 46
	.cfi_undefined 47
	.cfi_undefined 48
	.cfi_undefined 49
	.cfi_undefined 50
	.cfi_undefined 51
	.cfi_undefined 52
	.cfi_undefined 53
	.cfi_undefined 54
	.cfi_undefined 55
	.cfi_undefined 56
	.cfi_undefined 57
	.cfi_undefined 58
	.cfi_undefined 59
	.cfi_undefined 60
	.cfi_undefined 61
	.cfi_undefined 62
	.cfi_undefined 63
	s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
	s_or_saveexec_b64 s[16:17], -1
	buffer_store_dword v44, off, s[0:3], s32 offset:144 ; 4-byte Folded Spill
	.cfi_offset 2604, 9216
	s_mov_b64 exec, s[16:17]
	v_writelane_b32 v44, s30, 15
	v_writelane_b32 v44, s31, 16
	.cfi_escape 0x10, 0x10, 0x0e, 0x90, 0xac, 0x14, 0x9d, 0x20, 0xe0, 0x03, 0x90, 0xac, 0x14, 0x9d, 0x20, 0x80, 0x04 ;
	v_writelane_b32 v44, exec_lo, 17
	v_writelane_b32 v44, exec_hi, 18
	.cfi_escape 0x10, 0x11, 0x0e, 0x90, 0xac, 0x14, 0x9d, 0x20, 0xa0, 0x04, 0x90, 0xac, 0x14, 0x9d, 0x20, 0xc0, 0x04 ;
	v_writelane_b32 v44, s33, 19
	.cfi_escape 0x10, 0x41, 0x05, 0x90, 0xac, 0x14, 0xe4, 0x4c ;
	s_mov_b32 s33, s32
	.cfi_def_cfa_register 65
	v_writelane_b32 v44, s34, 0
	s_add_u32 s32, s32, 0x2800
	buffer_store_dword v40, off, s[0:3], s33 offset:12 ; 4-byte Folded Spill
	buffer_store_dword v41, off, s[0:3], s33 offset:8 ; 4-byte Folded Spill
	buffer_store_dword v42, off, s[0:3], s33 offset:4 ; 4-byte Folded Spill
	buffer_store_dword v43, off, s[0:3], s33 ; 4-byte Folded Spill
	.cfi_escape 0x10, 0x42, 0x05, 0x90, 0xac, 0x14, 0xe4, 0x00 ;
	v_writelane_b32 v44, s35, 1
	.cfi_escape 0x10, 0x43, 0x05, 0x90, 0xac, 0x14, 0xe4, 0x04 ;
	v_writelane_b32 v44, s36, 2
	.cfi_escape 0x10, 0x44, 0x05, 0x90, 0xac, 0x14, 0xe4, 0x08 ;
	v_writelane_b32 v44, s38, 3
	.cfi_escape 0x10, 0x46, 0x05, 0x90, 0xac, 0x14, 0xe4, 0x0c ;
	v_writelane_b32 v44, s39, 4
	.cfi_escape 0x10, 0x47, 0x05, 0x90, 0xac, 0x14, 0xe4, 0x10 ;
	v_writelane_b32 v44, s40, 5
	.cfi_escape 0x10, 0x48, 0x05, 0x90, 0xac, 0x14, 0xe4, 0x14 ;
	v_writelane_b32 v44, s41, 6
	.cfi_escape 0x10, 0x49, 0x05, 0x90, 0xac, 0x14, 0xe4, 0x18 ;
	v_writelane_b32 v44, s42, 7
	.cfi_escape 0x10, 0x4a, 0x05, 0x90, 0xac, 0x14, 0xe4, 0x1c ;
	v_writelane_b32 v44, s43, 8
	.cfi_escape 0x10, 0x4b, 0x05, 0x90, 0xac, 0x14, 0xe4, 0x20 ;
	v_writelane_b32 v44, s44, 9
	.cfi_escape 0x10, 0x4c, 0x05, 0x90, 0xac, 0x14, 0xe4, 0x24 ;
	v_writelane_b32 v44, s45, 10
	.cfi_escape 0x10, 0x4d, 0x05, 0x90, 0xac, 0x14, 0xe4, 0x28 ;
	v_writelane_b32 v44, s46, 11
	.cfi_escape 0x10, 0x4e, 0x05, 0x90, 0xac, 0x14, 0xe4, 0x2c ;
	v_writelane_b32 v44, s47, 12
	.cfi_escape 0x10, 0x4f, 0x05, 0x90, 0xac, 0x14, 0xe4, 0x30 ;
	s_mov_b32 s35, s13
	s_mov_b32 s36, s12
	s_getpc_b64 s[12:13]
	s_add_u32 s12, s12, __const.ChangeLocalContent.array@rel32@lo+116
	s_addc_u32 s13, s13, __const.ChangeLocalContent.array@rel32@hi+124
	s_mov_b32 s34, s14
	s_getpc_b64 s[14:15]
	s_add_u32 s14, s14, __const.ChangeLocalContent.array@rel32@lo+100
	s_addc_u32 s15, s15, __const.ChangeLocalContent.array@rel32@hi+108
	s_mov_b64 s[38:39], s[10:11]
	s_mov_b64 s[40:41], s[8:9]
	s_mov_b64 s[42:43], s[6:7]
	s_mov_b64 s[44:45], s[4:5]
.Ltmp2:
	.loc	1 57 7 prologue_end             ; lane-pc-vega20-kernel.cl:57:7
	s_load_dwordx4 s[4:7], s[12:13], 0x0
	s_load_dwordx4 s[8:11], s[14:15], 0x0
	v_mov_b32_e32 v41, v0
	s_getpc_b64 s[12:13]
	s_add_u32 s12, s12, __const.ChangeLocalContent.array@rel32@lo+84
	s_addc_u32 s13, s13, __const.ChangeLocalContent.array@rel32@hi+92
	s_getpc_b64 s[14:15]
	s_add_u32 s14, s14, __const.ChangeLocalContent.array@rel32@lo+68
	s_addc_u32 s15, s15, __const.ChangeLocalContent.array@rel32@hi+76
	s_waitcnt lgkmcnt(0)
	v_mov_b32_e32 v0, s7
	buffer_store_dword v0, off, s[0:3], s33 offset:140
	v_mov_b32_e32 v0, s6
	buffer_store_dword v0, off, s[0:3], s33 offset:136
	v_mov_b32_e32 v0, s5
	buffer_store_dword v0, off, s[0:3], s33 offset:132
	v_mov_b32_e32 v0, s4
	buffer_store_dword v0, off, s[0:3], s33 offset:128
	v_mov_b32_e32 v0, s11
	buffer_store_dword v0, off, s[0:3], s33 offset:124
	v_mov_b32_e32 v0, s10
	buffer_store_dword v0, off, s[0:3], s33 offset:120
	v_mov_b32_e32 v0, s9
	buffer_store_dword v0, off, s[0:3], s33 offset:116
	v_mov_b32_e32 v0, s8
	buffer_store_dword v0, off, s[0:3], s33 offset:112
	s_load_dwordx4 s[4:7], s[12:13], 0x0
	s_load_dwordx4 s[8:11], s[14:15], 0x0
	s_getpc_b64 s[12:13]
	s_add_u32 s12, s12, __const.ChangeLocalContent.array@rel32@lo+52
	s_addc_u32 s13, s13, __const.ChangeLocalContent.array@rel32@hi+60
	s_getpc_b64 s[14:15]
	s_add_u32 s14, s14, __const.ChangeLocalContent.array@rel32@lo+36
	s_addc_u32 s15, s15, __const.ChangeLocalContent.array@rel32@hi+44
	v_writelane_b32 v44, s30, 13
	s_waitcnt lgkmcnt(0)
	v_mov_b32_e32 v0, s7
	buffer_store_dword v0, off, s[0:3], s33 offset:108
	v_mov_b32_e32 v0, s6
	buffer_store_dword v0, off, s[0:3], s33 offset:104
	v_mov_b32_e32 v0, s5
	buffer_store_dword v0, off, s[0:3], s33 offset:100
	v_mov_b32_e32 v0, s4
	buffer_store_dword v0, off, s[0:3], s33 offset:96
	v_mov_b32_e32 v0, s11
	buffer_store_dword v0, off, s[0:3], s33 offset:92
	v_mov_b32_e32 v0, s10
	buffer_store_dword v0, off, s[0:3], s33 offset:88
	v_mov_b32_e32 v0, s9
	buffer_store_dword v0, off, s[0:3], s33 offset:84
	v_mov_b32_e32 v0, s8
	buffer_store_dword v0, off, s[0:3], s33 offset:80
	s_load_dwordx4 s[4:7], s[12:13], 0x0
	s_load_dwordx4 s[8:11], s[14:15], 0x0
	s_getpc_b64 s[12:13]
	s_add_u32 s12, s12, __const.ChangeLocalContent.array@rel32@lo+20
	s_addc_u32 s13, s13, __const.ChangeLocalContent.array@rel32@hi+28
	s_getpc_b64 s[14:15]
	s_add_u32 s14, s14, __const.ChangeLocalContent.array@rel32@lo+4
	s_addc_u32 s15, s15, __const.ChangeLocalContent.array@rel32@hi+12
	v_writelane_b32 v44, s31, 14
	s_waitcnt lgkmcnt(0)
	v_mov_b32_e32 v0, s7
	buffer_store_dword v0, off, s[0:3], s33 offset:76
	v_mov_b32_e32 v0, s6
	buffer_store_dword v0, off, s[0:3], s33 offset:72
	v_mov_b32_e32 v0, s5
	buffer_store_dword v0, off, s[0:3], s33 offset:68
	v_mov_b32_e32 v0, s4
	buffer_store_dword v0, off, s[0:3], s33 offset:64
	v_mov_b32_e32 v0, s11
	buffer_store_dword v0, off, s[0:3], s33 offset:60
	v_mov_b32_e32 v0, s10
	buffer_store_dword v0, off, s[0:3], s33 offset:56
	v_mov_b32_e32 v0, s9
	buffer_store_dword v0, off, s[0:3], s33 offset:52
	v_mov_b32_e32 v0, s8
	buffer_store_dword v0, off, s[0:3], s33 offset:48
	s_load_dwordx4 s[4:7], s[12:13], 0x0
	s_load_dwordx4 s[8:11], s[14:15], 0x0
Llist_start1:
	s_mov_b32 s12, s36
	s_mov_b32 s13, s35
	s_mov_b32 s14, s34
	s_waitcnt lgkmcnt(0)
	v_mov_b32_e32 v0, s7
	buffer_store_dword v0, off, s[0:3], s33 offset:44
	v_mov_b32_e32 v0, s6
	buffer_store_dword v0, off, s[0:3], s33 offset:40
	v_mov_b32_e32 v0, s5
	buffer_store_dword v0, off, s[0:3], s33 offset:36
	v_mov_b32_e32 v0, s4
	buffer_store_dword v0, off, s[0:3], s33 offset:32
	v_mov_b32_e32 v0, s11
	buffer_store_dword v0, off, s[0:3], s33 offset:28
	v_mov_b32_e32 v0, s10
	buffer_store_dword v0, off, s[0:3], s33 offset:24
	v_mov_b32_e32 v0, s9
	buffer_store_dword v0, off, s[0:3], s33 offset:20
	v_mov_b32_e32 v0, s8
	buffer_store_dword v0, off, s[0:3], s33 offset:16
	v_mov_b32_e32 v0, 0
	.loc	1 59 18                         ; lane-pc-vega20-kernel.cl:59:18
	s_mov_b64 s[4:5], s[44:45]
	s_mov_b64 s[6:7], s[42:43]
	s_mov_b64 s[8:9], s[40:41]
	s_mov_b64 s[10:11], s[38:39]
	s_getpc_b64 s[16:17]
	s_add_u32 s16, s16, _Z13get_global_idj@rel32@lo+4
	s_addc_u32 s17, s17, _Z13get_global_idj@rel32@hi+12
	v_mov_b32_e32 v40, v31
	s_swappc_b64 s[30:31], s[16:17]
	v_mov_b32_e32 v42, v0
	.loc	1 61 11                         ; lane-pc-vega20-kernel.cl:61:11
	v_and_b32_e32 v0, 3, v42
	v_cmp_eq_u32_e32 vcc, 0, v0
	s_and_saveexec_b64 s[4:5], vcc
	s_xor_b64 s[46:47], exec, s[4:5]
	s_cbranch_execz BB3_2
; %bb.1:                                ; %if.else13
	.loc	1 79 25                         ; lane-pc-vega20-kernel.cl:79:25
	s_mov_b64 s[4:5], s[44:45]
	s_mov_b64 s[6:7], s[42:43]
	s_mov_b64 s[8:9], s[40:41]
	s_mov_b64 s[10:11], s[38:39]
	s_mov_b32 s12, s36
	s_mov_b32 s13, s35
	s_mov_b32 s14, s34
	v_mov_b32_e32 v31, v40
	s_getpc_b64 s[16:17]
	s_add_u32 s16, s16, GenValue@rel32@lo+4
	s_addc_u32 s17, s17, GenValue@rel32@hi+12
	s_swappc_b64 s[30:31], s[16:17]
	.loc	1 79 23 is_stmt 0               ; lane-pc-vega20-kernel.cl:79:23
	ds_write_b32 v41, v0
BB3_2:                                  ; %Flow8
	.loc	1 0 23                          ; lane-pc-vega20-kernel.cl:0:23
	s_or_saveexec_b64 s[4:5], s[46:47]
	s_xor_b64 exec, exec, s[4:5]
VPC_LINE_61:
	s_cbranch_execz BB3_10
; %bb.3:                                ; %if.then
	.loc	1 66 17 is_stmt 1               ; lane-pc-vega20-kernel.cl:66:17
	v_and_b32_e32 v0, 1, v42
	v_cmp_eq_u32_e32 vcc, 1, v0
	s_xor_b64 s[8:9], vcc, -1
	s_getpc_b64 s[6:7]
	s_add_u32 s6, s6, const_array@rel32@lo+44
	s_addc_u32 s7, s7, const_array@rel32@hi+52
	.loc	1 0 0 is_stmt 0                 ; lane-pc-vega20-kernel.cl:0:0
	v_mov_b32_e32 v0, s6
	v_mov_b32_e32 v1, s7
	s_and_saveexec_b64 s[6:7], s[8:9]
; %bb.4:                                ; %if.else
	.loc	1 69 29 is_stmt 1               ; lane-pc-vega20-kernel.cl:69:29
	v_mov_b32_e32 v43, 0
	v_lshlrev_b64 v[0:1], 2, v[42:43]
	s_getpc_b64 s[8:9]
	s_add_u32 s8, s8, const_array@rel32@lo+4
	s_addc_u32 s9, s9, const_array@rel32@hi+12
	v_mov_b32_e32 v2, s9
	v_add_co_u32_e32 v0, vcc, s8, v0
	v_addc_co_u32_e32 v1, vcc, v1, v2, vcc
; %bb.5:                                ; %if.end
VPC_LINE_66:
	.loc	1 0 29 is_stmt 0                ; lane-pc-vega20-kernel.cl:0:29
	s_or_b64 exec, exec, s[6:7]
	s_mov_b32 s6, 0xaaaaaaab
	.loc	1 71 17 is_stmt 1               ; lane-pc-vega20-kernel.cl:71:17
	v_mul_lo_u32 v2, v42, s6
	s_mov_b32 s6, 0x55555554
	v_add_u32_e32 v2, 0x55555555, v2
	v_cmp_lt_u32_e32 vcc, s6, v2
                                        ; implicit-def: $vgpr2
	s_and_saveexec_b64 s[6:7], vcc
	s_xor_b64 s[6:7], exec, s[6:7]
	s_cbranch_execz BB3_7
; %bb.6:                                ; %if.else10
	.loc	1 74 14                         ; lane-pc-vega20-kernel.cl:74:14
	buffer_load_dword v2, off, s[0:3], s33 offset:56
BB3_7:                                  ; %Flow
	.loc	1 0 14 is_stmt 0                ; lane-pc-vega20-kernel.cl:0:14
VPC_LINE_71:
	s_or_saveexec_b64 s[6:7], s[6:7]
	s_xor_b64 exec, exec, s[6:7]
	s_cbranch_execz BB3_9
; %bb.8:                                ; %if.then7
	s_waitcnt vmcnt(0)
	v_lshrrev_b32_e64 v2, 6, s33
	v_add_u32_e32 v2, 16, v2
	.loc	1 72 14 is_stmt 1               ; lane-pc-vega20-kernel.cl:72:14
	v_lshl_add_u32 v2, v42, 2, v2
	buffer_load_dword v2, v2, s[0:3], 0 offen
BB3_9:                                  ; %if.end12
VPC_LINE_71_2:
	.loc	1 0 14 is_stmt 0                ; lane-pc-vega20-kernel.cl:0:14
	s_or_b64 exec, exec, s[6:7]
	.loc	1 76 32 is_stmt 1               ; lane-pc-vega20-kernel.cl:76:32
	global_load_dword v0, v[0:1], off
	.loc	1 76 30 is_stmt 0               ; lane-pc-vega20-kernel.cl:76:30
	s_waitcnt vmcnt(0)
	v_add_u32_e32 v0, v2, v0
	.loc	1 76 23                         ; lane-pc-vega20-kernel.cl:76:23
	ds_write_b32 v41, v0
BB3_10:                                 ; %Flow9
	.loc	1 0 23                          ; lane-pc-vega20-kernel.cl:0:23
	s_or_b64 exec, exec, s[4:5]
	.loc	1 80 1 is_stmt 1                ; lane-pc-vega20-kernel.cl:80:1
	buffer_load_dword v43, off, s[0:3], s33 ; 4-byte Folded Reload
	buffer_load_dword v42, off, s[0:3], s33 offset:4 ; 4-byte Folded Reload
	buffer_load_dword v41, off, s[0:3], s33 offset:8 ; 4-byte Folded Reload
	buffer_load_dword v40, off, s[0:3], s33 offset:12 ; 4-byte Folded Reload
	v_readlane_b32 s4, v44, 13
	v_readlane_b32 s5, v44, 14
	v_readlane_b32 s47, v44, 12
	v_readlane_b32 s46, v44, 11
	v_readlane_b32 s45, v44, 10
	v_readlane_b32 s44, v44, 9
	v_readlane_b32 s43, v44, 8
	v_readlane_b32 s42, v44, 7
	v_readlane_b32 s41, v44, 6
	v_readlane_b32 s40, v44, 5
	v_readlane_b32 s39, v44, 4
	v_readlane_b32 s38, v44, 3
	v_readlane_b32 s36, v44, 2
	v_readlane_b32 s35, v44, 1
	v_readlane_b32 s34, v44, 0
	s_sub_u32 s32, s32, 0x2800
	v_readlane_b32 s33, v44, 19
	.cfi_def_cfa_register 64
	s_or_saveexec_b64 s[6:7], -1
	buffer_load_dword v44, off, s[0:3], s32 offset:144 ; 4-byte Folded Reload
	s_mov_b64 exec, s[6:7]
	s_waitcnt vmcnt(0) lgkmcnt(0)
	s_setpc_b64 s[4:5]
.Ltmp3:
Lfunc_end3:
	.size	ChangeLocalContent, Lfunc_end3-ChangeLocalContent
	.cfi_endproc
                                        ; -- End function
	.section	.AMDGPU.csdata
; Function info:
; codeLenInByte = 1456
; NumSgprs: 52
; NumVgprs: 45
; ScratchSize: 320
; MemoryBound: 0
	.text
	.hidden SendResults
	.globl	SendResults                     ; -- Begin function SendResults
	.hidden Lfunc_end4
	.globl	Lfunc_end4
	.hidden VPC_LINE_86
	.globl	VPC_LINE_86
	.p2align	2
	.type	SendResults,@function
SendResults:                            ; @SendResults
SendResults$local:
.Lfunc_begin4:
	.loc	1 83 0                          ; lane-pc-vega20-kernel.cl:83:0
	.cfi_startproc
; %bb.0:                                ; %entry
	.cfi_llvm_def_aspace_cfa 64, 0, 6
	.cfi_escape 0x10, 0x10, 0x08, 0x90, 0x3e, 0x93, 0x04, 0x90, 0x3f, 0x93, 0x04 ;
	.cfi_undefined 2560
	.cfi_undefined 2561
	.cfi_undefined 2562
	.cfi_undefined 2563
	.cfi_undefined 2564
	.cfi_undefined 2565
	.cfi_undefined 2566
	.cfi_undefined 2567
	.cfi_undefined 2568
	.cfi_undefined 2569
	.cfi_undefined 2570
	.cfi_undefined 2571
	.cfi_undefined 2572
	.cfi_undefined 2573
	.cfi_undefined 2574
	.cfi_undefined 2575
	.cfi_undefined 2576
	.cfi_undefined 2577
	.cfi_undefined 2578
	.cfi_undefined 2579
	.cfi_undefined 2580
	.cfi_undefined 2581
	.cfi_undefined 2582
	.cfi_undefined 2583
	.cfi_undefined 2584
	.cfi_undefined 2585
	.cfi_undefined 2586
	.cfi_undefined 2587
	.cfi_undefined 2588
	.cfi_undefined 2589
	.cfi_undefined 2590
	.cfi_undefined 2591
	.cfi_undefined 2592
	.cfi_undefined 2593
	.cfi_undefined 2594
	.cfi_undefined 2595
	.cfi_undefined 2596
	.cfi_undefined 2597
	.cfi_undefined 2598
	.cfi_undefined 2599
	.cfi_undefined 2608
	.cfi_undefined 2609
	.cfi_undefined 2610
	.cfi_undefined 2611
	.cfi_undefined 2612
	.cfi_undefined 2613
	.cfi_undefined 2614
	.cfi_undefined 2615
	.cfi_undefined 2624
	.cfi_undefined 2625
	.cfi_undefined 2626
	.cfi_undefined 2627
	.cfi_undefined 2628
	.cfi_undefined 2629
	.cfi_undefined 2630
	.cfi_undefined 2631
	.cfi_undefined 2640
	.cfi_undefined 2641
	.cfi_undefined 2642
	.cfi_undefined 2643
	.cfi_undefined 2644
	.cfi_undefined 2645
	.cfi_undefined 2646
	.cfi_undefined 2647
	.cfi_undefined 2656
	.cfi_undefined 2657
	.cfi_undefined 2658
	.cfi_undefined 2659
	.cfi_undefined 2660
	.cfi_undefined 2661
	.cfi_undefined 2662
	.cfi_undefined 2663
	.cfi_undefined 2672
	.cfi_undefined 2673
	.cfi_undefined 2674
	.cfi_undefined 2675
	.cfi_undefined 2676
	.cfi_undefined 2677
	.cfi_undefined 2678
	.cfi_undefined 2679
	.cfi_undefined 2688
	.cfi_undefined 2689
	.cfi_undefined 2690
	.cfi_undefined 2691
	.cfi_undefined 2692
	.cfi_undefined 2693
	.cfi_undefined 2694
	.cfi_undefined 2695
	.cfi_undefined 2704
	.cfi_undefined 2705
	.cfi_undefined 2706
	.cfi_undefined 2707
	.cfi_undefined 2708
	.cfi_undefined 2709
	.cfi_undefined 2710
	.cfi_undefined 2711
	.cfi_undefined 2720
	.cfi_undefined 2721
	.cfi_undefined 2722
	.cfi_undefined 2723
	.cfi_undefined 2724
	.cfi_undefined 2725
	.cfi_undefined 2726
	.cfi_undefined 2727
	.cfi_undefined 2736
	.cfi_undefined 2737
	.cfi_undefined 2738
	.cfi_undefined 2739
	.cfi_undefined 2740
	.cfi_undefined 2741
	.cfi_undefined 2742
	.cfi_undefined 2743
	.cfi_undefined 2752
	.cfi_undefined 2753
	.cfi_undefined 2754
	.cfi_undefined 2755
	.cfi_undefined 2756
	.cfi_undefined 2757
	.cfi_undefined 2758
	.cfi_undefined 2759
	.cfi_undefined 2768
	.cfi_undefined 2769
	.cfi_undefined 2770
	.cfi_undefined 2771
	.cfi_undefined 2772
	.cfi_undefined 2773
	.cfi_undefined 2774
	.cfi_undefined 2775
	.cfi_undefined 2784
	.cfi_undefined 2785
	.cfi_undefined 2786
	.cfi_undefined 2787
	.cfi_undefined 2788
	.cfi_undefined 2789
	.cfi_undefined 2790
	.cfi_undefined 2791
	.cfi_undefined 2800
	.cfi_undefined 2801
	.cfi_undefined 2802
	.cfi_undefined 2803
	.cfi_undefined 2804
	.cfi_undefined 2805
	.cfi_undefined 2806
	.cfi_undefined 2807
	.cfi_undefined 32
	.cfi_undefined 33
	.cfi_undefined 34
	.cfi_undefined 35
	.cfi_undefined 36
	.cfi_undefined 37
	.cfi_undefined 38
	.cfi_undefined 39
	.cfi_undefined 40
	.cfi_undefined 41
	.cfi_undefined 42
	.cfi_undefined 43
	.cfi_undefined 44
	.cfi_undefined 45
	.cfi_undefined 46
	.cfi_undefined 47
	.cfi_undefined 48
	.cfi_undefined 49
	.cfi_undefined 50
	.cfi_undefined 51
	.cfi_undefined 52
	.cfi_undefined 53
	.cfi_undefined 54
	.cfi_undefined 55
	.cfi_undefined 56
	.cfi_undefined 57
	.cfi_undefined 58
	.cfi_undefined 59
	.cfi_undefined 60
	.cfi_undefined 61
	.cfi_undefined 62
	.cfi_undefined 63
	s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
	s_or_saveexec_b64 s[16:17], -1
	buffer_store_dword v43, off, s[0:3], s32 offset:12 ; 4-byte Folded Spill
	.cfi_offset 2603, 768
	s_mov_b64 exec, s[16:17]
	v_writelane_b32 v43, s30, 2
	v_writelane_b32 v43, s31, 3
	.cfi_escape 0x10, 0x10, 0x0c, 0x90, 0xab, 0x14, 0x9d, 0x20, 0x40, 0x90, 0xab, 0x14, 0x9d, 0x20, 0x60 ;
	v_writelane_b32 v43, exec_lo, 4
	v_writelane_b32 v43, exec_hi, 5
	.cfi_escape 0x10, 0x11, 0x0e, 0x90, 0xab, 0x14, 0x9d, 0x20, 0x80, 0x01, 0x90, 0xab, 0x14, 0x9d, 0x20, 0xa0, 0x01 ;
	v_writelane_b32 v43, s33, 6
	.cfi_escape 0x10, 0x41, 0x05, 0x90, 0xab, 0x14, 0xe4, 0x18 ;
	s_mov_b32 s33, s32
	.cfi_def_cfa_register 65
	v_writelane_b32 v43, s30, 0
	buffer_store_dword v40, off, s[0:3], s33 offset:8 ; 4-byte Folded Spill
	buffer_store_dword v41, off, s[0:3], s33 offset:4 ; 4-byte Folded Spill
	buffer_store_dword v42, off, s[0:3], s33 ; 4-byte Folded Spill
	s_add_u32 s32, s32, 0x800
	v_mov_b32_e32 v42, v0
	v_mov_b32_e32 v0, 0
.Ltmp4:
	.loc	1 84 18 prologue_end            ; lane-pc-vega20-kernel.cl:84:18
	s_getpc_b64 s[16:17]
	s_add_u32 s16, s16, _Z13get_global_idj@rel32@lo+4
	s_addc_u32 s17, s17, _Z13get_global_idj@rel32@hi+12
	v_writelane_b32 v43, s31, 1
	v_mov_b32_e32 v41, v2
	v_mov_b32_e32 v40, v1
	s_swappc_b64 s[30:31], s[16:17]
	.loc	1 86 8                          ; lane-pc-vega20-kernel.cl:86:8
	v_cmp_eq_u32_e32 vcc, 0, v0
	s_and_saveexec_b64 s[4:5], vcc
	s_cbranch_execz BB4_2
; %bb.1:                                ; %if.then
	.loc	1 88 36                         ; lane-pc-vega20-kernel.cl:88:36
	ds_read_b32 v1, v42
	ds_read_u8 v2, v42 offset:4
	s_waitcnt lgkmcnt(1)
	v_add_u32_e32 v1, 32, v1
	s_waitcnt lgkmcnt(0)
	v_add_u32_e32 v2, 50, v2
	global_store_dword v[40:41], v1, off
	.loc	1 89 20                         ; lane-pc-vega20-kernel.cl:89:20
	global_store_byte v[40:41], v2, off offset:4
BB4_2:                                  ; %if.end
VPC_LINE_86:
	.loc	1 0 20 is_stmt 0                ; lane-pc-vega20-kernel.cl:0:20
	s_or_b64 exec, exec, s[4:5]
	.loc	1 92 26 is_stmt 1               ; lane-pc-vega20-kernel.cl:92:26
	v_mov_b32_e32 v1, 0
	.loc	1 92 59 is_stmt 0               ; lane-pc-vega20-kernel.cl:92:59
	v_lshlrev_b64 v[1:2], 2, v[0:1]
	s_getpc_b64 s[4:5]
	s_add_u32 s4, s4, const_struct@rel32@lo+12
	s_addc_u32 s5, s5, const_struct@rel32@hi+20
	v_mov_b32_e32 v4, s5
	v_add_co_u32_e32 v3, vcc, s4, v1
	v_addc_co_u32_e32 v4, vcc, v2, v4, vcc
	.loc	1 92 90                         ; lane-pc-vega20-kernel.cl:92:90
	s_getpc_b64 s[4:5]
	s_add_u32 s4, s4, const_array@rel32@lo+4
	s_addc_u32 s5, s5, const_array@rel32@hi+12
	.loc	1 92 59                         ; lane-pc-vega20-kernel.cl:92:59
	global_load_dword v5, v[3:4], off
	.loc	1 92 90                         ; lane-pc-vega20-kernel.cl:92:90
	v_mov_b32_e32 v4, s5
	v_add_co_u32_e32 v3, vcc, s4, v1
	v_addc_co_u32_e32 v4, vcc, v2, v4, vcc
	global_load_dword v3, v[3:4], off
	.loc	1 92 41                         ; lane-pc-vega20-kernel.cl:92:41
	v_lshl_add_u32 v0, v0, 2, v42
	.loc	1 92 26                         ; lane-pc-vega20-kernel.cl:92:26
	ds_read_b32 v0, v0 offset:8
	v_readlane_b32 s4, v43, 0
	v_readlane_b32 s5, v43, 1
	s_sub_u32 s32, s32, 0x800
	.loc	1 92 88                         ; lane-pc-vega20-kernel.cl:92:88
	s_waitcnt vmcnt(0) lgkmcnt(0)
	v_add3_u32 v3, v0, v5, v3
	.loc	1 92 8 is_stmt 0                ; lane-pc-vega20-kernel.cl:92:8
	v_add_co_u32_e32 v0, vcc, v40, v1
	v_addc_co_u32_e32 v1, vcc, v41, v2, vcc
	.loc	1 92 24                         ; lane-pc-vega20-kernel.cl:92:24
	global_store_dword v[0:1], v3, off offset:8
	.loc	1 93 1 is_stmt 1                ; lane-pc-vega20-kernel.cl:93:1
	buffer_load_dword v42, off, s[0:3], s33 ; 4-byte Folded Reload
	buffer_load_dword v41, off, s[0:3], s33 offset:4 ; 4-byte Folded Reload
	buffer_load_dword v40, off, s[0:3], s33 offset:8 ; 4-byte Folded Reload
	v_readlane_b32 s33, v43, 6
	.cfi_def_cfa_register 64
	s_or_saveexec_b64 s[6:7], -1
	buffer_load_dword v43, off, s[0:3], s32 offset:12 ; 4-byte Folded Reload
	s_mov_b64 exec, s[6:7]
	s_waitcnt vmcnt(0)
	s_setpc_b64 s[4:5]
.Ltmp5:
Lfunc_end4:
	.size	SendResults, Lfunc_end4-SendResults
	.cfi_endproc
                                        ; -- End function
	.section	.AMDGPU.csdata
; Function info:
; codeLenInByte = 432
; NumSgprs: 38
; NumVgprs: 44
; ScratchSize: 48
; MemoryBound: 0
	.text
	.protected	_Z10atomic_addPU3AS3Vii ; -- Begin function _Z10atomic_addPU3AS3Vii
	.weak	_Z10atomic_addPU3AS3Vii
	.p2align	2
	.type	_Z10atomic_addPU3AS3Vii,@function
_Z10atomic_addPU3AS3Vii:                ; @_Z10atomic_addPU3AS3Vii
.Lfunc_begin5:
	.cfi_startproc
; %bb.0:
	.cfi_llvm_def_aspace_cfa 64, 0, 6
	.cfi_escape 0x10, 0x10, 0x08, 0x90, 0x3e, 0x93, 0x04, 0x90, 0x3f, 0x93, 0x04 ;
	.cfi_undefined 2560
	s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
	s_or_saveexec_b64 s[4:5], -1
	buffer_store_dword v2, off, s[0:3], s32 ; 4-byte Folded Spill
	.cfi_offset 2562, 0
	s_mov_b64 exec, s[4:5]
	v_writelane_b32 v2, s30, 0
	v_writelane_b32 v2, s31, 1
	.cfi_escape 0x10, 0x10, 0x0c, 0x90, 0x82, 0x14, 0x9d, 0x20, 0x00, 0x90, 0x82, 0x14, 0x9d, 0x20, 0x20 ;
	v_writelane_b32 v2, exec_lo, 2
	v_writelane_b32 v2, exec_hi, 3
	.cfi_escape 0x10, 0x11, 0x0c, 0x90, 0x82, 0x14, 0x9d, 0x20, 0x40, 0x90, 0x82, 0x14, 0x9d, 0x20, 0x60 ;
	ds_add_rtn_u32 v0, v0, v1
	s_or_saveexec_b64 s[4:5], -1
	buffer_load_dword v2, off, s[0:3], s32  ; 4-byte Folded Reload
	s_mov_b64 exec, s[4:5]
	s_waitcnt vmcnt(0) lgkmcnt(0)
	s_setpc_b64 s[30:31]
.Lfunc_end5:
	.size	_Z10atomic_addPU3AS3Vii, .Lfunc_end5-_Z10atomic_addPU3AS3Vii
	.cfi_endproc
                                        ; -- End function
	.section	.AMDGPU.csdata
; Function info:
; codeLenInByte = 84
; NumSgprs: 37
; NumVgprs: 3
; ScratchSize: 8
; MemoryBound: 0
	.text
	.globl	AddrClassTest                   ; -- Begin function AddrClassTest
	.hidden Lfunc_end6
	.globl	Lfunc_end6
	.hidden VPC_LINE_101
	.globl	VPC_LINE_101
	.hidden VPC_LINE_101_2
	.globl	VPC_LINE_101_2
	.hidden VPC_LINE_109
	.globl	VPC_LINE_109
	.hidden VPC_LINE_109_2
	.globl	VPC_LINE_109_2
	.p2align	8
	.type	AddrClassTest,@function
AddrClassTest:                          ; @AddrClassTest
AddrClassTest$local:
.Lfunc_begin6:
	.loc	1 96 0                          ; lane-pc-vega20-kernel.cl:96:0
	.cfi_startproc
; %bb.0:                                ; %entry
	.cfi_escape 0x0f, 0x03, 0x30, 0x36, 0xe1 ;
	.cfi_undefined 16
	s_mov_b64 s[52:53], s[6:7]
	s_mov_b64 s[54:55], s[4:5]
	s_load_dwordx4 s[4:7], s[8:9], 0x0
	v_mov_b32_e32 v42, v1
	v_mov_b32_e32 v41, v0
	s_add_u32 flat_scratch_lo, s12, s17
	s_addc_u32 flat_scratch_hi, s13, 0
	s_waitcnt lgkmcnt(0)
	v_mov_b32_e32 v0, s4
	v_mov_b32_e32 v1, s5
	s_load_dwordx2 s[4:5], s[54:55], 0x4
	s_add_u32 s0, s0, s17
	s_addc_u32 s1, s1, 0
	v_mov_b32_e32 v43, v2
	s_mov_b64 s[56:57], s[8:9]
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s4, s4, 16
	s_mul_i32 s4, s4, s5
	v_mul_lo_u32 v4, s4, v41
	v_mul_u32_u24_e32 v5, s5, v42
	s_movk_i32 s4, 0x88
	v_mov_b32_e32 v2, s6
	v_add3_u32 v4, v4, v5, v43
	v_lshlrev_b32_e32 v5, 3, v4
	v_lshlrev_b32_e32 v44, 2, v4
	v_mov_b32_e32 v3, s7
	v_add_u32_e32 v4, s4, v5
	ds_write2st64_b64 v4, v[2:3], v[0:1] offset0:2 offset1:6
.Ltmp6:
	.loc	1 97 18 prologue_end            ; lane-pc-vega20-kernel.cl:97:18
	s_add_u32 s8, s56, 16
	s_mov_b32 s37, s16
	v_lshlrev_b32_e32 v0, 20, v43
	v_lshlrev_b32_e32 v1, 10, v42
	v_mov_b32_e32 v40, 0
	s_addc_u32 s9, s57, 0
	v_add_u32_e32 v46, s4, v44
	v_or3_b32 v31, v41, v1, v0
	s_mov_b32 s49, s14
	s_mov_b32 s12, s14
	s_mov_b64 s[4:5], s[54:55]
	s_mov_b64 s[6:7], s[52:53]
	s_mov_b32 s13, s15
	s_mov_b32 s14, s37
	v_mov_b32_e32 v0, v40
	s_getpc_b64 s[16:17]
	s_add_u32 s16, s16, _Z13get_global_idj@rel32@lo+4
	s_addc_u32 s17, s17, _Z13get_global_idj@rel32@hi+12
	s_mov_b32 s32, 0
	s_mov_b32 s48, s15
	s_mov_b64 s[50:51], s[10:11]
	v_add_u32_e32 v47, 0xc88, v5
	v_add_u32_e32 v45, 0x488, v5
	s_swappc_b64 s[30:31], s[16:17]
	s_mov_b32 s4, 0xaaaaaaab
	v_mul_lo_u32 v1, v0, s4
	s_mov_b32 s4, 0x55555556
	v_mov_b32_e32 v2, 50
	.loc	1 97 12                         ; lane-pc-vega20-kernel.cl:97:12
	ds_write_b32 v44, v0 offset:136
	.loc	1 101 11                         ; lane-pc-vega20-kernel.cl:101:11
	v_cmp_gt_u32_e32 vcc, s4, v1
	.loc	1 99 26                         ; lane-pc-vega20-kernel.cl:99:26
	ds_write_b8 v40, v2 offset:4
	s_and_saveexec_b64 s[4:5], vcc
	s_xor_b64 s[4:5], exec, s[4:5]
	s_cbranch_execz BB6_2
; %bb.1:                                ; %if.else
	.loc	1 107 29                         ; lane-pc-vega20-kernel.cl:107:29
	ds_read_b64 v[2:3], v47
	v_mov_b32_e32 v1, 0
	v_lshlrev_b64 v[4:5], 2, v[0:1]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32_e32 v2, vcc, v2, v4
	v_addc_co_u32_e32 v3, vcc, v3, v5, vcc
	global_load_dword v0, v[2:3], off
	.loc	1 107 27 is_stmt 0               ; lane-pc-vega20-kernel.cl:107:27
	s_waitcnt vmcnt(0)
	ds_write_b32 v1, v0
BB6_2:                                  ; %Flow9
	.loc	1 0 27                          ; lane-pc-vega20-kernel.cl:0:27
VPC_LINE_101:
	s_or_saveexec_b64 s[38:39], s[4:5]
	s_xor_b64 exec, exec, s[38:39]
	s_cbranch_execz BB6_4
; %bb.3:                                ; %if.then
	.loc	1 103 40 is_stmt 1               ; lane-pc-vega20-kernel.cl:103:40
	s_add_u32 s34, s56, 16
	v_lshlrev_b32_e32 v0, 20, v43
	v_lshlrev_b32_e32 v1, 10, v42
	s_addc_u32 s35, s57, 0
	v_or3_b32 v44, v41, v1, v0
	s_mov_b64 s[4:5], s[54:55]
	s_mov_b64 s[6:7], s[52:53]
	s_mov_b64 s[8:9], s[34:35]
	s_mov_b64 s[10:11], s[50:51]
	s_mov_b32 s12, s49
	s_mov_b32 s13, s48
	s_mov_b32 s14, s37
	v_mov_b32_e32 v31, v44
	s_getpc_b64 s[16:17]
	s_add_u32 s16, s16, GenValue@rel32@lo+4
	s_addc_u32 s17, s17, GenValue@rel32@hi+12
	s_swappc_b64 s[30:31], s[16:17]
	v_mov_b32_e32 v40, 0
	v_mov_b32_e32 v1, v0
	.loc	1 103 5 is_stmt 0                ; lane-pc-vega20-kernel.cl:103:5
	s_mov_b64 s[4:5], s[54:55]
	s_mov_b64 s[6:7], s[52:53]
	s_mov_b64 s[8:9], s[34:35]
	s_mov_b64 s[10:11], s[50:51]
	s_mov_b32 s12, s49
	s_mov_b32 s13, s48
	s_mov_b32 s14, s37
	v_mov_b32_e32 v31, v44
	v_mov_b32_e32 v0, v40
	s_getpc_b64 s[16:17]
	s_add_u32 s16, s16, _Z10atomic_addPU3AS3Vii@rel32@lo+4
	s_addc_u32 s17, s17, _Z10atomic_addPU3AS3Vii@rel32@hi+12
	s_swappc_b64 s[30:31], s[16:17]
	.loc	1 104 39 is_stmt 1               ; lane-pc-vega20-kernel.cl:104:39
	ds_read_b32 v0, v46
	.loc	1 104 36 is_stmt 0               ; lane-pc-vega20-kernel.cl:104:36
	ds_read_b64 v[2:3], v47
	.loc	1 104 39                         ; lane-pc-vega20-kernel.cl:104:39
	v_mov_b32_e32 v1, 0
	.loc	1 104 36                         ; lane-pc-vega20-kernel.cl:104:36
	s_waitcnt lgkmcnt(1)
	v_lshlrev_b64 v[0:1], 2, v[0:1]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32_e32 v0, vcc, v2, v0
	v_addc_co_u32_e32 v1, vcc, v3, v1, vcc
	global_load_dword v0, v[0:1], off
	.loc	1 104 28                         ; lane-pc-vega20-kernel.cl:104:28
	s_waitcnt vmcnt(0)
	ds_write_b8 v40, v0 offset:4
VPC_LINE_101_2:
BB6_4:                                  ; %if.end
	.loc	1 0 28                          ; lane-pc-vega20-kernel.cl:0:28
	s_or_b64 exec, exec, s[38:39]
	.loc	1 109 7 is_stmt 1                ; lane-pc-vega20-kernel.cl:109:7
	ds_read_b32 v0, v46
	.loc	1 109 11 is_stmt 0               ; lane-pc-vega20-kernel.cl:109:11
	s_waitcnt lgkmcnt(0)
	v_and_b32_e32 v1, 7, v0
	v_cmp_eq_u32_e32 vcc, 0, v1
	s_and_saveexec_b64 s[4:5], vcc
	s_xor_b64 s[4:5], exec, s[4:5]
	s_cbranch_execz BB6_6
; %bb.5:                                ; %if.else11
	.loc	1 112 36 is_stmt 1               ; lane-pc-vega20-kernel.cl:112:36
	v_mov_b32_e32 v1, 0
	v_lshlrev_b64 v[1:2], 2, v[0:1]
	s_getpc_b64 s[6:7]
	s_add_u32 s6, s6, const_array@rel32@lo+4
	s_addc_u32 s7, s7, const_array@rel32@hi+12
	v_mov_b32_e32 v3, s7
	v_add_co_u32_e32 v1, vcc, s6, v1
	v_addc_co_u32_e32 v2, vcc, v2, v3, vcc
	global_load_dword v1, v[1:2], off
	.loc	1 112 5 is_stmt 0                ; lane-pc-vega20-kernel.cl:112:5
	v_lshlrev_b32_e32 v2, 2, v0
	.loc	1 112 34                         ; lane-pc-vega20-kernel.cl:112:34
	s_waitcnt vmcnt(0)
	ds_write_b32 v2, v1 offset:8
BB6_6:                                  ; %Flow
	.loc	1 0 34                          ; lane-pc-vega20-kernel.cl:0:34
	s_or_saveexec_b64 s[58:59], s[4:5]
	s_xor_b64 exec, exec, s[58:59]
VPC_LINE_109:
	s_cbranch_execz BB6_8
; %bb.7:                                ; %if.then8
	.loc	1 110 5 is_stmt 1                ; lane-pc-vega20-kernel.cl:110:5
	s_add_u32 s8, s56, 16
	v_lshlrev_b32_e32 v1, 20, v43
	v_lshlrev_b32_e32 v2, 10, v42
	s_addc_u32 s9, s57, 0
	.loc	1 110 26 is_stmt 0               ; lane-pc-vega20-kernel.cl:110:26
	v_lshl_add_u32 v0, v0, 2, 8
	v_or3_b32 v31, v41, v2, v1
	.loc	1 110 5                          ; lane-pc-vega20-kernel.cl:110:5
	s_mov_b64 s[4:5], s[54:55]
	s_mov_b64 s[6:7], s[52:53]
	s_mov_b64 s[10:11], s[50:51]
	s_mov_b32 s12, s49
	s_mov_b32 s13, s48
	s_mov_b32 s14, s37
	s_getpc_b64 s[16:17]
	s_add_u32 s16, s16, ChangeLocalContent@rel32@lo+4
	s_addc_u32 s17, s17, ChangeLocalContent@rel32@hi+12
	s_swappc_b64 s[30:31], s[16:17]
VPC_LINE_109_2:
BB6_8:                                  ; %if.end16
	.loc	1 0 5                           ; lane-pc-vega20-kernel.cl:0:5
	s_or_b64 exec, exec, s[58:59]
	.loc	1 114 31 is_stmt 1               ; lane-pc-vega20-kernel.cl:114:31
	ds_read_b64 v[1:2], v45
	.loc	1 114 3 is_stmt 0                ; lane-pc-vega20-kernel.cl:114:3
	s_add_u32 s8, s56, 16
	v_lshlrev_b32_e32 v0, 20, v43
	v_lshlrev_b32_e32 v3, 10, v42
	s_addc_u32 s9, s57, 0
	v_or3_b32 v31, v41, v3, v0
	s_mov_b64 s[4:5], s[54:55]
	s_mov_b64 s[6:7], s[52:53]
	s_mov_b64 s[10:11], s[50:51]
	s_mov_b32 s12, s49
	s_mov_b32 s13, s48
	s_mov_b32 s14, s37
	v_mov_b32_e32 v0, 0
	s_getpc_b64 s[16:17]
	s_add_u32 s16, s16, SendResults@rel32@lo+4
	s_addc_u32 s17, s17, SendResults@rel32@hi+12
	s_swappc_b64 s[30:31], s[16:17]
	.loc	1 115 1 is_stmt 1                ; lane-pc-vega20-kernel.cl:115:1
	s_endpgm
.Ltmp7:
	.section	.rodata,#alloc
	.p2align	6
	.amdhsa_kernel AddrClassTest
		.amdhsa_group_segment_fixed_size 5256
		.amdhsa_private_segment_fixed_size 320
		.amdhsa_kernarg_size 72
		.amdhsa_user_sgpr_private_segment_buffer 1
		.amdhsa_user_sgpr_dispatch_ptr 1
		.amdhsa_user_sgpr_queue_ptr 1
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 1
		.amdhsa_user_sgpr_flat_scratch_init 1
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_system_sgpr_private_segment_wavefront_offset 1
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 1
		.amdhsa_system_sgpr_workgroup_id_z 1
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 2
		.amdhsa_next_free_vgpr 48
		.amdhsa_next_free_sgpr 60
		.amdhsa_reserve_xnack_mask 1
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
Lfunc_end6:
	.size	AddrClassTest, Lfunc_end6-AddrClassTest
	.cfi_endproc
                                        ; -- End function
	.section	.AMDGPU.csdata
; Kernel info:
; codeLenInByte = 872
; NumSgprs: 66
; NumVgprs: 48
; ScratchSize: 320
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 5256 bytes/workgroup (compile time only)
; SGPRBlocks: 8
; VGPRBlocks: 11
; NumSGPRsForWavesPerEU: 66
; NumVGPRsForWavesPerEU: 48
; Occupancy: 5
; WaveLimiterHint : 1
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 1
; COMPUTE_PGM_RSRC2:USER_SGPR: 14
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 1
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 2
	.type	const_struct,@object            ; @const_struct
	.section	.rodata,#alloc
	.hidden const_struct
	.globl	const_struct
	.p2align	2
const_struct:
const_struct$local:
	.long	32                              ; 0x20
	.byte	50                              ; 0x32
	.zero	3
	.long	32                              ; 0x20
	.long	31                              ; 0x1f
	.long	30                              ; 0x1e
	.long	29                              ; 0x1d
	.long	28                              ; 0x1c
	.long	27                              ; 0x1b
	.long	26                              ; 0x1a
	.long	25                              ; 0x19
	.long	24                              ; 0x18
	.long	23                              ; 0x17
	.long	22                              ; 0x16
	.long	21                              ; 0x15
	.long	20                              ; 0x14
	.long	19                              ; 0x13
	.long	18                              ; 0x12
	.long	17                              ; 0x11
	.long	16                              ; 0x10
	.long	15                              ; 0xf
	.long	14                              ; 0xe
	.long	13                              ; 0xd
	.long	12                              ; 0xc
	.long	11                              ; 0xb
	.long	10                              ; 0xa
	.long	9                               ; 0x9
	.long	8                               ; 0x8
	.long	7                               ; 0x7
	.long	6                               ; 0x6
	.long	5                               ; 0x5
	.long	4                               ; 0x4
	.long	3                               ; 0x3
	.long	2                               ; 0x2
	.long	1                               ; 0x1
	.size	const_struct, 136

	.type	const_array,@object             ; @const_array
	.hidden const_array
	.globl	const_array
	.p2align	2
const_array:
const_array$local:
	.long	1                               ; 0x1
	.long	1                               ; 0x1
	.long	1                               ; 0x1
	.long	1                               ; 0x1
	.long	1                               ; 0x1
	.long	5                               ; 0x5
	.long	5                               ; 0x5
	.long	7                               ; 0x7
	.long	2                               ; 0x2
	.long	2                               ; 0x2
	.long	2                               ; 0x2
	.long	2                               ; 0x2
	.long	2                               ; 0x2
	.long	5                               ; 0x5
	.long	5                               ; 0x5
	.long	10                              ; 0xa
	.long	3                               ; 0x3
	.long	3                               ; 0x3
	.long	3                               ; 0x3
	.long	3                               ; 0x3
	.long	3                               ; 0x3
	.long	5                               ; 0x5
	.long	5                               ; 0x5
	.long	2                               ; 0x2
	.long	4                               ; 0x4
	.long	4                               ; 0x4
	.long	4                               ; 0x4
	.long	4                               ; 0x4
	.long	4                               ; 0x4
	.long	5                               ; 0x5
	.long	5                               ; 0x5
	.long	3                               ; 0x3
	.size	const_array, 128

	.type	__const.GenValue.array,@object  ; @__const.GenValue.array
	.p2align	2
__const.GenValue.array:
	.long	1                               ; 0x1
	.long	1                               ; 0x1
	.long	1                               ; 0x1
	.long	1                               ; 0x1
	.long	2                               ; 0x2
	.long	2                               ; 0x2
	.long	2                               ; 0x2
	.long	2                               ; 0x2
	.long	3                               ; 0x3
	.long	3                               ; 0x3
	.long	3                               ; 0x3
	.long	3                               ; 0x3
	.long	4                               ; 0x4
	.long	4                               ; 0x4
	.long	4                               ; 0x4
	.long	4                               ; 0x4
	.long	5                               ; 0x5
	.long	5                               ; 0x5
	.long	5                               ; 0x5
	.long	5                               ; 0x5
	.long	6                               ; 0x6
	.long	6                               ; 0x6
	.long	6                               ; 0x6
	.long	6                               ; 0x6
	.long	7                               ; 0x7
	.long	7                               ; 0x7
	.long	7                               ; 0x7
	.long	7                               ; 0x7
	.long	8                               ; 0x8
	.long	8                               ; 0x8
	.long	8                               ; 0x8
	.long	8                               ; 0x8
	.size	__const.GenValue.array, 128

	.type	__const.ChangeLocalContent.array,@object ; @__const.ChangeLocalContent.array
	.p2align	2
__const.ChangeLocalContent.array:
	.long	1                               ; 0x1
	.long	1                               ; 0x1
	.long	1                               ; 0x1
	.long	1                               ; 0x1
	.long	2                               ; 0x2
	.long	2                               ; 0x2
	.long	2                               ; 0x2
	.long	2                               ; 0x2
	.long	3                               ; 0x3
	.long	3                               ; 0x3
	.long	3                               ; 0x3
	.long	3                               ; 0x3
	.long	4                               ; 0x4
	.long	4                               ; 0x4
	.long	4                               ; 0x4
	.long	4                               ; 0x4
	.long	5                               ; 0x5
	.long	5                               ; 0x5
	.long	5                               ; 0x5
	.long	5                               ; 0x5
	.long	6                               ; 0x6
	.long	6                               ; 0x6
	.long	6                               ; 0x6
	.long	6                               ; 0x6
	.long	7                               ; 0x7
	.long	7                               ; 0x7
	.long	7                               ; 0x7
	.long	7                               ; 0x7
	.long	8                               ; 0x8
	.long	8                               ; 0x8
	.long	8                               ; 0x8
	.long	8                               ; 0x8
	.size	__const.ChangeLocalContent.array, 128

	.section	.debug_abbrev
	.byte	1                               ; Abbreviation Code
	.byte	17                              ; DW_TAG_compile_unit
	.byte	0                               ; DW_CHILDREN_no
	.byte	37                              ; DW_AT_producer
	.byte	14                              ; DW_FORM_strp
	.byte	19                              ; DW_AT_language
	.byte	5                               ; DW_FORM_data2
	.byte	3                               ; DW_AT_name
	.byte	14                              ; DW_FORM_strp
	.byte	16                              ; DW_AT_stmt_list
	.byte	23                              ; DW_FORM_sec_offset
	.byte	27                              ; DW_AT_comp_dir
	.byte	14                              ; DW_FORM_strp
	.ascii	"\211|"                         ; DW_AT_LLVM_augmentation
	.byte	14                              ; DW_FORM_strp
	.byte	17                              ; DW_AT_low_pc
	.byte	1                               ; DW_FORM_addr
	.byte	85                              ; DW_AT_ranges
	.byte	23                              ; DW_FORM_sec_offset
	.byte	0                               ; EOM(1)
	.byte	0                               ; EOM(2)
	.byte	0                               ; EOM(3)
	.section	.debug_ranges
.Ldebug_ranges0:
	.quad	.Lfunc_begin2
	.quad	Lfunc_end4
	.quad	.Lfunc_begin6
	.quad	Lfunc_end6
	.quad	0
	.quad	0
	.section	.debug_str,"MS",@progbits,1
.Linfo_string0:
	.asciz	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)" ; string offset=0
.Linfo_string1:
	.asciz	"<stdin>"                       ; string offset=130
.Linfo_string2:
	.asciz	"." ; string offset=138
.Linfo_string3:
	.asciz	"[llvm:v0.0]"                   ; string offset=183
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.ident	"clang version 13.0.0 (ssh://gerritgit/lightning/ec/llvm-project amd-mainline-open 21195 1e0323dcced173003e83b4c3d12047c6cb2b0003)"
	.section	".note.GNU-stack"
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .is_const:       true
        .name:           in
        .offset:         0
        .size:           8
        .type_name:      'int*'
        .value_kind:     global_buffer
      - .address_space:  global
        .name:           out
        .offset:         8
        .size:           8
        .type_name:      'struct test_struct*'
        .value_kind:     global_buffer
      - .offset:         16
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         24
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         32
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .address_space:  global
        .offset:         40
        .size:           8
        .value_kind:     hidden_none
      - .address_space:  global
        .offset:         48
        .size:           8
        .value_kind:     hidden_none
      - .address_space:  global
        .offset:         56
        .size:           8
        .value_kind:     hidden_none
      - .address_space:  global
        .offset:         64
        .size:           8
        .value_kind:     hidden_multigrid_sync_arg
    .group_segment_fixed_size: 5256
    .kernarg_segment_align: 8
    .kernarg_segment_size: 72
    .language:       OpenCL C
    .language_version:
      - 1
      - 0
    .max_flat_workgroup_size: 256
    .name:           AddrClassTest
    .private_segment_fixed_size: 320
    .sgpr_count:     66
    .sgpr_spill_count: 0
    .symbol:         AddrClassTest.kd
    .vgpr_count:     48
    .vgpr_spill_count: 0
    .wavefront_size: 64
amdhsa.target:   amdgcn-amd-amdhsa--gfx906
amdhsa.version:
  - 1
  - 1
...

	.end_amdgpu_metadata
	.section	.debug_line
.Lline_table_start0: