#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#include <cuda.h>

#define TPB 1024 //elements per thread
#define BN 64    //block number

/*function declarations*/
int getmax(int *, int);
__global__ void kernel_getmax(int *, int, int);
__device__ void thread_getmax(int *, int *, int, int);

//the sequential version of getmax
int getmax(int num[], int size){
    int i;
    int max=num[0];
    for(i=1;i<size;i++){
	if(num[i]>max){
	    max=num[i];
	}
    }
    return max;
}

/*kernel called by the host to getmax. 
  The high level idea is that each thread first find the max in its share of TPB elements. 
  Then each block use the reduction tree algorithm to find the max in the block. 
  After the two steps above, the max can be found among the first BN elements of the array num[]. 
  We move them back to CPU and use the sequential version to find the max among the last BN elements.*/
__global__ void kernel_getmax(int num[], int size, int workload)
{
    //first, we ask each thread to find the max in its assigned EPT random numbers
    __shared__ int max_each_thread[TPB]; 
    thread_getmax(num, max_each_thread, size, workload);
    __syncthreads();
    
    //next, we find the max in a block. note that the same tree algorithm for parallel summation applies to max as well
    int thread_id=threadIdx.x;
    //loop unrolling for efficiency
    if(thread_id<512){
        if(max_each_thread[thread_id]<max_each_thread[thread_id+512]){
            max_each_thread[thread_id]=max_each_thread[thread_id+512];
        }
    }
    __syncthreads();
    if(thread_id<256){
        if(max_each_thread[thread_id]<max_each_thread[thread_id+256]){
            max_each_thread[thread_id]=max_each_thread[thread_id+256];
        }
    }
    __syncthreads();
    if(thread_id<128){
        if(max_each_thread[thread_id]<max_each_thread[thread_id+128]){
            max_each_thread[thread_id]=max_each_thread[thread_id+128];
        }
    }
    __syncthreads();
    if(thread_id<64){
        if(max_each_thread[thread_id]<max_each_thread[thread_id+64]){
            max_each_thread[thread_id]=max_each_thread[thread_id+64];
        }
    }
    __syncthreads();
    if(thread_id<32){
        if(max_each_thread[thread_id]<max_each_thread[thread_id+32]){
            max_each_thread[thread_id]=max_each_thread[thread_id+32];
        }
    }
    __syncthreads();
    if(thread_id<16){
        if(max_each_thread[thread_id]<max_each_thread[thread_id+16]){
            max_each_thread[thread_id]=max_each_thread[thread_id+16];
        }
    }
    __syncthreads();
    if(thread_id<8){
        if(max_each_thread[thread_id]<max_each_thread[thread_id+8]){
            max_each_thread[thread_id]=max_each_thread[thread_id+8];
        }
    }
    __syncthreads();    
    if(thread_id<4){
        if(max_each_thread[thread_id]<max_each_thread[thread_id+4]){
            max_each_thread[thread_id]=max_each_thread[thread_id+4];
        }
    }
    __syncthreads();
    if(thread_id<2){
        if(max_each_thread[thread_id]<max_each_thread[thread_id+2]){
            max_each_thread[thread_id]=max_each_thread[thread_id+2];
        }
    }
    __syncthreads();
    if(thread_id<1){
        if(max_each_thread[thread_id]<max_each_thread[thread_id+1]){
            max_each_thread[thread_id]=max_each_thread[thread_id+1];
        }
    }
    __syncthreads();

    //we put the max of the i-th block at num[i]
    if(thread_id==0){
	num[blockIdx.x]=max_each_thread[0];
    }
}

/*The function called by the kernel. The sequential getmax version for each thread*/
__device__ void thread_getmax(int num[], int max_each_thread[], int size, int workload){
    int max=0;
    int index=workload*(blockIdx.x*blockDim.x+threadIdx.x);
    int i;
    for(i=index;(i<size)&&(i<index+workload);i++){
	    if(max<num[i]){
		max=num[i];
	    }
    }

    //store the max of this thread in the corresponding position of the shared array
    max_each_thread[threadIdx.x]=max;
}


int main(int argc, char *argv[])
{
    int size = 0;  // The size of the array
    int i;  // loop index
    int * numbers; //pointer to the array
    
    if(argc !=2)
    {
       printf("usage: maxgpu num\n");
       printf("num = size of the array\n");
       exit(1);
    }
   
    size = atol(argv[1]);

    numbers = (int *)malloc(size * sizeof(int));
    if( !numbers )
    {
       printf("Unable to allocate mem for an array of size %u\n", size);
       exit(1);
    }    

    srand(time(NULL)); // setting a seed for the random number generator
    // Fill-up the array with random numbers from 0 to size-1 
    for( i = 0; i < size; i++)
       numbers[i] = rand()  % size;    

    /*todo: 1)allocate memory and copy numbers from host to device 
	    2)invoke kernels to deal with the array
	    3)copy numbers from device to host and free memory*/
    
    int workload=ceil((double)size/(TPB*BN));
    //step1:memory setup
    int *gpu_numbers;
    cudaError_t err;
    err=cudaMalloc((void**)&gpu_numbers, sizeof(int)*size);
    //sometimes when x=100,000,000, we might fail to allocate/transfer memory
    if(err!=cudaSuccess){
	printf("Cannot allocate memory for the initial random array\n");
    }
    err=cudaMemcpy(gpu_numbers, numbers, sizeof(int)*size, cudaMemcpyHostToDevice);
    if(err!=cudaSuccess){
	printf("cannot pass the random array from cpu to gpu\n");
    }
    
    //step2:invoke kernal
    kernel_getmax<<<BN, TPB>>>(gpu_numbers, size, workload);
    
    //step3:copy the max from device & housekeeping(free the allocated pointers)
    cudaMemcpy(numbers,gpu_numbers, sizeof(int)*BN, cudaMemcpyDeviceToHost);
    printf(" The maximum number calculated from GPU is: %d\n", getmax(numbers, BN));
    cudaFree(gpu_numbers);
    free(numbers);
    exit(0);
}



	.file	"tmpxft_002366f9_00000000-6_maxgpu.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2032:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2032:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z6getmaxPii
	.type	_Z6getmaxPii, @function
_Z6getmaxPii:
.LFB2027:
	.cfi_startproc
	endbr64
	movl	(%rdi), %eax
	xorl	%edx, %edx
.L3:
	incq	%rdx
	cmpl	%edx, %esi
	jle	.L6
	movl	(%rdi,%rdx,4), %ecx
	cmpl	%ecx, %eax
	cmovl	%ecx, %eax
	jmp	.L3
.L6:
	ret
	.cfi_endproc
.LFE2027:
	.size	_Z6getmaxPii, .-_Z6getmaxPii
	.globl	_Z13thread_getmaxPiS_ii
	.type	_Z13thread_getmaxPiS_ii, @function
_Z13thread_getmaxPiS_ii:
.LFB2028:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2028:
	.size	_Z13thread_getmaxPiS_ii, .-_Z13thread_getmaxPiS_ii
	.globl	_Z35__device_stub__Z13kernel_getmaxPiiiPiii
	.type	_Z35__device_stub__Z13kernel_getmaxPiiiPiii, @function
_Z35__device_stub__Z13kernel_getmaxPiiiPiii:
.LFB2054:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 8(%rsp)
	leaq	24(%rsp), %rcx
	leaq	32(%rsp), %rdi
	movl	%esi, 4(%rsp)
	leaq	44(%rsp), %rsi
	movl	%edx, (%rsp)
	leaq	16(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 40(%rsp)
	movq	%rax, 80(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 32(%rsp)
	movq	%rax, 44(%rsp)
	movl	$1, 52(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L9
	pushq	24(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z13kernel_getmaxPiii(%rip), %rdi
	pushq	24(%rsp)
	.cfi_def_cfa_offset 144
	movq	60(%rsp), %rcx
	movl	68(%rsp), %r8d
	movq	48(%rsp), %rsi
	movl	56(%rsp), %edx
	leaq	96(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L9:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L11
	call	__stack_chk_fail@PLT
.L11:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2054:
	.size	_Z35__device_stub__Z13kernel_getmaxPiiiPiii, .-_Z35__device_stub__Z13kernel_getmaxPiiiPiii
	.globl	_Z13kernel_getmaxPiii
	.type	_Z13kernel_getmaxPiii, @function
_Z13kernel_getmaxPiii:
.LFB2055:
	.cfi_startproc
	endbr64
	jmp	_Z35__device_stub__Z13kernel_getmaxPiiiPiii
	.cfi_endproc
.LFE2055:
	.size	_Z13kernel_getmaxPiii, .-_Z13kernel_getmaxPiii
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"usage: maxgpu num\n"
.LC1:
	.string	"num = size of the array\n"
.LC2:
	.string	"Unable to allocate mem for an array of size %u\n"
.LC3:
	.string	"Cannot allocate memory for the initial random array\n"
.LC4:
	.string	"cannot pass the random array from cpu to gpu\n"
.LC6:
	.string	" The maximum number calculated from GPU is: %d\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2029:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	subq	$48, %rsp
	.cfi_def_cfa_offset 96
	movq	%fs:40, %rax
	movq	%rax, 40(%rsp)
	xorl	%eax, %eax
	cmpl	$2, %edi
	je	.L15
	leaq	.LC0(%rip), %rsi
	movl	$2, %edi
	call	__printf_chk@PLT
	leaq	.LC1(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	jmp	.L30
.L15:
	movq	8(%rsi), %rdi
	call	atol@PLT
	movslq	%eax, %r13
	movq	%rax, %r12
	movl	%eax, %ebp
	salq	$2, %r13
	movq	%r13, %rdi
	call	malloc@PLT
	movq	%rax, %rbx
	testq	%rax, %rax
	jne	.L16
	movl	%r12d, %edx
	leaq	.LC2(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L30:
	movl	$1, %edi
.L31:
	call	exit@PLT
.L16:
	xorl	%edi, %edi
	xorl	%r14d, %r14d
	call	time@PLT
	movl	%eax, %edi
	call	srand@PLT
.L17:
	cmpl	%r14d, %ebp
	jle	.L32
	call	rand@PLT
	cltd
	idivl	%ebp
	movl	%edx, (%rbx,%r14,4)
	incq	%r14
	jmp	.L17
.L32:
	leaq	8(%rsp), %rdi
	movq	%r13, %rsi
	call	cudaMalloc@PLT
	testl	%eax, %eax
	je	.L19
	leaq	.LC3(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L19:
	movq	8(%rsp), %rdi
	movl	$1, %ecx
	movq	%r13, %rdx
	movq	%rbx, %rsi
	call	cudaMemcpy@PLT
	testl	%eax, %eax
	je	.L20
	leaq	.LC4(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L20:
	movl	$4194305, %edx
	movl	$67108865, %edi
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	salq	$10, %rdx
	salq	$6, %rdi
	movl	$1, %ecx
	movl	$1, %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L21
	cvtsi2sdl	%r12d, %xmm0
	mulsd	.LC5(%rip), %xmm0
	call	ceil@PLT
	movq	8(%rsp), %rdi
	movl	%ebp, %esi
	cvttsd2sil	%xmm0, %edx
	call	_Z35__device_stub__Z13kernel_getmaxPiiiPiii
.L21:
	movq	8(%rsp), %rsi
	movl	$2, %ecx
	movl	$256, %edx
	movq	%rbx, %rdi
	call	cudaMemcpy@PLT
	movl	$64, %esi
	movq	%rbx, %rdi
	call	_Z6getmaxPii
	leaq	.LC6(%rip), %rsi
	movl	$2, %edi
	movl	%eax, %edx
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	movq	%rbx, %rdi
	call	free@PLT
	xorl	%edi, %edi
	jmp	.L31
	.cfi_endproc
.LFE2029:
	.size	main, .-main
	.section	.rodata.str1.1
.LC7:
	.string	"_Z13kernel_getmaxPiii"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2057:
	.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	.LC7(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z13kernel_getmaxPiii(%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
.LFE2057:
	.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
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC5:
	.long	0
	.long	1055916032
	.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:

#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#include <hip/hip_runtime.h>

#define TPB 1024 //elements per thread
#define BN 64    //block number

/*function declarations*/
int getmax(int *, int);
__global__ void kernel_getmax(int *, int, int);
__device__ void thread_getmax(int *, int *, int, int);

//the sequential version of getmax
int getmax(int num[], int size){
    int i;
    int max=num[0];
    for(i=1;i<size;i++){
	if(num[i]>max){
	    max=num[i];
	}
    }
    return max;
}

/*kernel called by the host to getmax. 
  The high level idea is that each thread first find the max in its share of TPB elements. 
  Then each block use the reduction tree algorithm to find the max in the block. 
  After the two steps above, the max can be found among the first BN elements of the array num[]. 
  We move them back to CPU and use the sequential version to find the max among the last BN elements.*/
__global__ void kernel_getmax(int num[], int size, int workload)
{
    //first, we ask each thread to find the max in its assigned EPT random numbers
    __shared__ int max_each_thread[TPB]; 
    thread_getmax(num, max_each_thread, size, workload);
    __syncthreads();
    
    //next, we find the max in a block. note that the same tree algorithm for parallel summation applies to max as well
    int thread_id=threadIdx.x;
    //loop unrolling for efficiency
    if(thread_id<512){
        if(max_each_thread[thread_id]<max_each_thread[thread_id+512]){
            max_each_thread[thread_id]=max_each_thread[thread_id+512];
        }
    }
    __syncthreads();
    if(thread_id<256){
        if(max_each_thread[thread_id]<max_each_thread[thread_id+256]){
            max_each_thread[thread_id]=max_each_thread[thread_id+256];
        }
    }
    __syncthreads();
    if(thread_id<128){
        if(max_each_thread[thread_id]<max_each_thread[thread_id+128]){
            max_each_thread[thread_id]=max_each_thread[thread_id+128];
        }
    }
    __syncthreads();
    if(thread_id<64){
        if(max_each_thread[thread_id]<max_each_thread[thread_id+64]){
            max_each_thread[thread_id]=max_each_thread[thread_id+64];
        }
    }
    __syncthreads();
    if(thread_id<32){
        if(max_each_thread[thread_id]<max_each_thread[thread_id+32]){
            max_each_thread[thread_id]=max_each_thread[thread_id+32];
        }
    }
    __syncthreads();
    if(thread_id<16){
        if(max_each_thread[thread_id]<max_each_thread[thread_id+16]){
            max_each_thread[thread_id]=max_each_thread[thread_id+16];
        }
    }
    __syncthreads();
    if(thread_id<8){
        if(max_each_thread[thread_id]<max_each_thread[thread_id+8]){
            max_each_thread[thread_id]=max_each_thread[thread_id+8];
        }
    }
    __syncthreads();    
    if(thread_id<4){
        if(max_each_thread[thread_id]<max_each_thread[thread_id+4]){
            max_each_thread[thread_id]=max_each_thread[thread_id+4];
        }
    }
    __syncthreads();
    if(thread_id<2){
        if(max_each_thread[thread_id]<max_each_thread[thread_id+2]){
            max_each_thread[thread_id]=max_each_thread[thread_id+2];
        }
    }
    __syncthreads();
    if(thread_id<1){
        if(max_each_thread[thread_id]<max_each_thread[thread_id+1]){
            max_each_thread[thread_id]=max_each_thread[thread_id+1];
        }
    }
    __syncthreads();

    //we put the max of the i-th block at num[i]
    if(thread_id==0){
	num[blockIdx.x]=max_each_thread[0];
    }
}

/*The function called by the kernel. The sequential getmax version for each thread*/
__device__ void thread_getmax(int num[], int max_each_thread[], int size, int workload){
    int max=0;
    int index=workload*(blockIdx.x*blockDim.x+threadIdx.x);
    int i;
    for(i=index;(i<size)&&(i<index+workload);i++){
	    if(max<num[i]){
		max=num[i];
	    }
    }

    //store the max of this thread in the corresponding position of the shared array
    max_each_thread[threadIdx.x]=max;
}


int main(int argc, char *argv[])
{
    int size = 0;  // The size of the array
    int i;  // loop index
    int * numbers; //pointer to the array
    
    if(argc !=2)
    {
       printf("usage: maxgpu num\n");
       printf("num = size of the array\n");
       exit(1);
    }
   
    size = atol(argv[1]);

    numbers = (int *)malloc(size * sizeof(int));
    if( !numbers )
    {
       printf("Unable to allocate mem for an array of size %u\n", size);
       exit(1);
    }    

    srand(time(NULL)); // setting a seed for the random number generator
    // Fill-up the array with random numbers from 0 to size-1 
    for( i = 0; i < size; i++)
       numbers[i] = rand()  % size;    

    /*todo: 1)allocate memory and copy numbers from host to device 
	    2)invoke kernels to deal with the array
	    3)copy numbers from device to host and free memory*/
    
    int workload=ceil((double)size/(TPB*BN));
    //step1:memory setup
    int *gpu_numbers;
    hipError_t err;
    err=hipMalloc((void**)&gpu_numbers, sizeof(int)*size);
    //sometimes when x=100,000,000, we might fail to allocate/transfer memory
    if(err!=hipSuccess){
	printf("Cannot allocate memory for the initial random array\n");
    }
    err=hipMemcpy(gpu_numbers, numbers, sizeof(int)*size, hipMemcpyHostToDevice);
    if(err!=hipSuccess){
	printf("cannot pass the random array from cpu to gpu\n");
    }
    
    //step2:invoke kernal
    kernel_getmax<<<BN, TPB>>>(gpu_numbers, size, workload);
    
    //step3:copy the max from device & housekeeping(free the allocated pointers)
    hipMemcpy(numbers,gpu_numbers, sizeof(int)*BN, hipMemcpyDeviceToHost);
    printf(" The maximum number calculated from GPU is: %d\n", getmax(numbers, BN));
    hipFree(gpu_numbers);
    free(numbers);
    exit(0);
}



	.text
	.file	"maxgpu.hip"
	.globl	_Z6getmaxPii                    # -- Begin function _Z6getmaxPii
	.type	_Z6getmaxPii,@function
_Z6getmaxPii:                           # @_Z6getmaxPii
	.cfi_startproc
# %bb.0:
	movl	(%rdi), %eax
	cmpl	$2, %esi
	jl	.LBB0_3
# %bb.1:                                # %.lr.ph.preheader
	movl	%esi, %ecx
	movl	$1, %edx
.LBB0_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movl	(%rdi,%rdx,4), %esi
	cmpl	%eax, %esi
	cmovgl	%esi, %eax
	incq	%rdx
	cmpq	%rdx, %rcx
	jne	.LBB0_2
.LBB0_3:                                # %._crit_edge
	retq
.Lfunc_end0:
	.size	_Z6getmaxPii, .Lfunc_end0-_Z6getmaxPii
	.cfi_endproc
                                        # -- End function
	.globl	_Z28__device_stub__kernel_getmaxPiii # -- Begin function _Z28__device_stub__kernel_getmaxPiii
	.type	_Z28__device_stub__kernel_getmaxPiii,@function
_Z28__device_stub__kernel_getmaxPiii:   # @_Z28__device_stub__kernel_getmaxPiii
	.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	$96, %rsp
	.cfi_def_cfa_offset 144
	.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)
	movq	%rsp, %rsi
	movl	%edx, (%rsi)
	leaq	64(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	movq	%rsi, 16(%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	$_Z13kernel_getmaxPiii, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$112, %rsp
	.cfi_adjust_cfa_offset -112
	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	_Z28__device_stub__kernel_getmaxPiii, .Lfunc_end1-_Z28__device_stub__kernel_getmaxPiii
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function main
.LCPI2_0:
	.quad	0x3ef0000000000000              # double 1.52587890625E-5
	.text
	.globl	main
	.type	main,@function
main:                                   # @main
	.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	$16, %rsp
	.cfi_def_cfa_offset 64
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	cmpl	$2, %edi
	jne	.LBB2_4
# %bb.1:
	movq	8(%rsi), %rdi
	callq	atol
	movq	%rax, %r14
	movslq	%r14d, %r15
	shlq	$2, %r15
	movq	%r15, %rdi
	callq	malloc
	testq	%rax, %rax
	jne	.LBB2_5
# %bb.2:
	movl	$.L.str.2, %edi
	movl	%r14d, %esi
	xorl	%eax, %eax
	callq	printf
	jmp	.LBB2_3
.LBB2_4:
	movl	$.Lstr.2, %edi
	callq	puts@PLT
	movl	$.Lstr.3, %edi
	callq	puts@PLT
.LBB2_3:
	movl	$1, %edi
	callq	exit
.LBB2_5:
	movq	%rax, %rbx
	xorl	%edi, %edi
	callq	time
	movl	%eax, %edi
	callq	srand
	testl	%r14d, %r14d
	jle	.LBB2_8
# %bb.6:                                # %.lr.ph.preheader
	movl	%r14d, %r12d
	xorl	%r13d, %r13d
.LBB2_7:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	callq	rand
	cltd
	idivl	%r14d
	movl	%edx, (%rbx,%r13,4)
	incq	%r13
	cmpq	%r13, %r12
	jne	.LBB2_7
.LBB2_8:                                # %._crit_edge
	cvtsi2sd	%r14d, %xmm0
	mulsd	.LCPI2_0(%rip), %xmm0
	callq	ceil@PLT
	movsd	%xmm0, 8(%rsp)                  # 8-byte Spill
	movq	%rsp, %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	testl	%eax, %eax
	je	.LBB2_10
# %bb.9:
	movl	$.Lstr, %edi
	callq	puts@PLT
.LBB2_10:
	movq	(%rsp), %rdi
	movq	%rbx, %rsi
	movq	%r15, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	testl	%eax, %eax
	je	.LBB2_12
# %bb.11:
	movl	$.Lstr.1, %edi
	callq	puts@PLT
.LBB2_12:
	movabsq	$4294967360, %rdi               # imm = 0x100000040
	leaq	960(%rdi), %rdx
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB2_14
# %bb.13:
	cvttsd2si	8(%rsp), %edx           # 8-byte Folded Reload
	movq	(%rsp), %rdi
	movl	%r14d, %esi
	callq	_Z28__device_stub__kernel_getmaxPiii
.LBB2_14:
	movq	(%rsp), %rsi
	movl	$256, %edx                      # imm = 0x100
	movq	%rbx, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movl	(%rbx), %esi
	movl	$1, %eax
.LBB2_15:                               # %.lr.ph.i
                                        # =>This Inner Loop Header: Depth=1
	movl	(%rbx,%rax,4), %ecx
	cmpl	%esi, %ecx
	cmovgl	%ecx, %esi
	incq	%rax
	cmpq	$64, %rax
	jne	.LBB2_15
# %bb.16:                               # %_Z6getmaxPii.exit
	movl	$.L.str.5, %edi
	xorl	%eax, %eax
	callq	printf
	movq	(%rsp), %rdi
	callq	hipFree
	movq	%rbx, %rdi
	callq	free
	xorl	%edi, %edi
	callq	exit
.Lfunc_end2:
	.size	main, .Lfunc_end2-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	.LBB3_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB3_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z13kernel_getmaxPiii, %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_end3:
	.size	__hip_module_ctor, .Lfunc_end3-__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	.LBB4_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
.LBB4_2:
	retq
.Lfunc_end4:
	.size	__hip_module_dtor, .Lfunc_end4-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z13kernel_getmaxPiii,@object   # @_Z13kernel_getmaxPiii
	.section	.rodata,"a",@progbits
	.globl	_Z13kernel_getmaxPiii
	.p2align	3, 0x0
_Z13kernel_getmaxPiii:
	.quad	_Z28__device_stub__kernel_getmaxPiii
	.size	_Z13kernel_getmaxPiii, 8

	.type	.L.str.2,@object                # @.str.2
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.2:
	.asciz	"Unable to allocate mem for an array of size %u\n"
	.size	.L.str.2, 48

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	" The maximum number calculated from GPU is: %d\n"
	.size	.L.str.5, 48

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z13kernel_getmaxPiii"
	.size	.L__unnamed_1, 22

	.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	"Cannot allocate memory for the initial random array"
	.size	.Lstr, 52

	.type	.Lstr.1,@object                 # @str.1
.Lstr.1:
	.asciz	"cannot pass the random array from cpu to gpu"
	.size	.Lstr.1, 45

	.type	.Lstr.2,@object                 # @str.2
.Lstr.2:
	.asciz	"usage: maxgpu num"
	.size	.Lstr.2, 18

	.type	.Lstr.3,@object                 # @str.3
.Lstr.3:
	.asciz	"num = size of the array"
	.size	.Lstr.3, 24

	.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 _Z28__device_stub__kernel_getmaxPiii
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z13kernel_getmaxPiii
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z13kernel_getmaxPiii   ; -- Begin function _Z13kernel_getmaxPiii
	.globl	_Z13kernel_getmaxPiii
	.p2align	8
	.type	_Z13kernel_getmaxPiii,@function
_Z13kernel_getmaxPiii:                  ; @_Z13kernel_getmaxPiii
; %bb.0:
	s_clause 0x1
	s_load_b32 s3, s[0:1], 0x1c
	s_load_b128 s[4:7], s[0:1], 0x0
	s_mov_b32 s2, s15
	v_mov_b32_e32 v4, 0
	s_mov_b32 s1, exec_lo
	s_waitcnt lgkmcnt(0)
	s_and_b32 s0, s3, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s2, s0, v[0:1]
	v_mul_lo_u32 v1, v1, s7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v2, s7, v1
	v_min_i32_e32 v5, s6, v2
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_lt_i32_e64 v1, v5
	s_cbranch_execz .LBB0_4
; %bb.1:                                ; %.lr.ph.i.preheader
	v_ashrrev_i32_e32 v2, 31, v1
	v_mov_b32_e32 v4, 0
	s_mov_b32 s3, 0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[2:3], 2, v[1:2]
	v_add_co_u32 v2, vcc_lo, s4, v2
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v3, vcc_lo, s5, v3, vcc_lo
.LBB0_2:                                ; %.lr.ph.i
                                        ; =>This Inner Loop Header: Depth=1
	global_load_b32 v6, v[2:3], off
	v_add_nc_u32_e32 v1, 1, v1
	v_add_co_u32 v2, vcc_lo, v2, 4
	v_add_co_ci_u32_e32 v3, vcc_lo, 0, v3, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_ge_i32_e64 s0, v1, v5
	s_or_b32 s3, s0, s3
	s_waitcnt vmcnt(0)
	v_max_i32_e32 v4, v4, v6
	s_and_not1_b32 exec_lo, exec_lo, s3
	s_cbranch_execnz .LBB0_2
; %bb.3:                                ; %Flow82
	s_or_b32 exec_lo, exec_lo, s3
.LBB0_4:                                ; %Flow83
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s1
	v_lshlrev_b32_e32 v3, 2, v0
	s_mov_b32 s0, exec_lo
	ds_store_b32 v3, v4
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	v_cmpx_gt_u32_e32 0x200, v0
	s_cbranch_execz .LBB0_7
; %bb.5:
	ds_load_2addr_stride64_b32 v[1:2], v3 offset1:8
	s_waitcnt lgkmcnt(0)
	v_cmp_lt_i32_e32 vcc_lo, v1, v2
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_7
; %bb.6:
	ds_store_b32 v3, v2
.LBB0_7:
	s_or_b32 exec_lo, exec_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_mov_b32 s0, exec_lo
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	v_cmpx_gt_u32_e32 0x100, v0
	s_cbranch_execz .LBB0_10
; %bb.8:
	ds_load_b32 v2, v3
	ds_load_b32 v1, v3 offset:1024
	s_waitcnt lgkmcnt(0)
	v_cmp_lt_i32_e32 vcc_lo, v2, v1
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_10
; %bb.9:
	ds_store_b32 v3, v1
.LBB0_10:
	s_or_b32 exec_lo, exec_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_mov_b32 s0, exec_lo
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	v_cmpx_gt_u32_e32 0x80, v0
	s_cbranch_execz .LBB0_13
; %bb.11:
	ds_load_b32 v2, v3
	ds_load_b32 v1, v3 offset:512
	s_waitcnt lgkmcnt(0)
	v_cmp_lt_i32_e32 vcc_lo, v2, v1
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_13
; %bb.12:
	ds_store_b32 v3, v1
.LBB0_13:
	s_or_b32 exec_lo, exec_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_mov_b32 s0, exec_lo
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	v_cmpx_gt_u32_e32 64, v0
	s_cbranch_execz .LBB0_16
; %bb.14:
	ds_load_b32 v2, v3
	ds_load_b32 v1, v3 offset:256
	s_waitcnt lgkmcnt(0)
	v_cmp_lt_i32_e32 vcc_lo, v2, v1
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_16
; %bb.15:
	ds_store_b32 v3, v1
.LBB0_16:
	s_or_b32 exec_lo, exec_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_mov_b32 s0, exec_lo
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	v_cmpx_gt_u32_e32 32, v0
	s_cbranch_execz .LBB0_19
; %bb.17:
	ds_load_b32 v2, v3
	ds_load_b32 v1, v3 offset:128
	s_waitcnt lgkmcnt(0)
	v_cmp_lt_i32_e32 vcc_lo, v2, v1
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_19
; %bb.18:
	ds_store_b32 v3, v1
.LBB0_19:
	s_or_b32 exec_lo, exec_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_mov_b32 s0, exec_lo
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	v_cmpx_gt_u32_e32 16, v0
	s_cbranch_execz .LBB0_22
; %bb.20:
	ds_load_b32 v2, v3
	ds_load_b32 v1, v3 offset:64
	s_waitcnt lgkmcnt(0)
	v_cmp_lt_i32_e32 vcc_lo, v2, v1
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_22
; %bb.21:
	ds_store_b32 v3, v1
.LBB0_22:
	s_or_b32 exec_lo, exec_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_mov_b32 s0, exec_lo
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	v_cmpx_gt_u32_e32 8, v0
	s_cbranch_execz .LBB0_25
; %bb.23:
	ds_load_b32 v2, v3
	ds_load_b32 v1, v3 offset:32
	s_waitcnt lgkmcnt(0)
	v_cmp_lt_i32_e32 vcc_lo, v2, v1
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_25
; %bb.24:
	ds_store_b32 v3, v1
.LBB0_25:
	s_or_b32 exec_lo, exec_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_mov_b32 s0, exec_lo
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	v_cmpx_gt_u32_e32 4, v0
	s_cbranch_execz .LBB0_28
; %bb.26:
	ds_load_b32 v2, v3
	ds_load_b32 v1, v3 offset:16
	s_waitcnt lgkmcnt(0)
	v_cmp_lt_i32_e32 vcc_lo, v2, v1
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_28
; %bb.27:
	ds_store_b32 v3, v1
.LBB0_28:
	s_or_b32 exec_lo, exec_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_mov_b32 s0, exec_lo
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	v_cmpx_gt_u32_e32 2, v0
	s_cbranch_execz .LBB0_31
; %bb.29:
	ds_load_b32 v2, v3
	ds_load_b32 v1, v3 offset:8
	s_waitcnt lgkmcnt(0)
	v_cmp_lt_i32_e32 vcc_lo, v2, v1
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_31
; %bb.30:
	ds_store_b32 v3, v1
.LBB0_31:
	s_or_b32 exec_lo, exec_lo, s0
	v_cmp_eq_u32_e64 s0, 0, v0
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	s_and_saveexec_b32 s1, s0
	s_cbranch_execz .LBB0_34
; %bb.32:
	v_mov_b32_e32 v2, 0
	ds_load_b64 v[0:1], v2
	s_waitcnt lgkmcnt(0)
	v_cmp_ge_i32_e32 vcc_lo, v0, v1
	s_cbranch_vccnz .LBB0_34
; %bb.33:
	ds_store_b32 v2, v1
.LBB0_34:
	s_or_b32 exec_lo, exec_lo, s1
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	s_and_saveexec_b32 s1, s0
	s_cbranch_execz .LBB0_36
; %bb.35:
	v_mov_b32_e32 v0, 0
	s_mov_b32 s3, 0
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_lshl_b64 s[0:1], s[2:3], 2
	s_add_u32 s0, s4, s0
	ds_load_b32 v1, v0
	s_addc_u32 s1, s5, s1
	s_waitcnt lgkmcnt(0)
	global_store_b32 v0, v1, s[0:1]
.LBB0_36:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z13kernel_getmaxPiii
		.amdhsa_group_segment_fixed_size 4096
		.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 7
		.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	_Z13kernel_getmaxPiii, .Lfunc_end0-_Z13kernel_getmaxPiii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 1044
; NumSgprs: 18
; NumVgprs: 7
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 4096 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 7
; 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:         12
        .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: 4096
    .kernarg_segment_align: 8
    .kernarg_segment_size: 272
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z13kernel_getmaxPiii
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z13kernel_getmaxPiii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     7
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
