/* 
This is an good trick to see what is going on for debugging purposes. It is terribly BAD
to print from GPU kernels in anything you want to be performance oriented though. It is not a
performance oriented feature!
*/
#include <stdio.h>
#include <stdlib.h>
#include <cuda.h>
#include <cuda_runtime.h>
 
__global__ void kernel()
{
	/* 
           this just gets some kernel specific parameters
	   this is just so you can see how non-deterministic thread timing is
	*/
        int tidx = threadIdx.x + blockIdx.x * blockDim.x;
	int tidy = threadIdx.y + blockIdx.y * blockDim.y;
 
	/* print some stuff out */
	int size = sizeof(int);
	printf("Hello, World! size=%d   tidx=%d, tidy=%d\n", size, tidx, tidy);
	return;
}
 
int main(int argc, char** argv)
{
        /*
           Keep this in mind. in Cuda 8 compute 2.0 was deprecated and it may be 
           removed by now. CDER only currently (11/2018) supports Cuda 7 so the below
           will work and may or may not warn you about this.
        */
	printf("You compile this with 'nvcc -arch sm_20 hello.cu -o hello'\n");
	printf("You need -at least- arch of sm_20 to print from kernels\n");
	dim3 dimBlock( 16, 16, 1 );
	dim3 dimGrid( 16, 16, 1 );
 
	kernel<<<dimGrid,dimBlock>>>();
	cudaDeviceSynchronize();  /* you also -need- this here to print from the kernel */
	return 1;
}

	.file	"tmpxft_002a052f_00000000-6_GPUPrint.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2030:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2030:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z24__device_stub__Z6kernelvv
	.type	_Z24__device_stub__Z6kernelvv, @function
_Z24__device_stub__Z6kernelvv:
.LFB2052:
	.cfi_startproc
	endbr64
	subq	$88, %rsp
	.cfi_def_cfa_offset 96
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	movabsq	$4294967297, %rax
	leaq	8(%rsp), %rcx
	movq	%rsp, %rdx
	leaq	28(%rsp), %rsi
	leaq	16(%rsp), %rdi
	movl	$1, 24(%rsp)
	movl	$1, 36(%rsp)
	movq	%rax, 16(%rsp)
	movq	%rax, 28(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	8(%rsp)
	.cfi_def_cfa_offset 104
	leaq	_Z6kernelv(%rip), %rdi
	pushq	8(%rsp)
	.cfi_def_cfa_offset 112
	movq	44(%rsp), %rcx
	movl	52(%rsp), %r8d
	movq	32(%rsp), %rsi
	movl	40(%rsp), %edx
	leaq	80(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 104
	popq	%rdx
	.cfi_def_cfa_offset 96
.L2:
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$88, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2052:
	.size	_Z24__device_stub__Z6kernelvv, .-_Z24__device_stub__Z6kernelvv
	.globl	_Z6kernelv
	.type	_Z6kernelv, @function
_Z6kernelv:
.LFB2053:
	.cfi_startproc
	endbr64
	jmp	_Z24__device_stub__Z6kernelvv
	.cfi_endproc
.LFE2053:
	.size	_Z6kernelv, .-_Z6kernelv
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"You compile this with 'nvcc -arch sm_20 hello.cu -o hello'\n"
.LC1:
	.string	"You need -at least- arch of sm_20 to print from kernels\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2027:
	.cfi_startproc
	endbr64
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	leaq	.LC0(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC1(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	movabsq	$68719476752, %rdi
	movl	$1, %esi
	movq	%rdi, %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L9
	call	_Z24__device_stub__Z6kernelvv
.L9:
	call	cudaDeviceSynchronize@PLT
	movl	$1, %eax
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2027:
	.size	main, .-main
	.section	.rodata.str1.1
.LC2:
	.string	"_Z6kernelv"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2055:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC2(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z6kernelv(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2055:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:

/* 
This is an good trick to see what is going on for debugging purposes. It is terribly BAD
to print from GPU kernels in anything you want to be performance oriented though. It is not a
performance oriented feature!
*/
#include <stdio.h>
#include <stdlib.h>
#include <hip/hip_runtime.h>

 
__global__ void kernel()
{
	/* 
           this just gets some kernel specific parameters
	   this is just so you can see how non-deterministic thread timing is
	*/
        int tidx = threadIdx.x + blockIdx.x * blockDim.x;
	int tidy = threadIdx.y + blockIdx.y * blockDim.y;
 
	/* print some stuff out */
	int size = sizeof(int);
	printf("Hello, World! size=%d   tidx=%d, tidy=%d\n", size, tidx, tidy);
	return;
}
 
int main(int argc, char** argv)
{
        /*
           Keep this in mind. in Cuda 8 compute 2.0 was deprecated and it may be 
           removed by now. CDER only currently (11/2018) supports Cuda 7 so the below
           will work and may or may not warn you about this.
        */
	printf("You compile this with 'nvcc -arch sm_20 hello.cu -o hello'\n");
	printf("You need -at least- arch of sm_20 to print from kernels\n");
	dim3 dimBlock( 16, 16, 1 );
	dim3 dimGrid( 16, 16, 1 );
 
	kernel<<<dimGrid,dimBlock>>>();
	hipDeviceSynchronize();  /* you also -need- this here to print from the kernel */
	return 1;
}

	.text
	.file	"GPUPrint.hip"
	.globl	_Z21__device_stub__kernelv      # -- Begin function _Z21__device_stub__kernelv
	.type	_Z21__device_stub__kernelv,@function
_Z21__device_stub__kernelv:             # @_Z21__device_stub__kernelv
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r12
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	subq	$56, %rsp
	.cfi_def_cfa_offset 96
	.cfi_offset %rbx, -40
	.cfi_offset %r12, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rbx
	leaq	24(%rsp), %r14
	leaq	16(%rsp), %r15
	leaq	8(%rsp), %r12
	movq	%rbx, %rdi
	movq	%r14, %rsi
	movq	%r15, %rdx
	movq	%r12, %rcx
	callq	__hipPopCallConfiguration
	movq	(%rbx), %rsi
	movl	8(%rbx), %edx
	movq	(%r14), %rcx
	movl	8(%r14), %r8d
	movq	%rsp, %r9
	movl	$_Z6kernelv, %edi
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	pushq	(%r15)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$72, %rsp
	.cfi_adjust_cfa_offset -72
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z21__device_stub__kernelv, .Lfunc_end0-_Z21__device_stub__kernelv
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rax
	.cfi_def_cfa_offset 16
	movl	$.Lstr, %edi
	callq	puts@PLT
	movl	$.Lstr.1, %edi
	callq	puts@PLT
	movabsq	$68719476752, %rdi              # imm = 0x1000000010
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB1_2
# %bb.1:
	callq	_Z21__device_stub__kernelv
.LBB1_2:
	callq	hipDeviceSynchronize
	movl	$1, %eax
	popq	%rcx
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	main, .Lfunc_end1-main
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	jne	.LBB2_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB2_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z6kernelv, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end2:
	.size	__hip_module_ctor, .Lfunc_end2-__hip_module_ctor
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_dtor
	.type	__hip_module_dtor,@function
__hip_module_dtor:                      # @__hip_module_dtor
	.cfi_startproc
# %bb.0:
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB3_2
# %bb.1:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__hipUnregisterFatBinary
	movq	$0, __hip_gpubin_handle(%rip)
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB3_2:
	retq
.Lfunc_end3:
	.size	__hip_module_dtor, .Lfunc_end3-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z6kernelv,@object              # @_Z6kernelv
	.section	.rodata,"a",@progbits
	.globl	_Z6kernelv
	.p2align	3, 0x0
_Z6kernelv:
	.quad	_Z21__device_stub__kernelv
	.size	_Z6kernelv, 8

	.type	.L__unnamed_1,@object           # @0
	.section	.rodata.str1.1,"aMS",@progbits,1
.L__unnamed_1:
	.asciz	"_Z6kernelv"
	.size	.L__unnamed_1, 11

	.type	__hip_fatbin_wrapper,@object    # @__hip_fatbin_wrapper
	.section	.hipFatBinSegment,"a",@progbits
	.p2align	3, 0x0
__hip_fatbin_wrapper:
	.long	1212764230                      # 0x48495046
	.long	1                               # 0x1
	.quad	__hip_fatbin
	.quad	0
	.size	__hip_fatbin_wrapper, 24

	.type	__hip_gpubin_handle,@object     # @__hip_gpubin_handle
	.local	__hip_gpubin_handle
	.comm	__hip_gpubin_handle,8,8
	.section	.init_array,"aw",@init_array
	.p2align	3, 0x0
	.quad	__hip_module_ctor
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.type	.Lstr,@object                   # @str
	.section	.rodata.str1.1,"aMS",@progbits,1
.Lstr:
	.asciz	"You compile this with 'nvcc -arch sm_20 hello.cu -o hello'"
	.size	.Lstr, 59

	.type	.Lstr.1,@object                 # @str.1
.Lstr.1:
	.asciz	"You need -at least- arch of sm_20 to print from kernels"
	.size	.Lstr.1, 56

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym _Z21__device_stub__kernelv
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z6kernelv
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z6kernelv              ; -- Begin function _Z6kernelv
	.globl	_Z6kernelv
	.p2align	8
	.type	_Z6kernelv,@function
_Z6kernelv:                             ; @_Z6kernelv
; %bb.0:
	s_clause 0x1
	s_load_b32 s16, s[0:1], 0xc
	s_load_b64 s[2:3], s[0:1], 0x50
	v_mbcnt_lo_u32_b32 v29, -1, 0
	v_mov_b32_e32 v7, 0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_dual_mov_b32 v8, 0 :: v_dual_mov_b32 v5, v29
	;;#ASMSTART
	;;#ASMEND
	v_readfirstlane_b32 s0, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_eq_u32_e64 s0, s0, v5
	s_and_saveexec_b32 s1, s0
	s_cbranch_execz .LBB0_6
; %bb.1:
	v_mov_b32_e32 v1, 0
	s_mov_b32 s4, exec_lo
	s_waitcnt lgkmcnt(0)
	global_load_b64 v[9:10], v1, s[2:3] offset:24 glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	s_clause 0x1
	global_load_b64 v[2:3], v1, s[2:3] offset:40
	global_load_b64 v[6:7], v1, s[2:3]
	s_waitcnt vmcnt(1)
	v_and_b32_e32 v2, v2, v9
	v_and_b32_e32 v3, v3, v10
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_hi_u32 v4, v2, 24
	v_mul_lo_u32 v3, v3, 24
	v_mul_lo_u32 v2, v2, 24
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v3, v4, v3
	s_waitcnt vmcnt(0)
	v_add_co_u32 v2, vcc_lo, v6, v2
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v3, vcc_lo, v7, v3, vcc_lo
	global_load_b64 v[7:8], v[2:3], off glc
	s_waitcnt vmcnt(0)
	global_atomic_cmpswap_b64 v[7:8], v1, v[7:10], s[2:3] offset:24 glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	v_cmpx_ne_u64_e64 v[7:8], v[9:10]
	s_cbranch_execz .LBB0_5
; %bb.2:                                ; %.preheader3.i.i.i.preheader
	s_mov_b32 s5, 0
.LBB0_3:                                ; %.preheader3.i.i.i
                                        ; =>This Inner Loop Header: Depth=1
	s_sleep 1
	s_clause 0x1
	global_load_b64 v[2:3], v1, s[2:3] offset:40
	global_load_b64 v[11:12], v1, s[2:3]
	v_dual_mov_b32 v10, v8 :: v_dual_mov_b32 v9, v7
	s_waitcnt vmcnt(1)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_and_b32_e32 v2, v2, v9
	v_and_b32_e32 v8, v3, v10
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[6:7], null, v2, 24, v[11:12]
	v_mov_b32_e32 v2, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[3:4], null, v8, 24, v[2:3]
	v_mov_b32_e32 v7, v3
	global_load_b64 v[7:8], v[6:7], off glc
	s_waitcnt vmcnt(0)
	global_atomic_cmpswap_b64 v[7:8], v1, v[7:10], s[2:3] offset:24 glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	v_cmp_eq_u64_e32 vcc_lo, v[7:8], v[9:10]
	s_or_b32 s5, vcc_lo, s5
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s5
	s_cbranch_execnz .LBB0_3
; %bb.4:                                ; %Flow418
	s_or_b32 exec_lo, exec_lo, s5
.LBB0_5:                                ; %Flow420
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s4
.LBB0_6:                                ; %.loopexit4.i.i.i
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s1
	v_mov_b32_e32 v6, 0
	v_readfirstlane_b32 s4, v7
	v_readfirstlane_b32 s5, v8
	s_mov_b32 s10, exec_lo
	s_waitcnt lgkmcnt(0)
	s_clause 0x1
	global_load_b64 v[9:10], v6, s[2:3] offset:40
	global_load_b128 v[1:4], v6, s[2:3]
	s_waitcnt vmcnt(1)
	v_readfirstlane_b32 s6, v9
	v_readfirstlane_b32 s7, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b64 s[6:7], s[4:5], s[6:7]
	s_mul_i32 s1, s7, 24
	s_mul_hi_u32 s8, s6, 24
	s_mul_i32 s9, s6, 24
	s_and_saveexec_b32 s11, s0
	s_cbranch_execz .LBB0_8
; %bb.7:
	v_dual_mov_b32 v7, s10 :: v_dual_mov_b32 v8, v6
	s_add_i32 s10, s8, s1
	s_waitcnt vmcnt(0)
	v_add_co_u32 v11, vcc_lo, v1, s9
	v_add_co_ci_u32_e32 v12, vcc_lo, s10, v2, vcc_lo
	v_dual_mov_b32 v9, 2 :: v_dual_mov_b32 v10, 1
	global_store_b128 v[11:12], v[7:10], off offset:8
.LBB0_8:
	s_or_b32 exec_lo, exec_lo, s11
	s_lshl_b64 s[6:7], s[6:7], 12
	v_lshlrev_b64 v[7:8], 6, v[5:6]
	s_waitcnt vmcnt(0)
	v_add_co_u32 v3, vcc_lo, v3, s6
	v_add_co_ci_u32_e32 v4, vcc_lo, s7, v4, vcc_lo
	s_mov_b32 s20, 0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_u32 v9, vcc_lo, v3, v7
	s_mov_b32 s21, s20
	s_mov_b32 s22, s20
	s_mov_b32 s23, s20
	v_add_co_ci_u32_e32 v10, vcc_lo, v4, v8, vcc_lo
	v_dual_mov_b32 v5, 33 :: v_dual_mov_b32 v8, v6
	v_mov_b32_e32 v7, v6
	v_dual_mov_b32 v11, s20 :: v_dual_mov_b32 v14, s23
	v_dual_mov_b32 v12, s21 :: v_dual_mov_b32 v13, s22
	s_clause 0x3
	global_store_b128 v[9:10], v[5:8], off
	global_store_b128 v[9:10], v[11:14], off offset:16
	global_store_b128 v[9:10], v[11:14], off offset:32
	global_store_b128 v[9:10], v[11:14], off offset:48
	s_and_saveexec_b32 s6, s0
	s_cbranch_execz .LBB0_16
; %bb.9:
	v_mov_b32_e32 v11, 0
	s_mov_b32 s7, exec_lo
	s_clause 0x1
	global_load_b64 v[14:15], v11, s[2:3] offset:32 glc
	global_load_b64 v[3:4], v11, s[2:3] offset:40
	v_dual_mov_b32 v12, s4 :: v_dual_mov_b32 v13, s5
	s_waitcnt vmcnt(0)
	v_and_b32_e32 v4, s5, v4
	v_and_b32_e32 v3, s4, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_lo_u32 v4, v4, 24
	v_mul_hi_u32 v5, v3, 24
	v_mul_lo_u32 v3, v3, 24
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v4, v5, v4
	v_add_co_u32 v7, vcc_lo, v1, v3
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v8, vcc_lo, v2, v4, vcc_lo
	global_store_b64 v[7:8], v[14:15], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[5:6], v11, v[12:15], s[2:3] offset:32 glc
	s_waitcnt vmcnt(0)
	v_cmpx_ne_u64_e64 v[5:6], v[14:15]
	s_cbranch_execz .LBB0_12
; %bb.10:                               ; %.preheader1.i.i.i.preheader
	s_mov_b32 s10, 0
.LBB0_11:                               ; %.preheader1.i.i.i
                                        ; =>This Inner Loop Header: Depth=1
	v_dual_mov_b32 v3, s4 :: v_dual_mov_b32 v4, s5
	s_sleep 1
	global_store_b64 v[7:8], v[5:6], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[3:4], v11, v[3:6], s[2:3] offset:32 glc
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, v[3:4], v[5:6]
	v_dual_mov_b32 v6, v4 :: v_dual_mov_b32 v5, v3
	s_or_b32 s10, vcc_lo, s10
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s10
	s_cbranch_execnz .LBB0_11
.LBB0_12:                               ; %Flow416
	s_or_b32 exec_lo, exec_lo, s7
	v_mov_b32_e32 v6, 0
	s_mov_b32 s10, exec_lo
	s_mov_b32 s7, exec_lo
	v_mbcnt_lo_u32_b32 v5, s10, 0
	global_load_b64 v[3:4], v6, s[2:3] offset:16
	v_cmpx_eq_u32_e32 0, v5
	s_cbranch_execz .LBB0_14
; %bb.13:
	s_bcnt1_i32_b32 s10, s10
	s_delay_alu instid0(SALU_CYCLE_1)
	v_mov_b32_e32 v5, s10
	s_waitcnt vmcnt(0)
	global_atomic_add_u64 v[3:4], v[5:6], off offset:8
.LBB0_14:
	s_or_b32 exec_lo, exec_lo, s7
	s_waitcnt vmcnt(0)
	global_load_b64 v[5:6], v[3:4], off offset:16
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, 0, v[5:6]
	s_cbranch_vccnz .LBB0_16
; %bb.15:
	global_load_b32 v3, v[3:4], off offset:24
	v_mov_b32_e32 v4, 0
	s_waitcnt vmcnt(0)
	v_readfirstlane_b32 s7, v3
	s_waitcnt_vscnt null, 0x0
	global_store_b64 v[5:6], v[3:4], off
	s_and_b32 m0, s7, 0xff
	s_sendmsg sendmsg(MSG_INTERRUPT)
.LBB0_16:                               ; %Flow417
	s_or_b32 exec_lo, exec_lo, s6
	s_add_i32 s8, s8, s1
	v_add_co_u32 v1, vcc_lo, v1, s9
	v_add_co_ci_u32_e32 v2, vcc_lo, s8, v2, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v1, vcc_lo, v1, 20
	v_add_co_ci_u32_e32 v2, vcc_lo, 0, v2, vcc_lo
.LBB0_17:                               ; =>This Inner Loop Header: Depth=1
	v_mov_b32_e32 v3, 1
	s_and_saveexec_b32 s1, s0
	s_cbranch_execz .LBB0_19
; %bb.18:                               ;   in Loop: Header=BB0_17 Depth=1
	global_load_b32 v3, v[1:2], off glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	v_and_b32_e32 v3, 1, v3
.LBB0_19:                               ;   in Loop: Header=BB0_17 Depth=1
	s_or_b32 exec_lo, exec_lo, s1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_readfirstlane_b32 s1, v3
	s_cmp_eq_u32 s1, 0
	s_cbranch_scc1 .LBB0_21
; %bb.20:                               ;   in Loop: Header=BB0_17 Depth=1
	s_mov_b32 s1, 0
	s_sleep 1
	s_branch .LBB0_22
.LBB0_21:                               ;   in Loop: Header=BB0_17 Depth=1
	s_mov_b32 s1, -1
.LBB0_22:                               ; %Flow411
                                        ;   in Loop: Header=BB0_17 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s1
	s_cbranch_vccnz .LBB0_17
; %bb.23:
	global_load_b64 v[1:2], v[9:10], off
	s_and_saveexec_b32 s1, s0
	s_cbranch_execz .LBB0_27
; %bb.24:
	v_mov_b32_e32 v9, 0
	s_clause 0x2
	global_load_b64 v[5:6], v9, s[2:3] offset:40
	global_load_b64 v[10:11], v9, s[2:3] offset:24 glc
	global_load_b64 v[7:8], v9, s[2:3]
	s_waitcnt vmcnt(2)
	v_add_co_u32 v12, vcc_lo, v5, 1
	v_add_co_ci_u32_e32 v13, vcc_lo, 0, v6, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v3, vcc_lo, v12, s4
	v_add_co_ci_u32_e32 v4, vcc_lo, s5, v13, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cmp_eq_u64_e32 vcc_lo, 0, v[3:4]
	v_dual_cndmask_b32 v4, v4, v13 :: v_dual_cndmask_b32 v3, v3, v12
	v_and_b32_e32 v6, v4, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_and_b32_e32 v5, v3, v5
	v_mul_lo_u32 v6, v6, 24
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_hi_u32 v12, v5, 24
	v_mul_lo_u32 v5, v5, 24
	v_add_nc_u32_e32 v6, v12, v6
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_co_u32 v7, vcc_lo, v7, v5
	v_mov_b32_e32 v5, v10
	v_add_co_ci_u32_e32 v8, vcc_lo, v8, v6, vcc_lo
	v_mov_b32_e32 v6, v11
	global_store_b64 v[7:8], v[10:11], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[5:6], v9, v[3:6], s[2:3] offset:24 glc
	s_waitcnt vmcnt(0)
	v_cmp_ne_u64_e32 vcc_lo, v[5:6], v[10:11]
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_27
; %bb.25:                               ; %.preheader.i.i.i.preheader
	s_mov_b32 s0, 0
.LBB0_26:                               ; %.preheader.i.i.i
                                        ; =>This Inner Loop Header: Depth=1
	s_sleep 1
	global_store_b64 v[7:8], v[5:6], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[10:11], v9, v[3:6], s[2:3] offset:24 glc
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, v[10:11], v[5:6]
	v_dual_mov_b32 v5, v10 :: v_dual_mov_b32 v6, v11
	s_or_b32 s0, vcc_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s0
	s_cbranch_execnz .LBB0_26
.LBB0_27:                               ; %__ockl_printf_begin.exit
	s_or_b32 exec_lo, exec_lo, s1
	s_getpc_b64 s[4:5]
	s_add_u32 s4, s4, .str@rel32@lo+4
	s_addc_u32 s5, s5, .str@rel32@hi+12
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_lg_u64 s[4:5], 0
	s_cbranch_scc0 .LBB0_113
; %bb.28:
	s_waitcnt vmcnt(0)
	v_dual_mov_b32 v31, 0 :: v_dual_and_b32 v32, 2, v1
	v_dual_mov_b32 v4, v2 :: v_dual_and_b32 v3, -3, v1
	v_dual_mov_b32 v7, 2 :: v_dual_mov_b32 v8, 1
	s_mov_b64 s[6:7], 42
.LBB0_29:                               ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_32 Depth 2
                                        ;     Child Loop BB0_39 Depth 2
                                        ;     Child Loop BB0_47 Depth 2
                                        ;     Child Loop BB0_55 Depth 2
                                        ;     Child Loop BB0_63 Depth 2
                                        ;     Child Loop BB0_71 Depth 2
                                        ;     Child Loop BB0_79 Depth 2
                                        ;     Child Loop BB0_87 Depth 2
                                        ;     Child Loop BB0_95 Depth 2
                                        ;     Child Loop BB0_101 Depth 2
                                        ;     Child Loop BB0_110 Depth 2
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_lt_u64_e64 s0, s[6:7], 56
                                        ; implicit-def: $vgpr11_vgpr12
                                        ; implicit-def: $sgpr18
	s_and_b32 s0, s0, exec_lo
	s_cselect_b32 s8, s6, 56
	s_cselect_b32 s9, s7, 0
	s_cmp_gt_u32 s8, 7
	s_mov_b32 s0, -1
	s_cbranch_scc1 .LBB0_34
; %bb.30:                               ;   in Loop: Header=BB0_29 Depth=1
	v_mov_b32_e32 v11, 0
	v_mov_b32_e32 v12, 0
	s_cmp_eq_u32 s8, 0
	s_cbranch_scc1 .LBB0_33
; %bb.31:                               ; %.preheader31.i.preheader
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_lshl_b64 s[0:1], s[8:9], 3
	s_mov_b64 s[10:11], 0
	s_mov_b64 s[12:13], s[4:5]
.LBB0_32:                               ; %.preheader31.i
                                        ;   Parent Loop BB0_29 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	global_load_u8 v5, v31, s[12:13]
	s_waitcnt vmcnt(0)
	v_and_b32_e32 v30, 0xffff, v5
	s_delay_alu instid0(VALU_DEP_1)
	v_lshlrev_b64 v[5:6], s10, v[30:31]
	s_add_u32 s10, s10, 8
	s_addc_u32 s11, s11, 0
	s_add_u32 s12, s12, 1
	s_addc_u32 s13, s13, 0
	s_cmp_lg_u32 s0, s10
	v_or_b32_e32 v11, v5, v11
	v_or_b32_e32 v12, v6, v12
	s_cbranch_scc1 .LBB0_32
.LBB0_33:                               ; %Flow387
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_mov_b32 s0, 0
	s_mov_b32 s18, 0
.LBB0_34:                               ; %Flow389
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_and_not1_b32 vcc_lo, exec_lo, s0
	s_mov_b64 s[0:1], s[4:5]
	s_cbranch_vccnz .LBB0_36
; %bb.35:                               ;   in Loop: Header=BB0_29 Depth=1
	global_load_b64 v[11:12], v31, s[4:5]
	s_add_i32 s18, s8, -8
	s_add_u32 s0, s4, 8
	s_addc_u32 s1, s5, 0
.LBB0_36:                               ; %.loopexit32.i
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_cmp_gt_u32 s18, 7
	s_cbranch_scc1 .LBB0_41
; %bb.37:                               ;   in Loop: Header=BB0_29 Depth=1
	v_mov_b32_e32 v13, 0
	v_mov_b32_e32 v14, 0
	s_cmp_eq_u32 s18, 0
	s_cbranch_scc1 .LBB0_40
; %bb.38:                               ; %.preheader29.i.preheader
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_mov_b64 s[10:11], 0
	s_mov_b64 s[12:13], 0
.LBB0_39:                               ; %.preheader29.i
                                        ;   Parent Loop BB0_29 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_u32 s20, s0, s12
	s_addc_u32 s21, s1, s13
	s_add_u32 s12, s12, 1
	global_load_u8 v5, v31, s[20:21]
	s_addc_u32 s13, s13, 0
	s_waitcnt vmcnt(0)
	v_and_b32_e32 v30, 0xffff, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[5:6], s10, v[30:31]
	s_add_u32 s10, s10, 8
	s_addc_u32 s11, s11, 0
	s_cmp_lg_u32 s18, s12
	v_or_b32_e32 v13, v5, v13
	s_delay_alu instid0(VALU_DEP_2)
	v_or_b32_e32 v14, v6, v14
	s_cbranch_scc1 .LBB0_39
.LBB0_40:                               ; %Flow382
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_mov_b32 s10, 0
	s_mov_b32 s17, 0
	s_branch .LBB0_42
.LBB0_41:                               ;   in Loop: Header=BB0_29 Depth=1
	s_mov_b32 s10, -1
                                        ; implicit-def: $vgpr13_vgpr14
                                        ; implicit-def: $sgpr17
.LBB0_42:                               ; %Flow384
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s10
	s_cbranch_vccnz .LBB0_44
; %bb.43:                               ;   in Loop: Header=BB0_29 Depth=1
	global_load_b64 v[13:14], v31, s[0:1]
	s_add_i32 s17, s18, -8
	s_add_u32 s0, s0, 8
	s_addc_u32 s1, s1, 0
.LBB0_44:                               ; %.loopexit30.i
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_cmp_gt_u32 s17, 7
	s_cbranch_scc1 .LBB0_49
; %bb.45:                               ;   in Loop: Header=BB0_29 Depth=1
	v_mov_b32_e32 v15, 0
	v_mov_b32_e32 v16, 0
	s_cmp_eq_u32 s17, 0
	s_cbranch_scc1 .LBB0_48
; %bb.46:                               ; %.preheader27.i.preheader
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_mov_b64 s[10:11], 0
	s_mov_b64 s[12:13], 0
.LBB0_47:                               ; %.preheader27.i
                                        ;   Parent Loop BB0_29 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_u32 s18, s0, s12
	s_addc_u32 s19, s1, s13
	s_add_u32 s12, s12, 1
	global_load_u8 v5, v31, s[18:19]
	s_addc_u32 s13, s13, 0
	s_waitcnt vmcnt(0)
	v_and_b32_e32 v30, 0xffff, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[5:6], s10, v[30:31]
	s_add_u32 s10, s10, 8
	s_addc_u32 s11, s11, 0
	s_cmp_lg_u32 s17, s12
	v_or_b32_e32 v15, v5, v15
	s_delay_alu instid0(VALU_DEP_2)
	v_or_b32_e32 v16, v6, v16
	s_cbranch_scc1 .LBB0_47
.LBB0_48:                               ; %Flow377
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_mov_b32 s10, 0
	s_mov_b32 s18, 0
	s_branch .LBB0_50
.LBB0_49:                               ;   in Loop: Header=BB0_29 Depth=1
	s_mov_b32 s10, -1
                                        ; implicit-def: $sgpr18
.LBB0_50:                               ; %Flow379
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s10
	s_cbranch_vccnz .LBB0_52
; %bb.51:                               ;   in Loop: Header=BB0_29 Depth=1
	global_load_b64 v[15:16], v31, s[0:1]
	s_add_i32 s18, s17, -8
	s_add_u32 s0, s0, 8
	s_addc_u32 s1, s1, 0
.LBB0_52:                               ; %.loopexit28.i
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_cmp_gt_u32 s18, 7
	s_cbranch_scc1 .LBB0_57
; %bb.53:                               ;   in Loop: Header=BB0_29 Depth=1
	v_mov_b32_e32 v17, 0
	v_mov_b32_e32 v18, 0
	s_cmp_eq_u32 s18, 0
	s_cbranch_scc1 .LBB0_56
; %bb.54:                               ; %.preheader25.i.preheader
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_mov_b64 s[10:11], 0
	s_mov_b64 s[12:13], 0
.LBB0_55:                               ; %.preheader25.i
                                        ;   Parent Loop BB0_29 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_u32 s20, s0, s12
	s_addc_u32 s21, s1, s13
	s_add_u32 s12, s12, 1
	global_load_u8 v5, v31, s[20:21]
	s_addc_u32 s13, s13, 0
	s_waitcnt vmcnt(0)
	v_and_b32_e32 v30, 0xffff, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[5:6], s10, v[30:31]
	s_add_u32 s10, s10, 8
	s_addc_u32 s11, s11, 0
	s_cmp_lg_u32 s18, s12
	v_or_b32_e32 v17, v5, v17
	s_delay_alu instid0(VALU_DEP_2)
	v_or_b32_e32 v18, v6, v18
	s_cbranch_scc1 .LBB0_55
.LBB0_56:                               ; %Flow372
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_mov_b32 s10, 0
	s_mov_b32 s17, 0
	s_branch .LBB0_58
.LBB0_57:                               ;   in Loop: Header=BB0_29 Depth=1
	s_mov_b32 s10, -1
                                        ; implicit-def: $vgpr17_vgpr18
                                        ; implicit-def: $sgpr17
.LBB0_58:                               ; %Flow374
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s10
	s_cbranch_vccnz .LBB0_60
; %bb.59:                               ;   in Loop: Header=BB0_29 Depth=1
	global_load_b64 v[17:18], v31, s[0:1]
	s_add_i32 s17, s18, -8
	s_add_u32 s0, s0, 8
	s_addc_u32 s1, s1, 0
.LBB0_60:                               ; %.loopexit26.i
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_cmp_gt_u32 s17, 7
	s_cbranch_scc1 .LBB0_65
; %bb.61:                               ;   in Loop: Header=BB0_29 Depth=1
	v_mov_b32_e32 v19, 0
	v_mov_b32_e32 v20, 0
	s_cmp_eq_u32 s17, 0
	s_cbranch_scc1 .LBB0_64
; %bb.62:                               ; %.preheader23.i.preheader
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_mov_b64 s[10:11], 0
	s_mov_b64 s[12:13], 0
.LBB0_63:                               ; %.preheader23.i
                                        ;   Parent Loop BB0_29 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_u32 s18, s0, s12
	s_addc_u32 s19, s1, s13
	s_add_u32 s12, s12, 1
	global_load_u8 v5, v31, s[18:19]
	s_addc_u32 s13, s13, 0
	s_waitcnt vmcnt(0)
	v_and_b32_e32 v30, 0xffff, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[5:6], s10, v[30:31]
	s_add_u32 s10, s10, 8
	s_addc_u32 s11, s11, 0
	s_cmp_lg_u32 s17, s12
	v_or_b32_e32 v19, v5, v19
	s_delay_alu instid0(VALU_DEP_2)
	v_or_b32_e32 v20, v6, v20
	s_cbranch_scc1 .LBB0_63
.LBB0_64:                               ; %Flow367
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_mov_b32 s10, 0
	s_mov_b32 s18, 0
	s_branch .LBB0_66
.LBB0_65:                               ;   in Loop: Header=BB0_29 Depth=1
	s_mov_b32 s10, -1
                                        ; implicit-def: $sgpr18
.LBB0_66:                               ; %Flow369
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s10
	s_cbranch_vccnz .LBB0_68
; %bb.67:                               ;   in Loop: Header=BB0_29 Depth=1
	global_load_b64 v[19:20], v31, s[0:1]
	s_add_i32 s18, s17, -8
	s_add_u32 s0, s0, 8
	s_addc_u32 s1, s1, 0
.LBB0_68:                               ; %.loopexit24.i
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_cmp_gt_u32 s18, 7
	s_cbranch_scc1 .LBB0_73
; %bb.69:                               ;   in Loop: Header=BB0_29 Depth=1
	v_mov_b32_e32 v21, 0
	v_mov_b32_e32 v22, 0
	s_cmp_eq_u32 s18, 0
	s_cbranch_scc1 .LBB0_72
; %bb.70:                               ; %.preheader21.i.preheader
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_mov_b64 s[10:11], 0
	s_mov_b64 s[12:13], 0
.LBB0_71:                               ; %.preheader21.i
                                        ;   Parent Loop BB0_29 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_u32 s20, s0, s12
	s_addc_u32 s21, s1, s13
	s_add_u32 s12, s12, 1
	global_load_u8 v5, v31, s[20:21]
	s_addc_u32 s13, s13, 0
	s_waitcnt vmcnt(0)
	v_and_b32_e32 v30, 0xffff, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[5:6], s10, v[30:31]
	s_add_u32 s10, s10, 8
	s_addc_u32 s11, s11, 0
	s_cmp_lg_u32 s18, s12
	v_or_b32_e32 v21, v5, v21
	s_delay_alu instid0(VALU_DEP_2)
	v_or_b32_e32 v22, v6, v22
	s_cbranch_scc1 .LBB0_71
.LBB0_72:                               ; %Flow362
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_mov_b32 s10, 0
	s_mov_b32 s17, 0
	s_branch .LBB0_74
.LBB0_73:                               ;   in Loop: Header=BB0_29 Depth=1
	s_mov_b32 s10, -1
                                        ; implicit-def: $vgpr21_vgpr22
                                        ; implicit-def: $sgpr17
.LBB0_74:                               ; %Flow364
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s10
	s_cbranch_vccnz .LBB0_76
; %bb.75:                               ;   in Loop: Header=BB0_29 Depth=1
	global_load_b64 v[21:22], v31, s[0:1]
	s_add_i32 s17, s18, -8
	s_add_u32 s0, s0, 8
	s_addc_u32 s1, s1, 0
.LBB0_76:                               ; %.loopexit22.i
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_cmp_gt_u32 s17, 7
	s_cbranch_scc1 .LBB0_81
; %bb.77:                               ;   in Loop: Header=BB0_29 Depth=1
	v_mov_b32_e32 v23, 0
	v_mov_b32_e32 v24, 0
	s_cmp_eq_u32 s17, 0
	s_cbranch_scc1 .LBB0_80
; %bb.78:                               ; %.preheader.i.preheader
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_mov_b64 s[10:11], 0
	s_mov_b64 s[12:13], s[0:1]
.LBB0_79:                               ; %.preheader.i
                                        ;   Parent Loop BB0_29 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	global_load_u8 v5, v31, s[12:13]
	s_add_i32 s17, s17, -1
	s_waitcnt vmcnt(0)
	v_and_b32_e32 v30, 0xffff, v5
	s_delay_alu instid0(VALU_DEP_1)
	v_lshlrev_b64 v[5:6], s10, v[30:31]
	s_add_u32 s10, s10, 8
	s_addc_u32 s11, s11, 0
	s_add_u32 s12, s12, 1
	s_addc_u32 s13, s13, 0
	s_cmp_lg_u32 s17, 0
	v_or_b32_e32 v23, v5, v23
	v_or_b32_e32 v24, v6, v24
	s_cbranch_scc1 .LBB0_79
.LBB0_80:                               ; %Flow357
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_mov_b32 s10, 0
	s_branch .LBB0_82
.LBB0_81:                               ;   in Loop: Header=BB0_29 Depth=1
	s_mov_b32 s10, -1
.LBB0_82:                               ; %Flow359
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s10
	s_cbranch_vccnz .LBB0_84
; %bb.83:                               ;   in Loop: Header=BB0_29 Depth=1
	global_load_b64 v[23:24], v31, s[0:1]
.LBB0_84:                               ; %.loopexit.i
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_waitcnt vmcnt(0)
	v_dual_mov_b32 v30, v29 :: v_dual_mov_b32 v5, 0
	;;#ASMSTART
	;;#ASMEND
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_readfirstlane_b32 s0, v30
	v_mov_b32_e32 v6, 0
	v_cmp_eq_u32_e64 s0, s0, v30
	s_delay_alu instid0(VALU_DEP_1)
	s_and_saveexec_b32 s1, s0
	s_cbranch_execz .LBB0_90
; %bb.85:                               ;   in Loop: Header=BB0_29 Depth=1
	global_load_b64 v[27:28], v31, s[2:3] offset:24 glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	s_clause 0x1
	global_load_b64 v[5:6], v31, s[2:3] offset:40
	global_load_b64 v[9:10], v31, s[2:3]
	s_mov_b32 s10, exec_lo
	s_waitcnt vmcnt(1)
	v_and_b32_e32 v6, v6, v28
	v_and_b32_e32 v5, v5, v27
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_lo_u32 v6, v6, 24
	v_mul_hi_u32 v25, v5, 24
	v_mul_lo_u32 v5, v5, 24
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v6, v25, v6
	s_waitcnt vmcnt(0)
	v_add_co_u32 v5, vcc_lo, v9, v5
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v6, vcc_lo, v10, v6, vcc_lo
	global_load_b64 v[25:26], v[5:6], off glc
	s_waitcnt vmcnt(0)
	global_atomic_cmpswap_b64 v[5:6], v31, v[25:28], s[2:3] offset:24 glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	v_cmpx_ne_u64_e64 v[5:6], v[27:28]
	s_cbranch_execz .LBB0_89
; %bb.86:                               ; %.preheader3.i.i19.i.preheader
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_mov_b32 s11, 0
.LBB0_87:                               ; %.preheader3.i.i19.i
                                        ;   Parent Loop BB0_29 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_sleep 1
	s_clause 0x1
	global_load_b64 v[9:10], v31, s[2:3] offset:40
	global_load_b64 v[25:26], v31, s[2:3]
	v_dual_mov_b32 v28, v6 :: v_dual_mov_b32 v27, v5
	s_waitcnt vmcnt(1)
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_and_b32_e32 v9, v9, v27
	s_waitcnt vmcnt(0)
	v_mad_u64_u32 v[5:6], null, v9, 24, v[25:26]
	v_and_b32_e32 v25, v10, v28
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[9:10], null, v25, 24, v[6:7]
	v_mov_b32_e32 v6, v9
	global_load_b64 v[25:26], v[5:6], off glc
	s_waitcnt vmcnt(0)
	global_atomic_cmpswap_b64 v[5:6], v31, v[25:28], s[2:3] offset:24 glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	v_cmp_eq_u64_e32 vcc_lo, v[5:6], v[27:28]
	s_or_b32 s11, vcc_lo, s11
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s11
	s_cbranch_execnz .LBB0_87
; %bb.88:                               ; %Flow352
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_or_b32 exec_lo, exec_lo, s11
.LBB0_89:                               ; %Flow354
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s10
.LBB0_90:                               ; %.loopexit4.i.i14.i
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s1
	s_clause 0x1
	global_load_b64 v[9:10], v31, s[2:3] offset:40
	global_load_b128 v[25:28], v31, s[2:3]
	v_readfirstlane_b32 s10, v5
	v_readfirstlane_b32 s11, v6
	s_mov_b32 s19, exec_lo
	s_waitcnt vmcnt(1)
	v_readfirstlane_b32 s12, v9
	v_readfirstlane_b32 s13, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b64 s[12:13], s[10:11], s[12:13]
	s_mul_i32 s1, s13, 24
	s_mul_hi_u32 s17, s12, 24
	s_mul_i32 s18, s12, 24
	s_and_saveexec_b32 s20, s0
	s_cbranch_execz .LBB0_92
; %bb.91:                               ;   in Loop: Header=BB0_29 Depth=1
	v_dual_mov_b32 v5, s19 :: v_dual_mov_b32 v6, v31
	s_add_i32 s19, s17, s1
	s_waitcnt vmcnt(0)
	v_add_co_u32 v9, vcc_lo, v25, s18
	v_add_co_ci_u32_e32 v10, vcc_lo, s19, v26, vcc_lo
	global_store_b128 v[9:10], v[5:8], off offset:8
.LBB0_92:                               ;   in Loop: Header=BB0_29 Depth=1
	s_or_b32 exec_lo, exec_lo, s20
	v_cmp_gt_u64_e64 vcc_lo, s[6:7], 56
	v_or_b32_e32 v5, 0, v4
	v_or_b32_e32 v6, v3, v32
	s_lshl_b64 s[12:13], s[12:13], 12
	s_lshl_b32 s19, s8, 2
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_add_i32 s19, s19, 28
	v_dual_cndmask_b32 v10, v5, v4 :: v_dual_cndmask_b32 v5, v6, v3
	v_lshlrev_b64 v[3:4], 6, v[30:31]
	s_waitcnt vmcnt(0)
	v_add_co_u32 v6, vcc_lo, v27, s12
	v_add_co_ci_u32_e32 v28, vcc_lo, s13, v28, vcc_lo
	s_and_b32 s19, s19, 0x1e0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_co_u32 v27, vcc_lo, v6, v3
	v_and_or_b32 v9, 0xffffff1f, v5, s19
	v_add_co_ci_u32_e32 v28, vcc_lo, v28, v4, vcc_lo
	s_clause 0x3
	global_store_b128 v[27:28], v[9:12], off
	global_store_b128 v[27:28], v[13:16], off offset:16
	global_store_b128 v[27:28], v[17:20], off offset:32
	global_store_b128 v[27:28], v[21:24], off offset:48
	s_and_saveexec_b32 s12, s0
	s_cbranch_execz .LBB0_100
; %bb.93:                               ;   in Loop: Header=BB0_29 Depth=1
	s_clause 0x1
	global_load_b64 v[13:14], v31, s[2:3] offset:32 glc
	global_load_b64 v[3:4], v31, s[2:3] offset:40
	v_dual_mov_b32 v11, s10 :: v_dual_mov_b32 v12, s11
	s_waitcnt vmcnt(0)
	v_readfirstlane_b32 s20, v3
	v_readfirstlane_b32 s21, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b64 s[20:21], s[20:21], s[10:11]
	s_mul_i32 s13, s21, 24
	s_mul_hi_u32 s19, s20, 24
	s_mul_i32 s20, s20, 24
	s_add_i32 s19, s19, s13
	v_add_co_u32 v9, vcc_lo, v25, s20
	v_add_co_ci_u32_e32 v10, vcc_lo, s19, v26, vcc_lo
	s_mov_b32 s13, exec_lo
	global_store_b64 v[9:10], v[13:14], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[5:6], v31, v[11:14], s[2:3] offset:32 glc
	s_waitcnt vmcnt(0)
	v_cmpx_ne_u64_e64 v[5:6], v[13:14]
	s_cbranch_execz .LBB0_96
; %bb.94:                               ; %.preheader1.i.i17.i.preheader
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_mov_b32 s19, 0
.LBB0_95:                               ; %.preheader1.i.i17.i
                                        ;   Parent Loop BB0_29 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	v_dual_mov_b32 v3, s10 :: v_dual_mov_b32 v4, s11
	s_sleep 1
	global_store_b64 v[9:10], v[5:6], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[3:4], v31, v[3:6], s[2:3] offset:32 glc
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, v[3:4], v[5:6]
	v_dual_mov_b32 v6, v4 :: v_dual_mov_b32 v5, v3
	s_or_b32 s19, vcc_lo, s19
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s19
	s_cbranch_execnz .LBB0_95
.LBB0_96:                               ; %Flow350
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_or_b32 exec_lo, exec_lo, s13
	global_load_b64 v[3:4], v31, s[2:3] offset:16
	s_mov_b32 s19, exec_lo
	s_mov_b32 s13, exec_lo
	v_mbcnt_lo_u32_b32 v5, s19, 0
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 0, v5
	s_cbranch_execz .LBB0_98
; %bb.97:                               ;   in Loop: Header=BB0_29 Depth=1
	s_bcnt1_i32_b32 s19, s19
	s_delay_alu instid0(SALU_CYCLE_1)
	v_mov_b32_e32 v30, s19
	s_waitcnt vmcnt(0)
	global_atomic_add_u64 v[3:4], v[30:31], off offset:8
.LBB0_98:                               ;   in Loop: Header=BB0_29 Depth=1
	s_or_b32 exec_lo, exec_lo, s13
	s_waitcnt vmcnt(0)
	global_load_b64 v[5:6], v[3:4], off offset:16
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, 0, v[5:6]
	s_cbranch_vccnz .LBB0_100
; %bb.99:                               ;   in Loop: Header=BB0_29 Depth=1
	global_load_b32 v30, v[3:4], off offset:24
	s_waitcnt vmcnt(0)
	v_readfirstlane_b32 s13, v30
	s_waitcnt_vscnt null, 0x0
	global_store_b64 v[5:6], v[30:31], off
	s_and_b32 m0, s13, 0xff
	s_sendmsg sendmsg(MSG_INTERRUPT)
.LBB0_100:                              ; %Flow351
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_or_b32 exec_lo, exec_lo, s12
	s_add_i32 s17, s17, s1
	v_add_co_u32 v3, vcc_lo, v25, s18
	v_add_co_ci_u32_e32 v4, vcc_lo, s17, v26, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v3, vcc_lo, v3, 20
	v_add_co_ci_u32_e32 v4, vcc_lo, 0, v4, vcc_lo
.LBB0_101:                              ;   Parent Loop BB0_29 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	v_mov_b32_e32 v5, 1
	s_and_saveexec_b32 s1, s0
	s_cbranch_execz .LBB0_103
; %bb.102:                              ;   in Loop: Header=BB0_101 Depth=2
	global_load_b32 v5, v[3:4], off glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	v_and_b32_e32 v5, 1, v5
.LBB0_103:                              ;   in Loop: Header=BB0_101 Depth=2
	s_or_b32 exec_lo, exec_lo, s1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_readfirstlane_b32 s1, v5
	s_cmp_eq_u32 s1, 0
	s_cbranch_scc1 .LBB0_105
; %bb.104:                              ;   in Loop: Header=BB0_101 Depth=2
	s_mov_b32 s1, 0
	s_sleep 1
	s_branch .LBB0_106
.LBB0_105:                              ;   in Loop: Header=BB0_101 Depth=2
	s_mov_b32 s1, -1
.LBB0_106:                              ; %Flow345
                                        ;   in Loop: Header=BB0_101 Depth=2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s1
	s_cbranch_vccnz .LBB0_101
; %bb.107:                              ;   in Loop: Header=BB0_29 Depth=1
	global_load_b128 v[3:6], v[27:28], off
	s_and_saveexec_b32 s1, s0
	s_cbranch_execz .LBB0_111
; %bb.108:                              ;   in Loop: Header=BB0_29 Depth=1
	s_clause 0x2
	global_load_b64 v[5:6], v31, s[2:3] offset:40
	global_load_b64 v[13:14], v31, s[2:3] offset:24 glc
	global_load_b64 v[11:12], v31, s[2:3]
	s_waitcnt vmcnt(2)
	v_add_co_u32 v15, vcc_lo, v5, 1
	v_add_co_ci_u32_e32 v16, vcc_lo, 0, v6, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v9, vcc_lo, v15, s10
	v_add_co_ci_u32_e32 v10, vcc_lo, s11, v16, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cmp_eq_u64_e32 vcc_lo, 0, v[9:10]
	v_dual_cndmask_b32 v10, v10, v16 :: v_dual_cndmask_b32 v9, v9, v15
	v_and_b32_e32 v6, v10, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_and_b32_e32 v5, v9, v5
	v_mul_hi_u32 v15, v5, 24
	v_mul_lo_u32 v5, v5, 24
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_add_co_u32 v5, vcc_lo, v11, v5
	v_mov_b32_e32 v11, v13
	v_mul_lo_u32 v6, v6, 24
	v_add_nc_u32_e32 v6, v15, v6
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e32 v6, vcc_lo, v12, v6, vcc_lo
	v_mov_b32_e32 v12, v14
	global_store_b64 v[5:6], v[13:14], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[11:12], v31, v[9:12], s[2:3] offset:24 glc
	s_waitcnt vmcnt(0)
	v_cmp_ne_u64_e32 vcc_lo, v[11:12], v[13:14]
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_111
; %bb.109:                              ; %.preheader.i.i16.i.preheader
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_mov_b32 s0, 0
.LBB0_110:                              ; %.preheader.i.i16.i
                                        ;   Parent Loop BB0_29 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_sleep 1
	global_store_b64 v[5:6], v[11:12], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[13:14], v31, v[9:12], s[2:3] offset:24 glc
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, v[13:14], v[11:12]
	v_dual_mov_b32 v11, v13 :: v_dual_mov_b32 v12, v14
	s_or_b32 s0, vcc_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s0
	s_cbranch_execnz .LBB0_110
.LBB0_111:                              ; %__ockl_hostcall_preview.exit20.i
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_or_b32 exec_lo, exec_lo, s1
	s_sub_u32 s6, s6, s8
	s_subb_u32 s7, s7, s9
	s_add_u32 s4, s4, s8
	s_addc_u32 s5, s5, s9
	s_cmp_lg_u64 s[6:7], 0
	s_cbranch_scc1 .LBB0_29
; %bb.112:                              ; %Flow390
	s_mov_b32 s0, 0
	s_branch .LBB0_114
.LBB0_113:
	s_mov_b32 s0, -1
                                        ; implicit-def: $vgpr3_vgpr4
.LBB0_114:                              ; %Flow405
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_b32 vcc_lo, exec_lo, s0
	s_cbranch_vccz .LBB0_143
; %bb.115:
	s_waitcnt vmcnt(0)
	v_mov_b32_e32 v3, v29
	v_mov_b32_e32 v9, 0
	v_mov_b32_e32 v10, 0
	;;#ASMSTART
	;;#ASMEND
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_readfirstlane_b32 s0, v3
	v_cmp_eq_u32_e64 s0, s0, v3
	s_delay_alu instid0(VALU_DEP_1)
	s_and_saveexec_b32 s1, s0
	s_cbranch_execz .LBB0_121
; %bb.116:
	v_mov_b32_e32 v4, 0
	s_mov_b32 s4, exec_lo
	global_load_b64 v[7:8], v4, s[2:3] offset:24 glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	s_clause 0x1
	global_load_b64 v[5:6], v4, s[2:3] offset:40
	global_load_b64 v[9:10], v4, s[2:3]
	s_waitcnt vmcnt(1)
	v_and_b32_e32 v5, v5, v7
	v_and_b32_e32 v6, v6, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_hi_u32 v11, v5, 24
	v_mul_lo_u32 v6, v6, 24
	v_mul_lo_u32 v5, v5, 24
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v6, v11, v6
	s_waitcnt vmcnt(0)
	v_add_co_u32 v5, vcc_lo, v9, v5
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v6, vcc_lo, v10, v6, vcc_lo
	global_load_b64 v[5:6], v[5:6], off glc
	s_waitcnt vmcnt(0)
	global_atomic_cmpswap_b64 v[9:10], v4, v[5:8], s[2:3] offset:24 glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	v_cmpx_ne_u64_e64 v[9:10], v[7:8]
	s_cbranch_execz .LBB0_120
; %bb.117:                              ; %.preheader3.i.i.i9.preheader
	s_mov_b32 s5, 0
.LBB0_118:                              ; %.preheader3.i.i.i9
                                        ; =>This Inner Loop Header: Depth=1
	s_sleep 1
	s_clause 0x1
	global_load_b64 v[5:6], v4, s[2:3] offset:40
	global_load_b64 v[11:12], v4, s[2:3]
	v_dual_mov_b32 v7, v9 :: v_dual_mov_b32 v8, v10
	s_waitcnt vmcnt(1)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_and_b32_e32 v5, v5, v7
	v_and_b32_e32 v6, v6, v8
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[9:10], null, v5, 24, v[11:12]
	v_mov_b32_e32 v5, v10
	s_delay_alu instid0(VALU_DEP_1)
	v_mad_u64_u32 v[10:11], null, v6, 24, v[5:6]
	global_load_b64 v[5:6], v[9:10], off glc
	s_waitcnt vmcnt(0)
	global_atomic_cmpswap_b64 v[9:10], v4, v[5:8], s[2:3] offset:24 glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	v_cmp_eq_u64_e32 vcc_lo, v[9:10], v[7:8]
	s_or_b32 s5, vcc_lo, s5
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s5
	s_cbranch_execnz .LBB0_118
; %bb.119:                              ; %Flow402
	s_or_b32 exec_lo, exec_lo, s5
.LBB0_120:                              ; %Flow404
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s4
.LBB0_121:                              ; %.loopexit4.i.i.i4
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s1
	v_mov_b32_e32 v4, 0
	v_readfirstlane_b32 s4, v9
	v_readfirstlane_b32 s5, v10
	s_mov_b32 s10, exec_lo
	s_clause 0x1
	global_load_b64 v[11:12], v4, s[2:3] offset:40
	global_load_b128 v[5:8], v4, s[2:3]
	s_waitcnt vmcnt(1)
	v_readfirstlane_b32 s6, v11
	v_readfirstlane_b32 s7, v12
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b64 s[6:7], s[4:5], s[6:7]
	s_mul_i32 s1, s7, 24
	s_mul_hi_u32 s8, s6, 24
	s_mul_i32 s9, s6, 24
	s_and_saveexec_b32 s11, s0
	s_cbranch_execz .LBB0_123
; %bb.122:
	v_dual_mov_b32 v9, s10 :: v_dual_mov_b32 v10, v4
	s_add_i32 s10, s8, s1
	s_waitcnt vmcnt(0)
	v_add_co_u32 v13, vcc_lo, v5, s9
	v_add_co_ci_u32_e32 v14, vcc_lo, s10, v6, vcc_lo
	v_dual_mov_b32 v11, 2 :: v_dual_mov_b32 v12, 1
	global_store_b128 v[13:14], v[9:12], off offset:8
.LBB0_123:
	s_or_b32 exec_lo, exec_lo, s11
	s_lshl_b64 s[6:7], s[6:7], 12
	v_lshlrev_b64 v[9:10], 6, v[3:4]
	s_waitcnt vmcnt(0)
	v_add_co_u32 v3, vcc_lo, v7, s6
	v_add_co_ci_u32_e32 v8, vcc_lo, s7, v8, vcc_lo
	s_mov_b32 s20, 0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_u32 v7, vcc_lo, v3, v9
	s_mov_b32 s21, s20
	s_mov_b32 s22, s20
	s_mov_b32 s23, s20
	v_and_or_b32 v1, 0xffffff1f, v1, 32
	v_add_co_ci_u32_e32 v8, vcc_lo, v8, v10, vcc_lo
	v_mov_b32_e32 v3, v4
	v_dual_mov_b32 v9, s20 :: v_dual_mov_b32 v12, s23
	v_dual_mov_b32 v10, s21 :: v_dual_mov_b32 v11, s22
	s_clause 0x3
	global_store_b128 v[7:8], v[1:4], off
	global_store_b128 v[7:8], v[9:12], off offset:16
	global_store_b128 v[7:8], v[9:12], off offset:32
	global_store_b128 v[7:8], v[9:12], off offset:48
	s_and_saveexec_b32 s6, s0
	s_cbranch_execz .LBB0_131
; %bb.124:
	v_dual_mov_b32 v11, 0 :: v_dual_mov_b32 v12, s4
	v_mov_b32_e32 v13, s5
	s_clause 0x1
	global_load_b64 v[14:15], v11, s[2:3] offset:32 glc
	global_load_b64 v[1:2], v11, s[2:3] offset:40
	s_waitcnt vmcnt(0)
	v_readfirstlane_b32 s10, v1
	v_readfirstlane_b32 s11, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b64 s[10:11], s[10:11], s[4:5]
	s_mul_i32 s7, s11, 24
	s_mul_hi_u32 s11, s10, 24
	s_mul_i32 s10, s10, 24
	s_add_i32 s11, s11, s7
	v_add_co_u32 v9, vcc_lo, v5, s10
	v_add_co_ci_u32_e32 v10, vcc_lo, s11, v6, vcc_lo
	s_mov_b32 s7, exec_lo
	global_store_b64 v[9:10], v[14:15], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[3:4], v11, v[12:15], s[2:3] offset:32 glc
	s_waitcnt vmcnt(0)
	v_cmpx_ne_u64_e64 v[3:4], v[14:15]
	s_cbranch_execz .LBB0_127
; %bb.125:                              ; %.preheader1.i.i.i7.preheader
	s_mov_b32 s10, 0
.LBB0_126:                              ; %.preheader1.i.i.i7
                                        ; =>This Inner Loop Header: Depth=1
	v_dual_mov_b32 v1, s4 :: v_dual_mov_b32 v2, s5
	s_sleep 1
	global_store_b64 v[9:10], v[3:4], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[1:2], v11, v[1:4], s[2:3] offset:32 glc
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, v[1:2], v[3:4]
	v_dual_mov_b32 v4, v2 :: v_dual_mov_b32 v3, v1
	s_or_b32 s10, vcc_lo, s10
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s10
	s_cbranch_execnz .LBB0_126
.LBB0_127:                              ; %Flow400
	s_or_b32 exec_lo, exec_lo, s7
	v_mov_b32_e32 v4, 0
	s_mov_b32 s10, exec_lo
	s_mov_b32 s7, exec_lo
	v_mbcnt_lo_u32_b32 v3, s10, 0
	global_load_b64 v[1:2], v4, s[2:3] offset:16
	v_cmpx_eq_u32_e32 0, v3
	s_cbranch_execz .LBB0_129
; %bb.128:
	s_bcnt1_i32_b32 s10, s10
	s_delay_alu instid0(SALU_CYCLE_1)
	v_mov_b32_e32 v3, s10
	s_waitcnt vmcnt(0)
	global_atomic_add_u64 v[1:2], v[3:4], off offset:8
.LBB0_129:
	s_or_b32 exec_lo, exec_lo, s7
	s_waitcnt vmcnt(0)
	global_load_b64 v[3:4], v[1:2], off offset:16
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, 0, v[3:4]
	s_cbranch_vccnz .LBB0_131
; %bb.130:
	global_load_b32 v1, v[1:2], off offset:24
	v_mov_b32_e32 v2, 0
	s_waitcnt vmcnt(0)
	v_readfirstlane_b32 s7, v1
	s_waitcnt_vscnt null, 0x0
	global_store_b64 v[3:4], v[1:2], off
	s_and_b32 m0, s7, 0xff
	s_sendmsg sendmsg(MSG_INTERRUPT)
.LBB0_131:                              ; %Flow401
	s_or_b32 exec_lo, exec_lo, s6
	s_add_i32 s8, s8, s1
	v_add_co_u32 v1, vcc_lo, v5, s9
	v_add_co_ci_u32_e32 v2, vcc_lo, s8, v6, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v1, vcc_lo, v1, 20
	v_add_co_ci_u32_e32 v2, vcc_lo, 0, v2, vcc_lo
.LBB0_132:                              ; =>This Inner Loop Header: Depth=1
	v_mov_b32_e32 v3, 1
	s_and_saveexec_b32 s1, s0
	s_cbranch_execz .LBB0_134
; %bb.133:                              ;   in Loop: Header=BB0_132 Depth=1
	global_load_b32 v3, v[1:2], off glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	v_and_b32_e32 v3, 1, v3
.LBB0_134:                              ;   in Loop: Header=BB0_132 Depth=1
	s_or_b32 exec_lo, exec_lo, s1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_readfirstlane_b32 s1, v3
	s_cmp_eq_u32 s1, 0
	s_cbranch_scc1 .LBB0_136
; %bb.135:                              ;   in Loop: Header=BB0_132 Depth=1
	s_mov_b32 s1, 0
	s_sleep 1
	s_branch .LBB0_137
.LBB0_136:                              ;   in Loop: Header=BB0_132 Depth=1
	s_mov_b32 s1, -1
.LBB0_137:                              ; %Flow395
                                        ;   in Loop: Header=BB0_132 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s1
	s_cbranch_vccnz .LBB0_132
; %bb.138:
	global_load_b128 v[3:6], v[7:8], off
	s_and_saveexec_b32 s1, s0
	s_cbranch_execz .LBB0_142
; %bb.139:
	v_mov_b32_e32 v9, 0
	s_clause 0x2
	global_load_b64 v[1:2], v9, s[2:3] offset:40
	global_load_b64 v[10:11], v9, s[2:3] offset:24 glc
	global_load_b64 v[7:8], v9, s[2:3]
	s_waitcnt vmcnt(2)
	v_add_co_u32 v12, vcc_lo, v1, 1
	v_add_co_ci_u32_e32 v13, vcc_lo, 0, v2, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v5, vcc_lo, v12, s4
	v_add_co_ci_u32_e32 v6, vcc_lo, s5, v13, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cmp_eq_u64_e32 vcc_lo, 0, v[5:6]
	v_dual_cndmask_b32 v6, v6, v13 :: v_dual_cndmask_b32 v5, v5, v12
	v_and_b32_e32 v2, v6, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_and_b32_e32 v1, v5, v1
	v_mul_lo_u32 v2, v2, 24
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_hi_u32 v12, v1, 24
	v_mul_lo_u32 v1, v1, 24
	v_add_nc_u32_e32 v2, v12, v2
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_co_u32 v1, vcc_lo, v7, v1
	v_mov_b32_e32 v7, v10
	v_add_co_ci_u32_e32 v2, vcc_lo, v8, v2, vcc_lo
	v_mov_b32_e32 v8, v11
	global_store_b64 v[1:2], v[10:11], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[7:8], v9, v[5:8], s[2:3] offset:24 glc
	s_waitcnt vmcnt(0)
	v_cmp_ne_u64_e32 vcc_lo, v[7:8], v[10:11]
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_142
; %bb.140:                              ; %.preheader.i.i.i6.preheader
	s_mov_b32 s0, 0
.LBB0_141:                              ; %.preheader.i.i.i6
                                        ; =>This Inner Loop Header: Depth=1
	s_sleep 1
	global_store_b64 v[1:2], v[7:8], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[10:11], v9, v[5:8], s[2:3] offset:24 glc
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, v[10:11], v[7:8]
	v_dual_mov_b32 v7, v10 :: v_dual_mov_b32 v8, v11
	s_or_b32 s0, vcc_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s0
	s_cbranch_execnz .LBB0_141
.LBB0_142:                              ; %__ockl_hostcall_preview.exit.i
	s_or_b32 exec_lo, exec_lo, s1
.LBB0_143:                              ; %__ockl_printf_append_string_n.exit
	s_waitcnt vmcnt(0)
	v_mov_b32_e32 v5, v29
	v_mov_b32_e32 v1, 0
	v_mov_b32_e32 v2, 0
	;;#ASMSTART
	;;#ASMEND
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_readfirstlane_b32 s0, v5
	v_cmp_eq_u32_e64 s0, s0, v5
	s_delay_alu instid0(VALU_DEP_1)
	s_and_saveexec_b32 s1, s0
	s_cbranch_execz .LBB0_149
; %bb.144:
	v_mov_b32_e32 v6, 0
	s_mov_b32 s4, exec_lo
	global_load_b64 v[9:10], v6, s[2:3] offset:24 glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	s_clause 0x1
	global_load_b64 v[1:2], v6, s[2:3] offset:40
	global_load_b64 v[7:8], v6, s[2:3]
	s_waitcnt vmcnt(1)
	v_and_b32_e32 v1, v1, v9
	v_and_b32_e32 v2, v2, v10
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_hi_u32 v11, v1, 24
	v_mul_lo_u32 v2, v2, 24
	v_mul_lo_u32 v1, v1, 24
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v2, v11, v2
	s_waitcnt vmcnt(0)
	v_add_co_u32 v1, vcc_lo, v7, v1
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v2, vcc_lo, v8, v2, vcc_lo
	global_load_b64 v[7:8], v[1:2], off glc
	s_waitcnt vmcnt(0)
	global_atomic_cmpswap_b64 v[1:2], v6, v[7:10], s[2:3] offset:24 glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	v_cmpx_ne_u64_e64 v[1:2], v[9:10]
	s_cbranch_execz .LBB0_148
; %bb.145:                              ; %.preheader3.i.i.i16.preheader
	s_mov_b32 s5, 0
.LBB0_146:                              ; %.preheader3.i.i.i16
                                        ; =>This Inner Loop Header: Depth=1
	s_sleep 1
	s_clause 0x1
	global_load_b64 v[7:8], v6, s[2:3] offset:40
	global_load_b64 v[11:12], v6, s[2:3]
	v_dual_mov_b32 v10, v2 :: v_dual_mov_b32 v9, v1
	s_waitcnt vmcnt(1)
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_and_b32_e32 v7, v7, v9
	s_waitcnt vmcnt(0)
	v_mad_u64_u32 v[1:2], null, v7, 24, v[11:12]
	v_and_b32_e32 v11, v8, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[7:8], null, v11, 24, v[2:3]
	v_mov_b32_e32 v2, v7
	global_load_b64 v[7:8], v[1:2], off glc
	s_waitcnt vmcnt(0)
	global_atomic_cmpswap_b64 v[1:2], v6, v[7:10], s[2:3] offset:24 glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	v_cmp_eq_u64_e32 vcc_lo, v[1:2], v[9:10]
	s_or_b32 s5, vcc_lo, s5
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s5
	s_cbranch_execnz .LBB0_146
; %bb.147:                              ; %Flow338
	s_or_b32 exec_lo, exec_lo, s5
.LBB0_148:                              ; %Flow340
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s4
.LBB0_149:                              ; %.loopexit4.i.i.i10
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s1
	v_mov_b32_e32 v6, 0
	v_readfirstlane_b32 s4, v1
	v_readfirstlane_b32 s5, v2
	s_mov_b32 s10, exec_lo
	s_clause 0x1
	global_load_b64 v[11:12], v6, s[2:3] offset:40
	global_load_b128 v[7:10], v6, s[2:3]
	s_waitcnt vmcnt(1)
	v_readfirstlane_b32 s6, v11
	v_readfirstlane_b32 s7, v12
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b64 s[6:7], s[4:5], s[6:7]
	s_mul_i32 s1, s7, 24
	s_mul_hi_u32 s8, s6, 24
	s_mul_i32 s9, s6, 24
	s_and_saveexec_b32 s11, s0
	s_cbranch_execz .LBB0_151
; %bb.150:
	v_dual_mov_b32 v11, s10 :: v_dual_mov_b32 v12, v6
	s_add_i32 s10, s8, s1
	s_waitcnt vmcnt(0)
	v_add_co_u32 v1, vcc_lo, v7, s9
	v_add_co_ci_u32_e32 v2, vcc_lo, s10, v8, vcc_lo
	v_dual_mov_b32 v13, 2 :: v_dual_mov_b32 v14, 1
	global_store_b128 v[1:2], v[11:14], off offset:8
.LBB0_151:
	s_or_b32 exec_lo, exec_lo, s11
	s_lshl_b64 s[6:7], s[6:7], 12
	v_lshlrev_b64 v[1:2], 6, v[5:6]
	s_waitcnt vmcnt(0)
	v_add_co_u32 v5, vcc_lo, v9, s6
	v_add_co_ci_u32_e32 v10, vcc_lo, s7, v10, vcc_lo
	s_mov_b32 s20, 0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_u32 v9, vcc_lo, v5, v1
	s_mov_b32 s21, s20
	s_mov_b32 s22, s20
	s_mov_b32 s23, s20
	v_and_or_b32 v3, 0xffffff1f, v3, 32
	v_add_co_ci_u32_e32 v10, vcc_lo, v10, v2, vcc_lo
	v_mov_b32_e32 v5, 4
	v_dual_mov_b32 v11, s20 :: v_dual_mov_b32 v14, s23
	v_dual_mov_b32 v12, s21 :: v_dual_mov_b32 v13, s22
	s_clause 0x3
	global_store_b128 v[9:10], v[3:6], off
	global_store_b128 v[9:10], v[11:14], off offset:16
	global_store_b128 v[9:10], v[11:14], off offset:32
	global_store_b128 v[9:10], v[11:14], off offset:48
	s_and_saveexec_b32 s6, s0
	s_cbranch_execz .LBB0_159
; %bb.152:
	v_dual_mov_b32 v11, 0 :: v_dual_mov_b32 v12, s4
	v_mov_b32_e32 v13, s5
	s_clause 0x1
	global_load_b64 v[14:15], v11, s[2:3] offset:32 glc
	global_load_b64 v[1:2], v11, s[2:3] offset:40
	s_waitcnt vmcnt(0)
	v_readfirstlane_b32 s10, v1
	v_readfirstlane_b32 s11, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b64 s[10:11], s[10:11], s[4:5]
	s_mul_i32 s7, s11, 24
	s_mul_hi_u32 s11, s10, 24
	s_mul_i32 s10, s10, 24
	s_add_i32 s11, s11, s7
	v_add_co_u32 v5, vcc_lo, v7, s10
	v_add_co_ci_u32_e32 v6, vcc_lo, s11, v8, vcc_lo
	s_mov_b32 s7, exec_lo
	global_store_b64 v[5:6], v[14:15], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[3:4], v11, v[12:15], s[2:3] offset:32 glc
	s_waitcnt vmcnt(0)
	v_cmpx_ne_u64_e64 v[3:4], v[14:15]
	s_cbranch_execz .LBB0_155
; %bb.153:                              ; %.preheader1.i.i.i14.preheader
	s_mov_b32 s10, 0
.LBB0_154:                              ; %.preheader1.i.i.i14
                                        ; =>This Inner Loop Header: Depth=1
	v_dual_mov_b32 v1, s4 :: v_dual_mov_b32 v2, s5
	s_sleep 1
	global_store_b64 v[5:6], v[3:4], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[1:2], v11, v[1:4], s[2:3] offset:32 glc
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, v[1:2], v[3:4]
	v_dual_mov_b32 v4, v2 :: v_dual_mov_b32 v3, v1
	s_or_b32 s10, vcc_lo, s10
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s10
	s_cbranch_execnz .LBB0_154
.LBB0_155:                              ; %Flow336
	s_or_b32 exec_lo, exec_lo, s7
	v_mov_b32_e32 v4, 0
	s_mov_b32 s10, exec_lo
	s_mov_b32 s7, exec_lo
	v_mbcnt_lo_u32_b32 v3, s10, 0
	global_load_b64 v[1:2], v4, s[2:3] offset:16
	v_cmpx_eq_u32_e32 0, v3
	s_cbranch_execz .LBB0_157
; %bb.156:
	s_bcnt1_i32_b32 s10, s10
	s_delay_alu instid0(SALU_CYCLE_1)
	v_mov_b32_e32 v3, s10
	s_waitcnt vmcnt(0)
	global_atomic_add_u64 v[1:2], v[3:4], off offset:8
.LBB0_157:
	s_or_b32 exec_lo, exec_lo, s7
	s_waitcnt vmcnt(0)
	global_load_b64 v[3:4], v[1:2], off offset:16
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, 0, v[3:4]
	s_cbranch_vccnz .LBB0_159
; %bb.158:
	global_load_b32 v1, v[1:2], off offset:24
	v_mov_b32_e32 v2, 0
	s_waitcnt vmcnt(0)
	v_readfirstlane_b32 s7, v1
	s_waitcnt_vscnt null, 0x0
	global_store_b64 v[3:4], v[1:2], off
	s_and_b32 m0, s7, 0xff
	s_sendmsg sendmsg(MSG_INTERRUPT)
.LBB0_159:                              ; %Flow337
	s_or_b32 exec_lo, exec_lo, s6
	s_add_i32 s8, s8, s1
	v_add_co_u32 v1, vcc_lo, v7, s9
	v_add_co_ci_u32_e32 v2, vcc_lo, s8, v8, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v1, vcc_lo, v1, 20
	v_add_co_ci_u32_e32 v2, vcc_lo, 0, v2, vcc_lo
.LBB0_160:                              ; =>This Inner Loop Header: Depth=1
	v_mov_b32_e32 v3, 1
	s_and_saveexec_b32 s1, s0
	s_cbranch_execz .LBB0_162
; %bb.161:                              ;   in Loop: Header=BB0_160 Depth=1
	global_load_b32 v3, v[1:2], off glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	v_and_b32_e32 v3, 1, v3
.LBB0_162:                              ;   in Loop: Header=BB0_160 Depth=1
	s_or_b32 exec_lo, exec_lo, s1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_readfirstlane_b32 s1, v3
	s_cmp_eq_u32 s1, 0
	s_cbranch_scc1 .LBB0_164
; %bb.163:                              ;   in Loop: Header=BB0_160 Depth=1
	s_mov_b32 s1, 0
	s_sleep 1
	s_branch .LBB0_165
.LBB0_164:                              ;   in Loop: Header=BB0_160 Depth=1
	s_mov_b32 s1, -1
.LBB0_165:                              ; %Flow331
                                        ;   in Loop: Header=BB0_160 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s1
	s_cbranch_vccnz .LBB0_160
; %bb.166:
	global_load_b64 v[1:2], v[9:10], off
	s_and_saveexec_b32 s1, s0
	s_cbranch_execz .LBB0_170
; %bb.167:
	v_mov_b32_e32 v9, 0
	s_clause 0x2
	global_load_b64 v[5:6], v9, s[2:3] offset:40
	global_load_b64 v[10:11], v9, s[2:3] offset:24 glc
	global_load_b64 v[7:8], v9, s[2:3]
	s_waitcnt vmcnt(2)
	v_add_co_u32 v12, vcc_lo, v5, 1
	v_add_co_ci_u32_e32 v13, vcc_lo, 0, v6, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v3, vcc_lo, v12, s4
	v_add_co_ci_u32_e32 v4, vcc_lo, s5, v13, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cmp_eq_u64_e32 vcc_lo, 0, v[3:4]
	v_dual_cndmask_b32 v4, v4, v13 :: v_dual_cndmask_b32 v3, v3, v12
	v_and_b32_e32 v6, v4, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_and_b32_e32 v5, v3, v5
	v_mul_lo_u32 v6, v6, 24
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_hi_u32 v12, v5, 24
	v_mul_lo_u32 v5, v5, 24
	v_add_nc_u32_e32 v6, v12, v6
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_co_u32 v7, vcc_lo, v7, v5
	v_mov_b32_e32 v5, v10
	v_add_co_ci_u32_e32 v8, vcc_lo, v8, v6, vcc_lo
	v_mov_b32_e32 v6, v11
	global_store_b64 v[7:8], v[10:11], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[5:6], v9, v[3:6], s[2:3] offset:24 glc
	s_waitcnt vmcnt(0)
	v_cmp_ne_u64_e32 vcc_lo, v[5:6], v[10:11]
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_170
; %bb.168:                              ; %.preheader.i.i.i13.preheader
	s_mov_b32 s0, 0
.LBB0_169:                              ; %.preheader.i.i.i13
                                        ; =>This Inner Loop Header: Depth=1
	s_sleep 1
	global_store_b64 v[7:8], v[5:6], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[10:11], v9, v[3:6], s[2:3] offset:24 glc
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, v[10:11], v[5:6]
	v_dual_mov_b32 v5, v10 :: v_dual_mov_b32 v6, v11
	s_or_b32 s0, vcc_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s0
	s_cbranch_execnz .LBB0_169
.LBB0_170:                              ; %__ockl_printf_append_args.exit
	s_or_b32 exec_lo, exec_lo, s1
	v_mov_b32_e32 v9, v29
	v_mov_b32_e32 v3, 0
	v_mov_b32_e32 v4, 0
	;;#ASMSTART
	;;#ASMEND
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_readfirstlane_b32 s0, v9
	v_cmp_eq_u32_e64 s0, s0, v9
	s_delay_alu instid0(VALU_DEP_1)
	s_and_saveexec_b32 s1, s0
	s_cbranch_execz .LBB0_176
; %bb.171:
	v_mov_b32_e32 v5, 0
	s_mov_b32 s4, exec_lo
	global_load_b64 v[12:13], v5, s[2:3] offset:24 glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	s_clause 0x1
	global_load_b64 v[3:4], v5, s[2:3] offset:40
	global_load_b64 v[6:7], v5, s[2:3]
	s_waitcnt vmcnt(1)
	v_and_b32_e32 v4, v4, v13
	v_and_b32_e32 v3, v3, v12
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_lo_u32 v4, v4, 24
	v_mul_hi_u32 v8, v3, 24
	v_mul_lo_u32 v3, v3, 24
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v4, v8, v4
	s_waitcnt vmcnt(0)
	v_add_co_u32 v3, vcc_lo, v6, v3
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v4, vcc_lo, v7, v4, vcc_lo
	global_load_b64 v[10:11], v[3:4], off glc
	s_waitcnt vmcnt(0)
	global_atomic_cmpswap_b64 v[3:4], v5, v[10:13], s[2:3] offset:24 glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	v_cmpx_ne_u64_e64 v[3:4], v[12:13]
	s_cbranch_execz .LBB0_175
; %bb.172:                              ; %.preheader3.i.i.i23.preheader
	s_mov_b32 s5, 0
.LBB0_173:                              ; %.preheader3.i.i.i23
                                        ; =>This Inner Loop Header: Depth=1
	s_sleep 1
	s_clause 0x1
	global_load_b64 v[6:7], v5, s[2:3] offset:40
	global_load_b64 v[10:11], v5, s[2:3]
	v_dual_mov_b32 v13, v4 :: v_dual_mov_b32 v12, v3
	s_waitcnt vmcnt(1)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_and_b32_e32 v6, v6, v12
	v_and_b32_e32 v8, v7, v13
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[3:4], null, v6, 24, v[10:11]
	v_mad_u64_u32 v[6:7], null, v8, 24, v[4:5]
	s_delay_alu instid0(VALU_DEP_1)
	v_mov_b32_e32 v4, v6
	global_load_b64 v[10:11], v[3:4], off glc
	s_waitcnt vmcnt(0)
	global_atomic_cmpswap_b64 v[3:4], v5, v[10:13], s[2:3] offset:24 glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	v_cmp_eq_u64_e32 vcc_lo, v[3:4], v[12:13]
	s_or_b32 s5, vcc_lo, s5
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s5
	s_cbranch_execnz .LBB0_173
; %bb.174:                              ; %Flow324
	s_or_b32 exec_lo, exec_lo, s5
.LBB0_175:                              ; %Flow326
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s4
.LBB0_176:                              ; %.loopexit4.i.i.i17
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s1
	v_mov_b32_e32 v10, 0
	v_readfirstlane_b32 s4, v3
	v_readfirstlane_b32 s5, v4
	s_mov_b32 s10, exec_lo
	s_clause 0x1
	global_load_b64 v[11:12], v10, s[2:3] offset:40
	global_load_b128 v[5:8], v10, s[2:3]
	s_waitcnt vmcnt(1)
	v_readfirstlane_b32 s6, v11
	v_readfirstlane_b32 s7, v12
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b64 s[6:7], s[4:5], s[6:7]
	s_mul_i32 s1, s7, 24
	s_mul_hi_u32 s8, s6, 24
	s_mul_i32 s9, s6, 24
	s_and_saveexec_b32 s11, s0
	s_cbranch_execz .LBB0_178
; %bb.177:
	v_dual_mov_b32 v11, s10 :: v_dual_mov_b32 v12, v10
	s_add_i32 s10, s8, s1
	s_waitcnt vmcnt(0)
	v_add_co_u32 v3, vcc_lo, v5, s9
	v_add_co_ci_u32_e32 v4, vcc_lo, s10, v6, vcc_lo
	v_dual_mov_b32 v13, 2 :: v_dual_mov_b32 v14, 1
	global_store_b128 v[3:4], v[11:14], off offset:8
.LBB0_178:
	s_or_b32 exec_lo, exec_lo, s11
	s_lshl_b64 s[6:7], s[6:7], 12
	v_and_b32_e32 v11, 0x3ff, v0
	s_waitcnt vmcnt(0)
	v_add_co_u32 v12, vcc_lo, v7, s6
	v_add_co_ci_u32_e32 v13, vcc_lo, s7, v8, vcc_lo
	v_lshlrev_b64 v[7:8], 6, v[9:10]
	s_and_b32 s6, s16, 0xffff
	s_mov_b32 s20, 0
	v_mad_u64_u32 v[3:4], null, s14, s6, v[11:12]
	v_mov_b32_e32 v4, v10
	s_delay_alu instid0(VALU_DEP_3)
	v_add_co_u32 v7, vcc_lo, v12, v7
	s_mov_b32 s21, s20
	s_mov_b32 s22, s20
	s_mov_b32 s23, s20
	v_and_or_b32 v1, 0xffffff1f, v1, 32
	v_add_co_ci_u32_e32 v8, vcc_lo, v13, v8, vcc_lo
	v_dual_mov_b32 v9, s20 :: v_dual_mov_b32 v10, s21
	v_dual_mov_b32 v11, s22 :: v_dual_mov_b32 v12, s23
	s_clause 0x3
	global_store_b128 v[7:8], v[1:4], off
	global_store_b128 v[7:8], v[9:12], off offset:16
	global_store_b128 v[7:8], v[9:12], off offset:32
	global_store_b128 v[7:8], v[9:12], off offset:48
	s_and_saveexec_b32 s6, s0
	s_cbranch_execz .LBB0_186
; %bb.179:
	v_dual_mov_b32 v11, 0 :: v_dual_mov_b32 v12, s4
	v_mov_b32_e32 v13, s5
	s_clause 0x1
	global_load_b64 v[14:15], v11, s[2:3] offset:32 glc
	global_load_b64 v[1:2], v11, s[2:3] offset:40
	s_waitcnt vmcnt(0)
	v_readfirstlane_b32 s10, v1
	v_readfirstlane_b32 s11, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b64 s[10:11], s[10:11], s[4:5]
	s_mul_i32 s7, s11, 24
	s_mul_hi_u32 s11, s10, 24
	s_mul_i32 s10, s10, 24
	s_add_i32 s11, s11, s7
	v_add_co_u32 v9, vcc_lo, v5, s10
	v_add_co_ci_u32_e32 v10, vcc_lo, s11, v6, vcc_lo
	s_mov_b32 s7, exec_lo
	global_store_b64 v[9:10], v[14:15], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[3:4], v11, v[12:15], s[2:3] offset:32 glc
	s_waitcnt vmcnt(0)
	v_cmpx_ne_u64_e64 v[3:4], v[14:15]
	s_cbranch_execz .LBB0_182
; %bb.180:                              ; %.preheader1.i.i.i21.preheader
	s_mov_b32 s10, 0
.LBB0_181:                              ; %.preheader1.i.i.i21
                                        ; =>This Inner Loop Header: Depth=1
	v_dual_mov_b32 v1, s4 :: v_dual_mov_b32 v2, s5
	s_sleep 1
	global_store_b64 v[9:10], v[3:4], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[1:2], v11, v[1:4], s[2:3] offset:32 glc
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, v[1:2], v[3:4]
	v_dual_mov_b32 v4, v2 :: v_dual_mov_b32 v3, v1
	s_or_b32 s10, vcc_lo, s10
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s10
	s_cbranch_execnz .LBB0_181
.LBB0_182:                              ; %Flow322
	s_or_b32 exec_lo, exec_lo, s7
	v_mov_b32_e32 v4, 0
	s_mov_b32 s10, exec_lo
	s_mov_b32 s7, exec_lo
	v_mbcnt_lo_u32_b32 v3, s10, 0
	global_load_b64 v[1:2], v4, s[2:3] offset:16
	v_cmpx_eq_u32_e32 0, v3
	s_cbranch_execz .LBB0_184
; %bb.183:
	s_bcnt1_i32_b32 s10, s10
	s_delay_alu instid0(SALU_CYCLE_1)
	v_mov_b32_e32 v3, s10
	s_waitcnt vmcnt(0)
	global_atomic_add_u64 v[1:2], v[3:4], off offset:8
.LBB0_184:
	s_or_b32 exec_lo, exec_lo, s7
	s_waitcnt vmcnt(0)
	global_load_b64 v[3:4], v[1:2], off offset:16
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, 0, v[3:4]
	s_cbranch_vccnz .LBB0_186
; %bb.185:
	global_load_b32 v1, v[1:2], off offset:24
	v_mov_b32_e32 v2, 0
	s_waitcnt vmcnt(0)
	v_readfirstlane_b32 s7, v1
	s_waitcnt_vscnt null, 0x0
	global_store_b64 v[3:4], v[1:2], off
	s_and_b32 m0, s7, 0xff
	s_sendmsg sendmsg(MSG_INTERRUPT)
.LBB0_186:                              ; %Flow323
	s_or_b32 exec_lo, exec_lo, s6
	s_add_i32 s8, s8, s1
	v_add_co_u32 v1, vcc_lo, v5, s9
	v_add_co_ci_u32_e32 v2, vcc_lo, s8, v6, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v1, vcc_lo, v1, 20
	v_add_co_ci_u32_e32 v2, vcc_lo, 0, v2, vcc_lo
.LBB0_187:                              ; =>This Inner Loop Header: Depth=1
	v_mov_b32_e32 v3, 1
	s_and_saveexec_b32 s1, s0
	s_cbranch_execz .LBB0_189
; %bb.188:                              ;   in Loop: Header=BB0_187 Depth=1
	global_load_b32 v3, v[1:2], off glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	v_and_b32_e32 v3, 1, v3
.LBB0_189:                              ;   in Loop: Header=BB0_187 Depth=1
	s_or_b32 exec_lo, exec_lo, s1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_readfirstlane_b32 s1, v3
	s_cmp_eq_u32 s1, 0
	s_cbranch_scc1 .LBB0_191
; %bb.190:                              ;   in Loop: Header=BB0_187 Depth=1
	s_mov_b32 s1, 0
	s_sleep 1
	s_branch .LBB0_192
.LBB0_191:                              ;   in Loop: Header=BB0_187 Depth=1
	s_mov_b32 s1, -1
.LBB0_192:                              ; %Flow317
                                        ;   in Loop: Header=BB0_187 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s1
	s_cbranch_vccnz .LBB0_187
; %bb.193:
	global_load_b64 v[1:2], v[7:8], off
	s_and_saveexec_b32 s1, s0
	s_cbranch_execz .LBB0_197
; %bb.194:
	v_mov_b32_e32 v9, 0
	s_clause 0x2
	global_load_b64 v[5:6], v9, s[2:3] offset:40
	global_load_b64 v[10:11], v9, s[2:3] offset:24 glc
	global_load_b64 v[7:8], v9, s[2:3]
	s_waitcnt vmcnt(2)
	v_add_co_u32 v12, vcc_lo, v5, 1
	v_add_co_ci_u32_e32 v13, vcc_lo, 0, v6, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v3, vcc_lo, v12, s4
	v_add_co_ci_u32_e32 v4, vcc_lo, s5, v13, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cmp_eq_u64_e32 vcc_lo, 0, v[3:4]
	v_dual_cndmask_b32 v4, v4, v13 :: v_dual_cndmask_b32 v3, v3, v12
	v_and_b32_e32 v6, v4, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_and_b32_e32 v5, v3, v5
	v_mul_lo_u32 v6, v6, 24
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_hi_u32 v12, v5, 24
	v_mul_lo_u32 v5, v5, 24
	v_add_nc_u32_e32 v6, v12, v6
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_co_u32 v7, vcc_lo, v7, v5
	v_mov_b32_e32 v5, v10
	v_add_co_ci_u32_e32 v8, vcc_lo, v8, v6, vcc_lo
	v_mov_b32_e32 v6, v11
	global_store_b64 v[7:8], v[10:11], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[5:6], v9, v[3:6], s[2:3] offset:24 glc
	s_waitcnt vmcnt(0)
	v_cmp_ne_u64_e32 vcc_lo, v[5:6], v[10:11]
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_197
; %bb.195:                              ; %.preheader.i.i.i20.preheader
	s_mov_b32 s0, 0
.LBB0_196:                              ; %.preheader.i.i.i20
                                        ; =>This Inner Loop Header: Depth=1
	s_sleep 1
	global_store_b64 v[7:8], v[5:6], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[10:11], v9, v[3:6], s[2:3] offset:24 glc
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, v[10:11], v[5:6]
	v_dual_mov_b32 v5, v10 :: v_dual_mov_b32 v6, v11
	s_or_b32 s0, vcc_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s0
	s_cbranch_execnz .LBB0_196
.LBB0_197:                              ; %__ockl_printf_append_args.exit24
	s_or_b32 exec_lo, exec_lo, s1
	;;#ASMSTART
	;;#ASMEND
	v_readfirstlane_b32 s0, v29
	v_mov_b32_e32 v3, 0
	v_mov_b32_e32 v4, 0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_eq_u32_e64 s0, s0, v29
	s_and_saveexec_b32 s1, s0
	s_cbranch_execz .LBB0_203
; %bb.198:
	v_mov_b32_e32 v5, 0
	s_mov_b32 s4, exec_lo
	global_load_b64 v[8:9], v5, s[2:3] offset:24 glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	s_clause 0x1
	global_load_b64 v[3:4], v5, s[2:3] offset:40
	global_load_b64 v[6:7], v5, s[2:3]
	s_waitcnt vmcnt(1)
	v_and_b32_e32 v4, v4, v9
	v_and_b32_e32 v3, v3, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_lo_u32 v4, v4, 24
	v_mul_hi_u32 v10, v3, 24
	v_mul_lo_u32 v3, v3, 24
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v4, v10, v4
	s_waitcnt vmcnt(0)
	v_add_co_u32 v3, vcc_lo, v6, v3
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v4, vcc_lo, v7, v4, vcc_lo
	global_load_b64 v[6:7], v[3:4], off glc
	s_waitcnt vmcnt(0)
	global_atomic_cmpswap_b64 v[3:4], v5, v[6:9], s[2:3] offset:24 glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	v_cmpx_ne_u64_e64 v[3:4], v[8:9]
	s_cbranch_execz .LBB0_202
; %bb.199:                              ; %.preheader3.i.i.i31.preheader
	s_mov_b32 s5, 0
.LBB0_200:                              ; %.preheader3.i.i.i31
                                        ; =>This Inner Loop Header: Depth=1
	s_sleep 1
	s_clause 0x1
	global_load_b64 v[6:7], v5, s[2:3] offset:40
	global_load_b64 v[10:11], v5, s[2:3]
	v_dual_mov_b32 v9, v4 :: v_dual_mov_b32 v8, v3
	s_waitcnt vmcnt(1)
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_and_b32_e32 v6, v6, v8
	s_waitcnt vmcnt(0)
	v_mad_u64_u32 v[3:4], null, v6, 24, v[10:11]
	v_and_b32_e32 v10, v7, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[6:7], null, v10, 24, v[4:5]
	v_mov_b32_e32 v4, v6
	global_load_b64 v[6:7], v[3:4], off glc
	s_waitcnt vmcnt(0)
	global_atomic_cmpswap_b64 v[3:4], v5, v[6:9], s[2:3] offset:24 glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	v_cmp_eq_u64_e32 vcc_lo, v[3:4], v[8:9]
	s_or_b32 s5, vcc_lo, s5
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s5
	s_cbranch_execnz .LBB0_200
; %bb.201:                              ; %Flow310
	s_or_b32 exec_lo, exec_lo, s5
.LBB0_202:                              ; %Flow312
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s4
.LBB0_203:                              ; %.loopexit4.i.i.i25
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s1
	v_mov_b32_e32 v30, 0
	v_readfirstlane_b32 s4, v3
	v_readfirstlane_b32 s5, v4
	s_mov_b32 s10, exec_lo
	s_clause 0x1
	global_load_b64 v[9:10], v30, s[2:3] offset:40
	global_load_b128 v[5:8], v30, s[2:3]
	s_waitcnt vmcnt(1)
	v_readfirstlane_b32 s6, v9
	v_readfirstlane_b32 s7, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b64 s[6:7], s[4:5], s[6:7]
	s_mul_i32 s1, s7, 24
	s_mul_hi_u32 s8, s6, 24
	s_mul_i32 s9, s6, 24
	s_and_saveexec_b32 s11, s0
	s_cbranch_execz .LBB0_205
; %bb.204:
	v_dual_mov_b32 v9, s10 :: v_dual_mov_b32 v10, v30
	s_add_i32 s10, s8, s1
	s_waitcnt vmcnt(0)
	v_add_co_u32 v3, vcc_lo, v5, s9
	v_add_co_ci_u32_e32 v4, vcc_lo, s10, v6, vcc_lo
	v_dual_mov_b32 v11, 2 :: v_dual_mov_b32 v12, 1
	global_store_b128 v[3:4], v[9:12], off offset:8
.LBB0_205:
	s_or_b32 exec_lo, exec_lo, s11
	s_lshl_b64 s[6:7], s[6:7], 12
	v_bfe_u32 v0, v0, 10, 10
	s_waitcnt vmcnt(0)
	v_add_co_u32 v9, vcc_lo, v7, s6
	v_add_co_ci_u32_e32 v10, vcc_lo, s7, v8, vcc_lo
	v_lshlrev_b64 v[7:8], 6, v[29:30]
	s_lshr_b32 s6, s16, 16
	s_mov_b32 s12, 0
	v_mad_u64_u32 v[3:4], null, s15, s6, v[0:1]
	s_mov_b32 s13, s12
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_u32 v11, vcc_lo, v9, v7
	s_mov_b32 s14, s12
	s_mov_b32 s15, s12
	v_and_or_b32 v1, 0xffffff1d, v1, 34
	v_add_co_ci_u32_e32 v12, vcc_lo, v10, v8, vcc_lo
	v_dual_mov_b32 v4, v30 :: v_dual_mov_b32 v7, s12
	v_dual_mov_b32 v8, s13 :: v_dual_mov_b32 v9, s14
	v_mov_b32_e32 v10, s15
	s_clause 0x3
	global_store_b128 v[11:12], v[1:4], off
	global_store_b128 v[11:12], v[7:10], off offset:16
	global_store_b128 v[11:12], v[7:10], off offset:32
	global_store_b128 v[11:12], v[7:10], off offset:48
	s_and_saveexec_b32 s6, s0
	s_cbranch_execz .LBB0_213
; %bb.206:
	v_dual_mov_b32 v4, 0 :: v_dual_mov_b32 v9, s4
	v_mov_b32_e32 v10, s5
	s_clause 0x1
	global_load_b64 v[11:12], v4, s[2:3] offset:32 glc
	global_load_b64 v[0:1], v4, s[2:3] offset:40
	s_waitcnt vmcnt(0)
	v_readfirstlane_b32 s10, v0
	v_readfirstlane_b32 s11, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b64 s[10:11], s[10:11], s[4:5]
	s_mul_i32 s7, s11, 24
	s_mul_hi_u32 s11, s10, 24
	s_mul_i32 s10, s10, 24
	s_add_i32 s11, s11, s7
	v_add_co_u32 v7, vcc_lo, v5, s10
	v_add_co_ci_u32_e32 v8, vcc_lo, s11, v6, vcc_lo
	s_mov_b32 s7, exec_lo
	global_store_b64 v[7:8], v[11:12], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[2:3], v4, v[9:12], s[2:3] offset:32 glc
	s_waitcnt vmcnt(0)
	v_cmpx_ne_u64_e64 v[2:3], v[11:12]
	s_cbranch_execz .LBB0_209
; %bb.207:                              ; %.preheader1.i.i.i29.preheader
	s_mov_b32 s10, 0
.LBB0_208:                              ; %.preheader1.i.i.i29
                                        ; =>This Inner Loop Header: Depth=1
	v_dual_mov_b32 v0, s4 :: v_dual_mov_b32 v1, s5
	s_sleep 1
	global_store_b64 v[7:8], v[2:3], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[0:1], v4, v[0:3], s[2:3] offset:32 glc
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, v[0:1], v[2:3]
	v_dual_mov_b32 v3, v1 :: v_dual_mov_b32 v2, v0
	s_or_b32 s10, vcc_lo, s10
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s10
	s_cbranch_execnz .LBB0_208
.LBB0_209:                              ; %Flow308
	s_or_b32 exec_lo, exec_lo, s7
	v_mov_b32_e32 v3, 0
	s_mov_b32 s10, exec_lo
	s_mov_b32 s7, exec_lo
	v_mbcnt_lo_u32_b32 v2, s10, 0
	global_load_b64 v[0:1], v3, s[2:3] offset:16
	v_cmpx_eq_u32_e32 0, v2
	s_cbranch_execz .LBB0_211
; %bb.210:
	s_bcnt1_i32_b32 s10, s10
	s_delay_alu instid0(SALU_CYCLE_1)
	v_mov_b32_e32 v2, s10
	s_waitcnt vmcnt(0)
	global_atomic_add_u64 v[0:1], v[2:3], off offset:8
.LBB0_211:
	s_or_b32 exec_lo, exec_lo, s7
	s_waitcnt vmcnt(0)
	global_load_b64 v[2:3], v[0:1], off offset:16
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, 0, v[2:3]
	s_cbranch_vccnz .LBB0_213
; %bb.212:
	global_load_b32 v0, v[0:1], off offset:24
	v_mov_b32_e32 v1, 0
	s_waitcnt vmcnt(0)
	v_readfirstlane_b32 s7, v0
	s_waitcnt_vscnt null, 0x0
	global_store_b64 v[2:3], v[0:1], off
	s_and_b32 m0, s7, 0xff
	s_sendmsg sendmsg(MSG_INTERRUPT)
.LBB0_213:                              ; %Flow309
	s_or_b32 exec_lo, exec_lo, s6
	s_add_i32 s8, s8, s1
	v_add_co_u32 v0, vcc_lo, v5, s9
	v_add_co_ci_u32_e32 v1, vcc_lo, s8, v6, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v0, vcc_lo, v0, 20
	v_add_co_ci_u32_e32 v1, vcc_lo, 0, v1, vcc_lo
.LBB0_214:                              ; =>This Inner Loop Header: Depth=1
	v_mov_b32_e32 v2, 1
	s_and_saveexec_b32 s1, s0
	s_cbranch_execz .LBB0_216
; %bb.215:                              ;   in Loop: Header=BB0_214 Depth=1
	global_load_b32 v2, v[0:1], off glc
	s_waitcnt vmcnt(0)
	buffer_gl1_inv
	buffer_gl0_inv
	v_and_b32_e32 v2, 1, v2
.LBB0_216:                              ;   in Loop: Header=BB0_214 Depth=1
	s_or_b32 exec_lo, exec_lo, s1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_readfirstlane_b32 s1, v2
	s_cmp_eq_u32 s1, 0
	s_cbranch_scc1 .LBB0_218
; %bb.217:                              ;   in Loop: Header=BB0_214 Depth=1
	s_mov_b32 s1, 0
	s_sleep 1
	s_branch .LBB0_219
.LBB0_218:                              ;   in Loop: Header=BB0_214 Depth=1
	s_mov_b32 s1, -1
.LBB0_219:                              ; %Flow303
                                        ;   in Loop: Header=BB0_214 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s1
	s_cbranch_vccnz .LBB0_214
; %bb.220:
	s_and_saveexec_b32 s1, s0
	s_cbranch_execz .LBB0_224
; %bb.221:
	v_mov_b32_e32 v6, 0
	s_clause 0x2
	global_load_b64 v[2:3], v6, s[2:3] offset:40
	global_load_b64 v[7:8], v6, s[2:3] offset:24 glc
	global_load_b64 v[4:5], v6, s[2:3]
	s_waitcnt vmcnt(2)
	v_add_co_u32 v9, vcc_lo, v2, 1
	v_add_co_ci_u32_e32 v10, vcc_lo, 0, v3, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v0, vcc_lo, v9, s4
	v_add_co_ci_u32_e32 v1, vcc_lo, s5, v10, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cmp_eq_u64_e32 vcc_lo, 0, v[0:1]
	v_dual_cndmask_b32 v1, v1, v10 :: v_dual_cndmask_b32 v0, v0, v9
	v_and_b32_e32 v3, v1, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_and_b32_e32 v2, v0, v2
	v_mul_lo_u32 v3, v3, 24
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_hi_u32 v9, v2, 24
	v_mul_lo_u32 v2, v2, 24
	v_add_nc_u32_e32 v3, v9, v3
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_co_u32 v4, vcc_lo, v4, v2
	v_mov_b32_e32 v2, v7
	v_add_co_ci_u32_e32 v5, vcc_lo, v5, v3, vcc_lo
	v_mov_b32_e32 v3, v8
	global_store_b64 v[4:5], v[7:8], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[2:3], v6, v[0:3], s[2:3] offset:24 glc
	s_waitcnt vmcnt(0)
	v_cmp_ne_u64_e32 vcc_lo, v[2:3], v[7:8]
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_224
; %bb.222:                              ; %.preheader.i.i.i28.preheader
	s_mov_b32 s0, 0
.LBB0_223:                              ; %.preheader.i.i.i28
                                        ; =>This Inner Loop Header: Depth=1
	s_sleep 1
	global_store_b64 v[4:5], v[2:3], off
	s_waitcnt_vscnt null, 0x0
	global_atomic_cmpswap_b64 v[7:8], v6, v[0:3], s[2:3] offset:24 glc
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, v[7:8], v[2:3]
	v_dual_mov_b32 v2, v7 :: v_dual_mov_b32 v3, v8
	s_or_b32 s0, vcc_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s0
	s_cbranch_execnz .LBB0_223
.LBB0_224:                              ; %__ockl_printf_append_args.exit32
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z6kernelv
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 256
		.amdhsa_user_sgpr_count 14
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 1
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 1
		.amdhsa_next_free_vgpr 33
		.amdhsa_next_free_sgpr 24
		.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_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 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_end0:
	.size	_Z6kernelv, .Lfunc_end0-_Z6kernelv
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 9440
; NumSgprs: 26
; NumVgprs: 33
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 3
; VGPRBlocks: 4
; NumSGPRsForWavesPerEU: 26
; NumVGPRsForWavesPerEU: 33
; Occupancy: 16
; WaveLimiterHint : 1
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; 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: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 1
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	.str,@object                    ; @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.str:
	.asciz	"Hello, World! size=%d   tidx=%d, tidy=%d\n"
	.size	.str, 42

	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .offset:         0
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         4
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         8
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         12
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         14
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         16
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         18
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         20
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         22
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         40
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         48
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         56
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         64
        .size:           2
        .value_kind:     hidden_grid_dims
      - .offset:         80
        .size:           8
        .value_kind:     hidden_hostcall_buffer
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 256
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z6kernelv
    .private_segment_fixed_size: 0
    .sgpr_count:     26
    .sgpr_spill_count: 0
    .symbol:         _Z6kernelv.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     33
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
