/*
Simple non visual example to attempt creating a makefile for cuda
This example will increment the contents of an array by 1
This program accepts command line argument to set the 
size of the array and the block size otherwise it uses default values.
No validation checks on these inputs for non integer args or 
if they exceed the maximum block size for device
*/

#include <stdio.h>
#include <assert.h>
#include <cuda.h>
#include <sstream>

// increment array on Host
void hostIncrement(int* a, int N)
{
	for(int i=0;i<N;i++)
	{
		a[i]=a[i]+1;
	}

	return;
}

// increment array on Device
__global__ void deviceIncrement(int* a, int N)
{
	int idx = blockIdx.x*blockDim.x + threadIdx.x;
	
	if(idx <N)
	{
		a[idx] = a[idx]+1;
	}
}

// main function
int main(int argc, char** argv)
{
	int N;
	int blockSize;
	// Check command line arguments
	if(argc == 1)
	{
		// none given, use default;
		N=16;
		blockSize=4;
	}
	else
	{
		std::istringstream(argv[1]) >> N;	
		std::istringstream(argv[2]) >> blockSize;
	}

	// calculate size of arrays
	size_t size=N*sizeof(int);

	// Allocate memory for arrays on host	
	int* a_h = (int*)malloc(size);
	int* b_h = (int*)malloc(size);

	// Allocate memory on device
	int* a_d;	
	cudaMalloc((void**) &a_d, size);

	// Initialise array data
	for(int i=0;i<N;i++)
	{
		a_h[i] = i;
	}

	// Copy data from host to device
	cudaMemcpy(a_d, a_h, sizeof(int)*N, cudaMemcpyHostToDevice);
	
	// Do calculation on host
	hostIncrement(a_h, N);

	// Do calculation on device
	int nBlocks = N/blockSize + (N%blockSize == 0?0:1);
	
	deviceIncrement <<< nBlocks, blockSize >>> (a_d, N);

	// Retrieve results from device and store in b_h
	cudaMemcpy(b_h, a_d, sizeof(int)*N, cudaMemcpyDeviceToHost);

	for(int i=0;i<N;i++)
	{
		assert(a_h[i] == b_h[i]);
		printf("host - %d, device - %d\n", a_h[i], b_h[i]);
	}

	// Release memory
	free(a_h);
	free(b_h);
	cudaFree(a_d);
	
	return 0;
}

	.file	"tmpxft_0038958e_00000000-6_incrementArray.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB3697:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE3697:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z13hostIncrementPii
	.type	_Z13hostIncrementPii, @function
_Z13hostIncrementPii:
.LFB3693:
	.cfi_startproc
	endbr64
	xorl	%eax, %eax
.L3:
	cmpl	%eax, %esi
	jle	.L6
	incl	(%rdi,%rax,4)
	incq	%rax
	jmp	.L3
.L6:
	ret
	.cfi_endproc
.LFE3693:
	.size	_Z13hostIncrementPii, .-_Z13hostIncrementPii
	.globl	_Z36__device_stub__Z15deviceIncrementPiiPii
	.type	_Z36__device_stub__Z15deviceIncrementPiiPii, @function
_Z36__device_stub__Z15deviceIncrementPiiPii:
.LFB3719:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 8(%rsp)
	leaq	32(%rsp), %rcx
	leaq	24(%rsp), %rdx
	movl	%esi, 4(%rsp)
	leaq	40(%rsp), %rdi
	leaq	52(%rsp), %rsi
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 48(%rsp)
	movq	%rax, 88(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 40(%rsp)
	movq	%rax, 52(%rsp)
	movl	$1, 60(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L7
	pushq	32(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z15deviceIncrementPii(%rip), %rdi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 144
	movq	68(%rsp), %rcx
	movl	76(%rsp), %r8d
	movq	56(%rsp), %rsi
	movl	64(%rsp), %edx
	leaq	104(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L7:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L9
	call	__stack_chk_fail@PLT
.L9:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3719:
	.size	_Z36__device_stub__Z15deviceIncrementPiiPii, .-_Z36__device_stub__Z15deviceIncrementPiiPii
	.globl	_Z15deviceIncrementPii
	.type	_Z15deviceIncrementPii, @function
_Z15deviceIncrementPii:
.LFB3720:
	.cfi_startproc
	endbr64
	jmp	_Z36__device_stub__Z15deviceIncrementPiiPii
	.cfi_endproc
.LFE3720:
	.size	_Z15deviceIncrementPii, .-_Z15deviceIncrementPii
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"_Z15deviceIncrementPii"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB3722:
	.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	.LC0(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z15deviceIncrementPii(%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
.LFE3722:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.rodata._ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_.str1.1,"aMS",@progbits,1
.LC1:
	.string	"basic_string: construction from null is not valid"
	.section	.text._ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_,"axG",@progbits,_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC5IS3_EEPKcRKS3_,comdat
	.align 2
	.weak	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_
	.type	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_, @function
_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_:
.LFB4023:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	subq	$16, %rsp
	.cfi_def_cfa_offset 48
	movq	%fs:40, %rax
	movq	%rax, 8(%rsp)
	xorl	%eax, %eax
	leaq	16(%rdi), %rax
	movq	%rax, (%rdi)
	testq	%rsi, %rsi
	jne	.L15
	movq	8(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L20
	leaq	.LC1(%rip), %rdi
	call	_ZSt19__throw_logic_errorPKc@PLT
.L15:
	movq	%rdi, %rbx
	movq	%rsi, %rdi
	movq	%rsi, %rbp
	call	strlen@PLT
	movq	%rax, (%rsp)
	leaq	0(%rbp,%rax), %r12
	cmpq	$15, %rax
	jbe	.L17
	movq	%rsp, %rsi
	xorl	%edx, %edx
	movq	%rbx, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE9_M_createERmm@PLT
	movq	%rax, (%rbx)
	movq	(%rsp), %rax
	movq	%rax, 16(%rbx)
.L17:
	movq	(%rbx), %rdi
	movq	%r12, %rdx
	movq	%rbp, %rsi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE13_S_copy_charsEPcPKcS7_@PLT
	movq	(%rsp), %rax
	movq	(%rbx), %rdx
	movq	%rax, 8(%rbx)
	movb	$0, (%rdx,%rax)
	movq	8(%rsp), %rax
	subq	%fs:40, %rax
	je	.L18
.L20:
	call	__stack_chk_fail@PLT
.L18:
	addq	$16, %rsp
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4023:
	.size	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_, .-_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_
	.weak	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC1IS3_EEPKcRKS3_
	.set	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC1IS3_EEPKcRKS3_,_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_
	.section	.rodata.str1.1
.LC2:
	.string	"host - %d, device - %d\n"
	.section	.text.startup
	.globl	main
	.type	main, @function
main:
.LFB3694:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA3694
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	subq	$472, %rsp
	.cfi_def_cfa_offset 512
	movq	%fs:40, %rax
	movq	%rax, 456(%rsp)
	xorl	%eax, %eax
	decl	%edi
	jne	.L22
	movl	$16, (%rsp)
	movl	$4, 4(%rsp)
	jmp	.L23
.L22:
	leaq	40(%rsp), %rbp
	movq	%rsi, %rbx
	leaq	28(%rsp), %r13
	movq	8(%rsi), %rsi
	movq	%r13, %rdx
	movq	%rbp, %rdi
	leaq	72(%rsp), %r12
.LEHB0:
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC1IS3_EEPKcRKS3_
.LEHE0:
	movl	$8, %edx
	movq	%rbp, %rsi
	movq	%r12, %rdi
.LEHB1:
	call	_ZNSt7__cxx1119basic_istringstreamIcSt11char_traitsIcESaIcEEC1ERKNS_12basic_stringIcS2_S3_EESt13_Ios_Openmode@PLT
.LEHE1:
	movq	%rsp, %rsi
	movq	%r12, %rdi
.LEHB2:
	call	_ZNSirsERi@PLT
.LEHE2:
	movq	%r12, %rdi
	call	_ZNSt7__cxx1119basic_istringstreamIcSt11char_traitsIcESaIcEED1Ev@PLT
	movq	%rbp, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	movq	16(%rbx), %rsi
	movq	%r13, %rdx
	movq	%rbp, %rdi
.LEHB3:
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC1IS3_EEPKcRKS3_
.LEHE3:
	movl	$8, %edx
	movq	%rbp, %rsi
	movq	%r12, %rdi
.LEHB4:
	call	_ZNSt7__cxx1119basic_istringstreamIcSt11char_traitsIcESaIcEEC1ERKNS_12basic_stringIcS2_S3_EESt13_Ios_Openmode@PLT
.LEHE4:
	leaq	4(%rsp), %rsi
	movq	%r12, %rdi
.LEHB5:
	call	_ZNSirsERi@PLT
.LEHE5:
	movq	%r12, %rdi
	call	_ZNSt7__cxx1119basic_istringstreamIcSt11char_traitsIcESaIcEED1Ev@PLT
	movq	%rbp, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
.L23:
	movslq	(%rsp), %r12
	salq	$2, %r12
	movq	%r12, %rdi
	call	malloc@PLT
	movq	%r12, %rdi
	movq	%rax, %rbx
	call	malloc@PLT
	leaq	8(%rsp), %rdi
	movq	%r12, %rsi
	movq	%rax, %rbp
.LEHB6:
	call	cudaMalloc@PLT
	movslq	(%rsp), %rdx
	xorl	%eax, %eax
.L24:
	cmpl	%eax, %edx
	jle	.L45
	movl	%eax, (%rbx,%rax,4)
	incq	%rax
	jmp	.L24
.L45:
	movq	8(%rsp), %rdi
	salq	$2, %rdx
	movl	$1, %ecx
	movq	%rbx, %rsi
	call	cudaMemcpy@PLT
	movl	(%rsp), %edx
	movq	%rbx, %rdi
	movl	%edx, %esi
	call	_Z13hostIncrementPii
	movl	4(%rsp), %ecx
	movl	%edx, %eax
	cltd
	idivl	%ecx
	movl	%ecx, 28(%rsp)
	movabsq	$4294967297, %rcx
	movq	%rcx, 20(%rsp)
	movl	24(%rsp), %esi
	movq	%rcx, 32(%rsp)
	movl	36(%rsp), %ecx
	cmpl	$1, %edx
	movq	28(%rsp), %rdx
	sbbl	$-1, %eax
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	%eax, 16(%rsp)
	movq	16(%rsp), %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L26
	movl	(%rsp), %esi
	movq	8(%rsp), %rdi
	call	_Z36__device_stub__Z15deviceIncrementPiiPii
.L26:
	movslq	(%rsp), %rdx
	movq	8(%rsp), %rsi
	movq	%rbp, %rdi
	xorl	%r12d, %r12d
	movl	$2, %ecx
	leaq	.LC2(%rip), %r13
	salq	$2, %rdx
	call	cudaMemcpy@PLT
.L27:
	cmpl	%r12d, (%rsp)
	jle	.L46
	movl	0(%rbp,%r12,4), %ecx
	movl	(%rbx,%r12,4), %edx
	movq	%r13, %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	incq	%r12
	call	__printf_chk@PLT
	jmp	.L27
.L46:
	movq	%rbx, %rdi
	call	free@PLT
	movq	%rbp, %rdi
	call	free@PLT
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	movq	456(%rsp), %rax
	subq	%fs:40, %rax
	je	.L35
	jmp	.L40
.L37:
	endbr64
	jmp	.L44
.L36:
	endbr64
	jmp	.L43
.L39:
	endbr64
.L44:
	movq	%rax, %rbx
.L32:
	movq	%r12, %rdi
	call	_ZNSt7__cxx1119basic_istringstreamIcSt11char_traitsIcESaIcEED1Ev@PLT
	jmp	.L33
.L38:
	endbr64
.L43:
	movq	%rax, %rbx
.L33:
	movq	%rbp, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	movq	456(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L40
	movq	%rbx, %rdi
	call	_Unwind_Resume@PLT
.LEHE6:
.L40:
	call	__stack_chk_fail@PLT
.L35:
	addq	$472, %rsp
	.cfi_def_cfa_offset 40
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3694:
	.globl	__gxx_personality_v0
	.section	.gcc_except_table,"a",@progbits
.LLSDA3694:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE3694-.LLSDACSB3694
.LLSDACSB3694:
	.uleb128 .LEHB0-.LFB3694
	.uleb128 .LEHE0-.LEHB0
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB1-.LFB3694
	.uleb128 .LEHE1-.LEHB1
	.uleb128 .L36-.LFB3694
	.uleb128 0
	.uleb128 .LEHB2-.LFB3694
	.uleb128 .LEHE2-.LEHB2
	.uleb128 .L37-.LFB3694
	.uleb128 0
	.uleb128 .LEHB3-.LFB3694
	.uleb128 .LEHE3-.LEHB3
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB4-.LFB3694
	.uleb128 .LEHE4-.LEHB4
	.uleb128 .L38-.LFB3694
	.uleb128 0
	.uleb128 .LEHB5-.LFB3694
	.uleb128 .LEHE5-.LEHB5
	.uleb128 .L39-.LFB3694
	.uleb128 0
	.uleb128 .LEHB6-.LFB3694
	.uleb128 .LEHE6-.LEHB6
	.uleb128 0
	.uleb128 0
.LLSDACSE3694:
	.section	.text.startup
	.size	main, .-main
	.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
	.hidden	DW.ref.__gxx_personality_v0
	.weak	DW.ref.__gxx_personality_v0
	.section	.data.rel.local.DW.ref.__gxx_personality_v0,"awG",@progbits,DW.ref.__gxx_personality_v0,comdat
	.align 8
	.type	DW.ref.__gxx_personality_v0, @object
	.size	DW.ref.__gxx_personality_v0, 8
DW.ref.__gxx_personality_v0:
	.quad	__gxx_personality_v0
	.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:

/*
Simple non visual example to attempt creating a makefile for cuda
This example will increment the contents of an array by 1
This program accepts command line argument to set the 
size of the array and the block size otherwise it uses default values.
No validation checks on these inputs for non integer args or 
if they exceed the maximum block size for device
*/

#include <stdio.h>
#include <assert.h>
#include <hip/hip_runtime.h>
#include <sstream>

// increment array on Host
void hostIncrement(int* a, int N)
{
	for(int i=0;i<N;i++)
	{
		a[i]=a[i]+1;
	}

	return;
}

// increment array on Device
__global__ void deviceIncrement(int* a, int N)
{
	int idx = blockIdx.x*blockDim.x + threadIdx.x;
	
	if(idx <N)
	{
		a[idx] = a[idx]+1;
	}
}

// main function
int main(int argc, char** argv)
{
	int N;
	int blockSize;
	// Check command line arguments
	if(argc == 1)
	{
		// none given, use default;
		N=16;
		blockSize=4;
	}
	else
	{
		std::istringstream(argv[1]) >> N;	
		std::istringstream(argv[2]) >> blockSize;
	}

	// calculate size of arrays
	size_t size=N*sizeof(int);

	// Allocate memory for arrays on host	
	int* a_h = (int*)malloc(size);
	int* b_h = (int*)malloc(size);

	// Allocate memory on device
	int* a_d;	
	hipMalloc((void**) &a_d, size);

	// Initialise array data
	for(int i=0;i<N;i++)
	{
		a_h[i] = i;
	}

	// Copy data from host to device
	hipMemcpy(a_d, a_h, sizeof(int)*N, hipMemcpyHostToDevice);
	
	// Do calculation on host
	hostIncrement(a_h, N);

	// Do calculation on device
	int nBlocks = N/blockSize + (N%blockSize == 0?0:1);
	
	deviceIncrement <<< nBlocks, blockSize >>> (a_d, N);

	// Retrieve results from device and store in b_h
	hipMemcpy(b_h, a_d, sizeof(int)*N, hipMemcpyDeviceToHost);

	for(int i=0;i<N;i++)
	{
		assert(a_h[i] == b_h[i]);
		printf("host - %d, device - %d\n", a_h[i], b_h[i]);
	}

	// Release memory
	free(a_h);
	free(b_h);
	hipFree(a_d);
	
	return 0;
}

	.text
	.file	"incrementArray.hip"
	.globl	_Z13hostIncrementPii            # -- Begin function _Z13hostIncrementPii
	.type	_Z13hostIncrementPii,@function
_Z13hostIncrementPii:                   # @_Z13hostIncrementPii
	.cfi_startproc
# %bb.0:
	testl	%esi, %esi
	jle	.LBB0_3
# %bb.1:                                # %.lr.ph.preheader
	movl	%esi, %eax
	xorl	%ecx, %ecx
.LBB0_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	incl	(%rdi,%rcx,4)
	incq	%rcx
	cmpq	%rcx, %rax
	jne	.LBB0_2
.LBB0_3:                                # %._crit_edge
	retq
.Lfunc_end0:
	.size	_Z13hostIncrementPii, .Lfunc_end0-_Z13hostIncrementPii
	.cfi_endproc
                                        # -- End function
	.globl	_Z30__device_stub__deviceIncrementPii # -- Begin function _Z30__device_stub__deviceIncrementPii
	.type	_Z30__device_stub__deviceIncrementPii,@function
_Z30__device_stub__deviceIncrementPii:  # @_Z30__device_stub__deviceIncrementPii
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$80, %rsp
	.cfi_def_cfa_offset 128
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	24(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	4(%rsp), %rcx
	movl	%esi, (%rcx)
	leaq	64(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z15deviceIncrementPii, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$96, %rsp
	.cfi_adjust_cfa_offset -96
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	_Z30__device_stub__deviceIncrementPii, .Lfunc_end1-_Z30__device_stub__deviceIncrementPii
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
.Lfunc_begin0:
	.cfi_startproc
	.cfi_personality 3, __gxx_personality_v0
	.cfi_lsda 3, .Lexception0
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	subq	$432, %rsp                      # imm = 0x1B0
	.cfi_def_cfa_offset 464
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	cmpl	$1, %edi
	jne	.LBB2_2
# %bb.1:
	movl	$16, 4(%rsp)
	movl	$4, 12(%rsp)
	movl	$64, %r15d
	jmp	.LBB2_13
.LBB2_2:
	movq	%rsi, %rbx
	movq	8(%rsi), %rsi
.Ltmp0:
	leaq	16(%rsp), %rdi
	leaq	11(%rsp), %rdx
	callq	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_
.Ltmp1:
# %bb.3:
.Ltmp3:
	leaq	48(%rsp), %rdi
	leaq	16(%rsp), %rsi
	movl	$8, %edx
	callq	_ZNSt7__cxx1119basic_istringstreamIcSt11char_traitsIcESaIcEEC1ERKNS_12basic_stringIcS2_S3_EESt13_Ios_Openmode
.Ltmp4:
# %bb.4:
.Ltmp6:
	leaq	48(%rsp), %rdi
	leaq	4(%rsp), %rsi
	callq	_ZNSirsERi
.Ltmp7:
# %bb.5:                                # %_ZStrsINSt7__cxx1119basic_istringstreamIcSt11char_traitsIcESaIcEEERiEOT_S8_OT0_.exit
	leaq	48(%rsp), %rdi
	movl	$_ZTTNSt7__cxx1119basic_istringstreamIcSt11char_traitsIcESaIcEEE, %esi
	callq	_ZNSt7__cxx1119basic_istringstreamIcSt11char_traitsIcESaIcEED2Ev
	leaq	168(%rsp), %rdi
	callq	_ZNSt8ios_baseD2Ev
	leaq	32(%rsp), %r14
	movq	-16(%r14), %rdi
	cmpq	%r14, %rdi
	je	.LBB2_7
# %bb.6:                                # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i
	callq	_ZdlPv
.LBB2_7:                                # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev.exit
	movq	16(%rbx), %rsi
.Ltmp9:
	leaq	16(%rsp), %rdi
	leaq	11(%rsp), %rdx
	callq	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_
.Ltmp10:
# %bb.8:
.Ltmp12:
	leaq	48(%rsp), %rdi
	leaq	16(%rsp), %rsi
	movl	$8, %edx
	callq	_ZNSt7__cxx1119basic_istringstreamIcSt11char_traitsIcESaIcEEC1ERKNS_12basic_stringIcS2_S3_EESt13_Ios_Openmode
.Ltmp13:
# %bb.9:
.Ltmp15:
	leaq	48(%rsp), %rdi
	leaq	12(%rsp), %rsi
	callq	_ZNSirsERi
.Ltmp16:
# %bb.10:                               # %_ZStrsINSt7__cxx1119basic_istringstreamIcSt11char_traitsIcESaIcEEERiEOT_S8_OT0_.exit40
	leaq	48(%rsp), %rdi
	movl	$_ZTTNSt7__cxx1119basic_istringstreamIcSt11char_traitsIcESaIcEEE, %esi
	callq	_ZNSt7__cxx1119basic_istringstreamIcSt11char_traitsIcESaIcEED2Ev
	leaq	168(%rsp), %rdi
	callq	_ZNSt8ios_baseD2Ev
	movq	16(%rsp), %rdi
	cmpq	%r14, %rdi
	je	.LBB2_12
# %bb.11:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i41
	callq	_ZdlPv
.LBB2_12:                               # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev.exit43
	movslq	4(%rsp), %r15
	shlq	$2, %r15
.LBB2_13:
	movq	%r15, %rdi
	callq	malloc
	movq	%rax, %rbx
	movq	%r15, %rdi
	callq	malloc
	movq	%rax, %r14
	leaq	48(%rsp), %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	movl	4(%rsp), %eax
	testl	%eax, %eax
	jle	.LBB2_16
# %bb.14:                               # %.lr.ph.preheader
	xorl	%ecx, %ecx
.LBB2_15:                               # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movl	%ecx, (%rbx,%rcx,4)
	incq	%rcx
	cmpq	%rcx, %rax
	jne	.LBB2_15
.LBB2_16:                               # %._crit_edge
	movq	48(%rsp), %rdi
	movslq	%eax, %rdx
	shlq	$2, %rdx
	movq	%rbx, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movl	4(%rsp), %eax
	testl	%eax, %eax
	jle	.LBB2_19
# %bb.17:                               # %.lr.ph.preheader.i
	xorl	%ecx, %ecx
.LBB2_18:                               # %.lr.ph.i
                                        # =>This Inner Loop Header: Depth=1
	incl	(%rbx,%rcx,4)
	incq	%rcx
	cmpq	%rcx, %rax
	jne	.LBB2_18
.LBB2_19:                               # %_Z13hostIncrementPii.exit
	movl	12(%rsp), %ecx
                                        # kill: def $eax killed $eax killed $rax
	cltd
	idivl	%ecx
                                        # kill: def $eax killed $eax def $rax
	cmpl	$1, %edx
	sbbl	$-1, %eax
	btsq	$32, %rax
	btsq	$32, %rcx
	movq	%rax, %rdi
	movl	$1, %esi
	movq	%rcx, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB2_21
# %bb.20:
	movq	48(%rsp), %rdi
	movl	4(%rsp), %esi
	callq	_Z30__device_stub__deviceIncrementPii
.LBB2_21:
	movq	48(%rsp), %rsi
	movslq	4(%rsp), %rdx
	shlq	$2, %rdx
	movq	%r14, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	cmpl	$0, 4(%rsp)
	jle	.LBB2_24
# %bb.22:                               # %.lr.ph56.preheader
	xorl	%r15d, %r15d
.LBB2_23:                               # %.lr.ph56
                                        # =>This Inner Loop Header: Depth=1
	movl	(%rbx,%r15,4), %esi
	movl	(%r14,%r15,4), %edx
	movl	$.L.str, %edi
	xorl	%eax, %eax
	callq	printf
	incq	%r15
	movslq	4(%rsp), %rax
	cmpq	%rax, %r15
	jl	.LBB2_23
.LBB2_24:                               # %._crit_edge57
	movq	%rbx, %rdi
	callq	free
	movq	%r14, %rdi
	callq	free
	movq	48(%rsp), %rdi
	callq	hipFree
	xorl	%eax, %eax
	addq	$432, %rsp                      # imm = 0x1B0
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.LBB2_34:
	.cfi_def_cfa_offset 464
.Ltmp17:
	movq	%rax, %rbx
	leaq	48(%rsp), %rdi
	movl	$_ZTTNSt7__cxx1119basic_istringstreamIcSt11char_traitsIcESaIcEEE, %esi
	callq	_ZNSt7__cxx1119basic_istringstreamIcSt11char_traitsIcESaIcEED2Ev
	leaq	168(%rsp), %rdi
	callq	_ZNSt8ios_baseD2Ev
	jmp	.LBB2_35
.LBB2_33:
.Ltmp14:
	movq	%rax, %rbx
.LBB2_35:
	movq	16(%rsp), %rdi
	cmpq	%r14, %rdi
	jne	.LBB2_31
	jmp	.LBB2_27
.LBB2_32:
.Ltmp11:
	jmp	.LBB2_26
.LBB2_29:
.Ltmp8:
	movq	%rax, %rbx
	leaq	48(%rsp), %rdi
	movl	$_ZTTNSt7__cxx1119basic_istringstreamIcSt11char_traitsIcESaIcEEE, %esi
	callq	_ZNSt7__cxx1119basic_istringstreamIcSt11char_traitsIcESaIcEED2Ev
	leaq	168(%rsp), %rdi
	callq	_ZNSt8ios_baseD2Ev
	jmp	.LBB2_30
.LBB2_28:
.Ltmp5:
	movq	%rax, %rbx
.LBB2_30:
	leaq	32(%rsp), %rax
	movq	-16(%rax), %rdi
	cmpq	%rax, %rdi
	je	.LBB2_27
.LBB2_31:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i44
	callq	_ZdlPv
	jmp	.LBB2_27
.LBB2_25:
.Ltmp2:
.LBB2_26:
	movq	%rax, %rbx
.LBB2_27:
	movq	%rbx, %rdi
	callq	_Unwind_Resume@PLT
.Lfunc_end2:
	.size	main, .Lfunc_end2-main
	.cfi_endproc
	.section	.gcc_except_table,"a",@progbits
	.p2align	2, 0x0
GCC_except_table2:
.Lexception0:
	.byte	255                             # @LPStart Encoding = omit
	.byte	255                             # @TType Encoding = omit
	.byte	1                               # Call site Encoding = uleb128
	.uleb128 .Lcst_end0-.Lcst_begin0
.Lcst_begin0:
	.uleb128 .Ltmp0-.Lfunc_begin0           # >> Call Site 1 <<
	.uleb128 .Ltmp1-.Ltmp0                  #   Call between .Ltmp0 and .Ltmp1
	.uleb128 .Ltmp2-.Lfunc_begin0           #     jumps to .Ltmp2
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp3-.Lfunc_begin0           # >> Call Site 2 <<
	.uleb128 .Ltmp4-.Ltmp3                  #   Call between .Ltmp3 and .Ltmp4
	.uleb128 .Ltmp5-.Lfunc_begin0           #     jumps to .Ltmp5
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp6-.Lfunc_begin0           # >> Call Site 3 <<
	.uleb128 .Ltmp7-.Ltmp6                  #   Call between .Ltmp6 and .Ltmp7
	.uleb128 .Ltmp8-.Lfunc_begin0           #     jumps to .Ltmp8
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp9-.Lfunc_begin0           # >> Call Site 4 <<
	.uleb128 .Ltmp10-.Ltmp9                 #   Call between .Ltmp9 and .Ltmp10
	.uleb128 .Ltmp11-.Lfunc_begin0          #     jumps to .Ltmp11
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp12-.Lfunc_begin0          # >> Call Site 5 <<
	.uleb128 .Ltmp13-.Ltmp12                #   Call between .Ltmp12 and .Ltmp13
	.uleb128 .Ltmp14-.Lfunc_begin0          #     jumps to .Ltmp14
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp15-.Lfunc_begin0          # >> Call Site 6 <<
	.uleb128 .Ltmp16-.Ltmp15                #   Call between .Ltmp15 and .Ltmp16
	.uleb128 .Ltmp17-.Lfunc_begin0          #     jumps to .Ltmp17
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp16-.Lfunc_begin0          # >> Call Site 7 <<
	.uleb128 .Lfunc_end2-.Ltmp16            #   Call between .Ltmp16 and .Lfunc_end2
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
.Lcst_end0:
	.p2align	2, 0x0
                                        # -- End function
	.section	.text._ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_,"axG",@progbits,_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_,comdat
	.weak	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_ # -- Begin function _ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_
	.p2align	1, 0x90
	.type	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_,@function
_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_: # @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_
	.cfi_startproc
# %bb.0:
	pushq	%r14
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	pushq	%rax
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -24
	.cfi_offset %r14, -16
	leaq	16(%rdi), %rax
	movq	%rax, (%rdi)
	testq	%rsi, %rsi
	je	.LBB3_1
# %bb.2:
	movq	%rsi, %rbx
	movq	%rdi, %r14
	movq	%rsi, %rdi
	callq	strlen
	leaq	(%rax,%rbx), %rdx
	movq	%r14, %rdi
	movq	%rbx, %rsi
	addq	$8, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	jmp	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructIPKcEEvT_S8_St20forward_iterator_tag # TAILCALL
.LBB3_1:
	.cfi_def_cfa_offset 32
	movl	$.L.str.1, %edi
	callq	_ZSt19__throw_logic_errorPKc
.Lfunc_end3:
	.size	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_, .Lfunc_end3-_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_
	.cfi_endproc
                                        # -- End function
	.section	.text._ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructIPKcEEvT_S8_St20forward_iterator_tag,"axG",@progbits,_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructIPKcEEvT_S8_St20forward_iterator_tag,comdat
	.weak	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructIPKcEEvT_S8_St20forward_iterator_tag # -- Begin function _ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructIPKcEEvT_S8_St20forward_iterator_tag
	.p2align	1, 0x90
	.type	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructIPKcEEvT_S8_St20forward_iterator_tag,@function
_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructIPKcEEvT_S8_St20forward_iterator_tag: # @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructIPKcEEvT_S8_St20forward_iterator_tag
	.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
	pushq	%rax
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -40
	.cfi_offset %r12, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movq	%rdx, %r14
	movq	%rsi, %r15
	movq	%rdi, %rbx
	subq	%rsi, %r14
	movq	%r14, (%rsp)
	cmpq	$15, %r14
	jbe	.LBB4_1
# %bb.2:
	movq	%rsp, %r12
	movq	%rbx, %rdi
	movq	%r12, %rsi
	xorl	%edx, %edx
	callq	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE9_M_createERmm
	movq	%rax, (%rbx)
	movq	(%r12), %rcx
	movq	%rcx, 16(%rbx)
	jmp	.LBB4_3
.LBB4_1:                                # %._crit_edge
	movq	(%rbx), %rax
.LBB4_3:
	testq	%r14, %r14
	je	.LBB4_7
# %bb.4:
	cmpq	$1, %r14
	jne	.LBB4_6
# %bb.5:
	movb	(%r15), %cl
	movb	%cl, (%rax)
	jmp	.LBB4_7
.LBB4_6:
	movq	%rax, %rdi
	movq	%r15, %rsi
	movq	%r14, %rdx
	callq	memcpy@PLT
.LBB4_7:                                # %_ZZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructIPKcEEvT_S8_St20forward_iterator_tagEN6_GuardD2Ev.exit
	movq	(%rsp), %rax
	movq	%rax, 8(%rbx)
	movq	(%rbx), %rcx
	movb	$0, (%rcx,%rax)
	addq	$8, %rsp
	.cfi_def_cfa_offset 40
	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_end4:
	.size	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructIPKcEEvT_S8_St20forward_iterator_tag, .Lfunc_end4-_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructIPKcEEvT_S8_St20forward_iterator_tag
	.cfi_endproc
                                        # -- End function
	.text
	.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	.LBB5_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB5_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z15deviceIncrementPii, %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_end5:
	.size	__hip_module_ctor, .Lfunc_end5-__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	.LBB6_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
.LBB6_2:
	retq
.Lfunc_end6:
	.size	__hip_module_dtor, .Lfunc_end6-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z15deviceIncrementPii,@object  # @_Z15deviceIncrementPii
	.section	.rodata,"a",@progbits
	.globl	_Z15deviceIncrementPii
	.p2align	3, 0x0
_Z15deviceIncrementPii:
	.quad	_Z30__device_stub__deviceIncrementPii
	.size	_Z15deviceIncrementPii, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"host - %d, device - %d\n"
	.size	.L.str, 24

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"basic_string: construction from null is not valid"
	.size	.L.str.1, 50

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z15deviceIncrementPii"
	.size	.L__unnamed_1, 23

	.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

	.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 _Z30__device_stub__deviceIncrementPii
	.addrsig_sym __gxx_personality_v0
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Unwind_Resume
	.addrsig_sym _Z15deviceIncrementPii
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z15deviceIncrementPii  ; -- Begin function _Z15deviceIncrementPii
	.globl	_Z15deviceIncrementPii
	.p2align	8
	.type	_Z15deviceIncrementPii,@function
_Z15deviceIncrementPii:                 ; @_Z15deviceIncrementPii
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x1c
	s_load_b32 s3, s[0:1], 0x8
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
	s_mov_b32 s2, exec_lo
	v_cmpx_gt_i32_e64 s3, v1
	s_cbranch_execz .LBB0_2
; %bb.1:
	s_load_b64 s[0:1], s[0:1], 0x0
	v_ashrrev_i32_e32 v2, 31, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 2, v[1:2]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v0, vcc_lo, s0, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
	global_load_b32 v2, v[0:1], off
	s_waitcnt vmcnt(0)
	v_add_nc_u32_e32 v2, 1, v2
	global_store_b32 v[0:1], v2, off
.LBB0_2:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z15deviceIncrementPii
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 272
		.amdhsa_user_sgpr_count 15
		.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 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 3
		.amdhsa_next_free_sgpr 16
		.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	_Z15deviceIncrementPii, .Lfunc_end0-_Z15deviceIncrementPii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 140
; NumSgprs: 18
; NumVgprs: 3
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 3
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.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:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .offset:         8
        .size:           4
        .value_kind:     by_value
      - .offset:         16
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         20
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         24
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         28
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         30
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         32
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         34
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         36
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         38
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         56
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         80
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 272
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z15deviceIncrementPii
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z15deviceIncrementPii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     3
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
