/*
========================================================================================
Name: cuda_demo.cu
Author: Mingran Peng
Class: EECS 587, Fall 2020
Description : Demo program for HW4
P.S. Fell free to use or modify this code for future terms of EECS 587 or other courses
Add you name if you modify it and preserve all author names
========================================================================================
*/
#include <algorithm>
#include <iostream>
#include <cstdio>
#include <cmath>
#include <numeric>
#include <iostream>
#include <cstring>
#include <sstream>
#include <string>

using namespace std;

cudaError_t addWithCuda(double *c, unsigned int size, int n, double *Along, int t);

__global__ void addKernel(double *c, int size, int n, double *dev_Along, int t)
{
    double arr[4];
    double secondSmallest=0;
    double temp;
    double update;
    // calculate index here
    int localro = threadIdx.x;
    int localco = threadIdx.y;
    int globalro = blockIdx.x * blockDim.x + threadIdx.x;
    int globalco = blockIdx.y * blockDim.y + threadIdx.y;
    int globalind = globalro * n + globalco;
    // Synchronize all threads in the block to make sure copy is done
    __syncthreads();
    if ((globalro>0) && (globalco>0) && (globalro<(n-1)) && (globalco<(n-1))){
        int il = globalro+1;
        int is = globalro-1;
        int jl = globalco+1;
        int js = globalco-1;
        arr[0] = dev_Along[il*n+jl];
        arr[1] = dev_Along[il*n+js];
        arr[2] = dev_Along[is*n+jl];
        arr[3] = dev_Along[is*n+js];
        for(int i=0;i<4;i++)
            {
                for(int j=i+1;j<4;j++)
                {
                    if(arr[i]>arr[j])
                    {
                        temp  =arr[i];
                        arr[i]=arr[j];
                        arr[j]=temp;
                    }
                }
            }
        secondSmallest = arr[1];
        update = dev_Along[globalind] + secondSmallest;
    }
    else if ((globalro==0)||(globalco==0)||(globalro==(n-1))||(globalco==(n-1))){
        update = dev_Along[globalind];
    }
    else{}
    if((globalro<n)&&(globalco<n))
        c[globalind] = update;
}

int main(int argc, char* argv[])
{
    int n;
    int t;
    n = atoi(argv[1]);
    t = atoi(argv[2]);

    int size = n*n;//number of elements
    double *c; // returned array
    //initiallize
    double *Along;
    Along = new double [size];
    c = new double [size];
    for (int i=0;i<size;i++){
        int ro = floor(i/n);
        int co = floor(i%n);
        Along[i] = pow((1+cos(2*ro)+sin(co)),2);
    }


    cudaError_t cudaStatus = addWithCuda(c, size, n, Along, t);
    if (cudaStatus != cudaSuccess) {
        cout<<"addWithCuda failed!"<<endl;
        return -1;
    }
    
    //here we get the c array then we can do the sum and check the certain element
    // examine
    double initial_sum = 0;
    double sumc = accumulate(c, c+size, initial_sum);
    cout<<"Sum: "<<sumc<<endl;
    cout<<"A(37,47): "<<c[37*n+47]<<endl;
    return 0;
}

// Helper function for using CUDA to add vectors in parallel.
cudaError_t addWithCuda(double *c, unsigned int size, int n, double *Along, int t)
{
    double *dev_Along = 0;
    //dev_Along = new double[size];
    double *dev_c = 0;

    dev_c = new double[size];
    for (int i =0;i<size;i++){
        dev_c[i]=0;
    }
    cudaError_t cudaStatus;
    cudaEvent_t start, stop;
    float gpu_time = 0.0f;
    dim3 gridSize(ceil(n/32)+1,ceil(n/32)+1,1);
    dim3 blockSize(32,32,1);


    // Choose which GPU to run on, 0 if you have only one GPU
    // on-chip GPU does not count
    cudaStatus = cudaSetDevice(0);
    if (cudaStatus != cudaSuccess) {
        cout<<"cudaSetDevice failed!"<<endl;
        goto Error;
    }
    // Malloc memory on GPU
    
    
    cudaStatus = cudaMalloc((void**)&dev_Along, size * sizeof(double));
    if (cudaStatus != cudaSuccess) {
        cout<<"cudaMalloc failed!"<<endl;
        goto Error;
    }
    cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(double));
    if (cudaStatus != cudaSuccess) {
        cout<<"cudaMalloc failed!"<<endl;
        goto Error;
    }
    
    // Copy memory from Host to Device
    cudaStatus = cudaMemcpy(dev_Along, Along, size * sizeof(double), cudaMemcpyHostToDevice);
    if (cudaStatus != cudaSuccess) {
        cout<<"cudaMemcpy failed!"<<endl;
        goto Error;
    }
    // Set up timing
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    cudaEventRecord(start);
    // Launch a kernel on the GPU with one thread for each element
    //cout<<"allocating "<<gridSize<<" blocks, "<<blockSize<<" threads per block"<<endl;
    for (int tt=0;tt<t;tt++){
        addKernel<<<gridSize, blockSize>>>(dev_c, size, n,dev_Along,t);

        if (cudaStatus != cudaSuccess) {
            cout<<"cudaMemcpy failed!"<<endl;
            goto Error;
        }
        cudaStatus = cudaMemcpy(dev_Along, dev_c, size * sizeof(double), cudaMemcpyDeviceToDevice);
    }
    
    
    
      // Check for any errors launching the kernel
    cudaStatus = cudaGetLastError();
    if (cudaStatus != cudaSuccess) {
        cout<<"addKernel failed: "<<cudaGetErrorString(cudaStatus)<<endl;
        goto Error;
    }
    // cudaDeviceSynchronize waits for the kernel to finish, and returns
    // any errors encountered during the launch.
    cudaStatus = cudaDeviceSynchronize();
    if (cudaStatus != cudaSuccess) {
        cout<<"cudaDeviceSynchronize failed: "<<cudaGetErrorString(cudaStatus)<<endl;
        goto Error;
    }
    // Close timing
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);
    cudaEventElapsedTime(&gpu_time, start, stop);
    cout<<"Time spent: "<<gpu_time<<"ms"<<endl;
    cudaEventDestroy(start);
    cudaEventDestroy(stop);

    // Copy memory from devide to host
    cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(double), cudaMemcpyDeviceToHost);
    if (cudaStatus != cudaSuccess) {
        cout<<"cudaMemcpy failed!"<<endl;
        goto Error;
    }

Error:
    cudaFree(dev_c);
    cudaFree(dev_Along);


    return cudaStatus;
}

	.file	"tmpxft_0021fb7a_00000000-6_hw4_4.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB3978:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE3978:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.align 2
	.type	_ZNSolsEPFRSoS_E.isra.0, @function
_ZNSolsEPFRSoS_E.isra.0:
.LFB4640:
	.cfi_startproc
	jmp	*%rsi
	.cfi_endproc
.LFE4640:
	.size	_ZNSolsEPFRSoS_E.isra.0, .-_ZNSolsEPFRSoS_E.isra.0
	.globl	_Z33__device_stub__Z9addKernelPdiiS_iPdiiS_i
	.type	_Z33__device_stub__Z9addKernelPdiiS_iPdiiS_i, @function
_Z33__device_stub__Z9addKernelPdiiS_iPdiiS_i:
.LFB4000:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rdi
	movl	%esi, 20(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 16(%rsp)
	leaq	32(%rsp), %rdx
	movq	%rcx, 8(%rsp)
	leaq	40(%rsp), %rcx
	movl	%r8d, 4(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L3
	pushq	40(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z9addKernelPdiiS_i(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 176
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L3:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L5
	call	__stack_chk_fail@PLT
.L5:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4000:
	.size	_Z33__device_stub__Z9addKernelPdiiS_iPdiiS_i, .-_Z33__device_stub__Z9addKernelPdiiS_iPdiiS_i
	.globl	_Z9addKernelPdiiS_i
	.type	_Z9addKernelPdiiS_i, @function
_Z9addKernelPdiiS_i:
.LFB4001:
	.cfi_startproc
	endbr64
	jmp	_Z33__device_stub__Z9addKernelPdiiS_iPdiiS_i
	.cfi_endproc
.LFE4001:
	.size	_Z9addKernelPdiiS_i, .-_Z9addKernelPdiiS_i
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC3:
	.string	"cudaSetDevice failed!"
.LC4:
	.string	"cudaMalloc failed!"
.LC5:
	.string	"cudaMemcpy failed!"
.LC6:
	.string	"addKernel failed: "
.LC7:
	.string	"cudaDeviceSynchronize failed: "
.LC8:
	.string	"Time spent: "
.LC9:
	.string	"ms"
	.text
	.globl	_Z11addWithCudaPdjiS_i
	.type	_Z11addWithCudaPdjiS_i, @function
_Z11addWithCudaPdjiS_i:
.LFB3975:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movq	%rcx, %r15
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movq	%rdi, %r14
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movl	%r8d, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movl	%edx, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movl	%esi, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$104, %rsp
	.cfi_def_cfa_offset 160
	movq	%fs:40, %rax
	movq	%rax, 88(%rsp)
	xorl	%eax, %eax
	xorl	%eax, %eax
	movq	%rax, 32(%rsp)
	movq	%rax, 40(%rsp)
	movl	%esi, %eax
	salq	$3, %rax
	movq	%rax, %rdi
	movq	%rax, 8(%rsp)
	call	_Znam@PLT
	xorl	%edx, %edx
	movq	%rax, 40(%rsp)
.L10:
	cmpl	%ebp, %edx
	jnb	.L33
	movq	$0x000000000, (%rax,%rdx,8)
	incq	%rdx
	jmp	.L10
.L33:
	movl	%r12d, %eax
	movl	$32, %ecx
	xorl	%edi, %edi
	movl	$0x00000000, 28(%rsp)
	cltd
	movl	$1, 72(%rsp)
	idivl	%ecx
	movl	$1, 84(%rsp)
	cvtsi2sdl	%eax, %xmm0
	addsd	.LC2(%rip), %xmm0
	cvttsd2siq	%xmm0, %rax
	movl	%eax, 64(%rsp)
	movl	%eax, 68(%rsp)
	movabsq	$137438953504, %rax
	movq	%rax, 76(%rsp)
	call	cudaSetDevice@PLT
	leaq	.LC3(%rip), %rsi
	movl	%eax, %ebx
	testl	%eax, %eax
	jne	.L30
	movq	8(%rsp), %rsi
	leaq	32(%rsp), %rdi
	call	cudaMalloc@PLT
	movl	%eax, %ebx
	testl	%eax, %eax
	je	.L14
.L31:
	leaq	.LC4(%rip), %rsi
.L30:
	leaq	_ZSt4cout(%rip), %rdi
	jmp	.L27
.L14:
	movq	8(%rsp), %rsi
	leaq	40(%rsp), %rdi
	call	cudaMalloc@PLT
	movl	%eax, %ebx
	testl	%eax, %eax
	jne	.L31
	movq	8(%rsp), %rdx
	movq	32(%rsp), %rdi
	movl	$1, %ecx
	movq	%r15, %rsi
	call	cudaMemcpy@PLT
	movl	%eax, %ebx
	testl	%eax, %eax
	jne	.L32
	leaq	48(%rsp), %rdi
	xorl	%r15d, %r15d
	call	cudaEventCreate@PLT
	leaq	56(%rsp), %rdi
	call	cudaEventCreate@PLT
	movq	48(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	jmp	.L17
.L32:
	leaq	.LC5(%rip), %rsi
	jmp	.L30
.L20:
	movl	84(%rsp), %ecx
	movq	76(%rsp), %rdx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	64(%rsp), %rdi
	movl	72(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L18
	movq	32(%rsp), %rcx
	movq	40(%rsp), %rdi
	movl	%r13d, %r8d
	movl	%r12d, %edx
	movl	%ebp, %esi
	call	_Z33__device_stub__Z9addKernelPdiiS_iPdiiS_i
.L18:
	testl	%ebx, %ebx
	jne	.L32
	movq	8(%rsp), %rdx
	movq	40(%rsp), %rsi
	movl	$3, %ecx
	incl	%r15d
	movq	32(%rsp), %rdi
	call	cudaMemcpy@PLT
	movl	%eax, %ebx
.L17:
	cmpl	%r13d, %r15d
	jl	.L20
	call	cudaGetLastError@PLT
	leaq	_ZSt4cout(%rip), %rbp
	leaq	.LC6(%rip), %rsi
	testl	%eax, %eax
	movl	%eax, %ebx
	jne	.L29
	call	cudaDeviceSynchronize@PLT
	movl	%eax, %ebx
	testl	%eax, %eax
	je	.L22
	leaq	.LC7(%rip), %rsi
.L29:
	movq	%rbp, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%ebx, %edi
	movq	%rax, %rbp
	call	cudaGetErrorString@PLT
	movq	%rax, %rsi
	jmp	.L28
.L22:
	movq	56(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	56(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movq	56(%rsp), %rdx
	movq	48(%rsp), %rsi
	leaq	28(%rsp), %rdi
	call	cudaEventElapsedTime@PLT
	leaq	.LC8(%rip), %rsi
	movq	%rbp, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	cvtss2sd	28(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	leaq	.LC9(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movq	48(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movq	56(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movq	8(%rsp), %rdx
	movq	40(%rsp), %rsi
	movq	%r14, %rdi
	movl	$2, %ecx
	call	cudaMemcpy@PLT
	movl	%eax, %ebx
	testl	%eax, %eax
	je	.L13
	leaq	.LC5(%rip), %rsi
.L28:
	movq	%rbp, %rdi
.L27:
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
.L13:
	movq	40(%rsp), %rdi
	call	cudaFree@PLT
	movq	32(%rsp), %rdi
	call	cudaFree@PLT
	movq	88(%rsp), %rax
	subq	%fs:40, %rax
	je	.L25
	call	__stack_chk_fail@PLT
.L25:
	addq	$104, %rsp
	.cfi_def_cfa_offset 56
	movl	%ebx, %eax
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.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
	ret
	.cfi_endproc
.LFE3975:
	.size	_Z11addWithCudaPdjiS_i, .-_Z11addWithCudaPdjiS_i
	.section	.rodata.str1.1
.LC10:
	.string	"addWithCuda failed!"
.LC11:
	.string	"Sum: "
.LC12:
	.string	"A(37,47): "
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB3974:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	xorl	%r15d, %r15d
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%rsi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$40, %rsp
	.cfi_def_cfa_offset 96
	movq	8(%rsi), %rdi
	call	atoi@PLT
	movq	16(%rbp), %rdi
	movl	%eax, %ebx
	call	atoi@PLT
	movl	%ebx, %r14d
	imull	%ebx, %r14d
	movl	%eax, 8(%rsp)
	movslq	%r14d, %rbp
	salq	$3, %rbp
	movq	%rbp, %rdi
	call	_Znam@PLT
	movq	%rbp, %rdi
	movq	%rax, %r13
	call	_Znam@PLT
	movq	%rax, %r12
.L35:
	movl	%r15d, %eax
	cmpl	%r15d, %r14d
	jle	.L42
	cltd
	idivl	%ebx
	addl	%eax, %eax
	movl	%edx, 28(%rsp)
	cvtsi2sdl	%eax, %xmm0
	call	cos@PLT
	movl	28(%rsp), %edx
	movsd	%xmm0, 16(%rsp)
	cvtsi2sdl	%edx, %xmm0
	call	sin@PLT
	movsd	16(%rsp), %xmm1
	movaps	%xmm0, %xmm2
	movsd	.LC2(%rip), %xmm0
	addsd	%xmm1, %xmm0
	addsd	%xmm2, %xmm0
	mulsd	%xmm0, %xmm0
	movsd	%xmm0, 0(%r13,%r15,8)
	incq	%r15
	jmp	.L35
.L42:
	movl	8(%rsp), %r8d
	movq	%r13, %rcx
	movl	%ebx, %edx
	movl	%r14d, %esi
	movq	%r12, %rdi
	call	_Z11addWithCudaPdjiS_i
	testl	%eax, %eax
	je	.L37
	leaq	.LC10(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	orl	$-1, %eax
	jmp	.L34
.L37:
	addq	%r12, %rbp
	movq	%r12, %rax
	xorps	%xmm0, %xmm0
.L39:
	cmpq	%rax, %rbp
	je	.L43
	addsd	(%rax), %xmm0
	addq	$8, %rax
	jmp	.L39
.L43:
	movsd	%xmm0, 8(%rsp)
	imull	$37, %ebx, %ebx
	leaq	_ZSt4cout(%rip), %rbp
	leaq	.LC11(%rip), %rsi
	movq	%rbp, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movsd	8(%rsp), %xmm0
	movq	%rax, %rdi
	movslq	%ebx, %rbx
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	.LC12(%rip), %rsi
	movq	%rbp, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movsd	376(%r12,%rbx,8), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	xorl	%eax, %eax
.L34:
	addq	$40, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.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
	ret
	.cfi_endproc
.LFE3974:
	.size	main, .-main
	.section	.rodata.str1.1
.LC13:
	.string	"_Z9addKernelPdiiS_i"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB4003:
	.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	.LC13(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z9addKernelPdiiS_i(%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
.LFE4003:
	.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
.LC2:
	.long	0
	.long	1072693248
	.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:

/*
========================================================================================
Name: cuda_demo.cu
Author: Mingran Peng
Class: EECS 587, Fall 2020
Description : Demo program for HW4
P.S. Fell free to use or modify this code for future terms of EECS 587 or other courses
Add you name if you modify it and preserve all author names
========================================================================================
*/

#include <hip/hip_runtime.h>
#include <algorithm>
#include <iostream>
#include <cstdio>
#include <cmath>
#include <numeric>
#include <iostream>
#include <cstring>
#include <sstream>
#include <string>

using namespace std;

hipError_t addWithCuda(double *c, unsigned int size, int n, double *Along, int t);

__global__ void addKernel(double *c, int size, int n, double *dev_Along, int t)
{
    double arr[4];
    double secondSmallest=0;
    double temp;
    double update;
    // calculate index here
    int localro = threadIdx.x;
    int localco = threadIdx.y;
    int globalro = blockIdx.x * blockDim.x + threadIdx.x;
    int globalco = blockIdx.y * blockDim.y + threadIdx.y;
    int globalind = globalro * n + globalco;
    // Synchronize all threads in the block to make sure copy is done
    __syncthreads();
    if ((globalro>0) && (globalco>0) && (globalro<(n-1)) && (globalco<(n-1))){
        int il = globalro+1;
        int is = globalro-1;
        int jl = globalco+1;
        int js = globalco-1;
        arr[0] = dev_Along[il*n+jl];
        arr[1] = dev_Along[il*n+js];
        arr[2] = dev_Along[is*n+jl];
        arr[3] = dev_Along[is*n+js];
        for(int i=0;i<4;i++)
            {
                for(int j=i+1;j<4;j++)
                {
                    if(arr[i]>arr[j])
                    {
                        temp  =arr[i];
                        arr[i]=arr[j];
                        arr[j]=temp;
                    }
                }
            }
        secondSmallest = arr[1];
        update = dev_Along[globalind] + secondSmallest;
    }
    else if ((globalro==0)||(globalco==0)||(globalro==(n-1))||(globalco==(n-1))){
        update = dev_Along[globalind];
    }
    else{}
    if((globalro<n)&&(globalco<n))
        c[globalind] = update;
}

int main(int argc, char* argv[])
{
    int n;
    int t;
    n = atoi(argv[1]);
    t = atoi(argv[2]);

    int size = n*n;//number of elements
    double *c; // returned array
    //initiallize
    double *Along;
    Along = new double [size];
    c = new double [size];
    for (int i=0;i<size;i++){
        int ro = floor(i/n);
        int co = floor(i%n);
        Along[i] = pow((1+cos(2*ro)+sin(co)),2);
    }


    hipError_t cudaStatus = addWithCuda(c, size, n, Along, t);
    if (cudaStatus != hipSuccess) {
        cout<<"addWithCuda failed!"<<endl;
        return -1;
    }
    
    //here we get the c array then we can do the sum and check the certain element
    // examine
    double initial_sum = 0;
    double sumc = accumulate(c, c+size, initial_sum);
    cout<<"Sum: "<<sumc<<endl;
    cout<<"A(37,47): "<<c[37*n+47]<<endl;
    return 0;
}

// Helper function for using CUDA to add vectors in parallel.
hipError_t addWithCuda(double *c, unsigned int size, int n, double *Along, int t)
{
    double *dev_Along = 0;
    //dev_Along = new double[size];
    double *dev_c = 0;

    dev_c = new double[size];
    for (int i =0;i<size;i++){
        dev_c[i]=0;
    }
    hipError_t cudaStatus;
    hipEvent_t start, stop;
    float gpu_time = 0.0f;
    dim3 gridSize(ceil(n/32)+1,ceil(n/32)+1,1);
    dim3 blockSize(32,32,1);


    // Choose which GPU to run on, 0 if you have only one GPU
    // on-chip GPU does not count
    cudaStatus = hipSetDevice(0);
    if (cudaStatus != hipSuccess) {
        cout<<"hipSetDevice failed!"<<endl;
        goto Error;
    }
    // Malloc memory on GPU
    
    
    cudaStatus = hipMalloc((void**)&dev_Along, size * sizeof(double));
    if (cudaStatus != hipSuccess) {
        cout<<"hipMalloc failed!"<<endl;
        goto Error;
    }
    cudaStatus = hipMalloc((void**)&dev_c, size * sizeof(double));
    if (cudaStatus != hipSuccess) {
        cout<<"hipMalloc failed!"<<endl;
        goto Error;
    }
    
    // Copy memory from Host to Device
    cudaStatus = hipMemcpy(dev_Along, Along, size * sizeof(double), hipMemcpyHostToDevice);
    if (cudaStatus != hipSuccess) {
        cout<<"hipMemcpy failed!"<<endl;
        goto Error;
    }
    // Set up timing
    hipEventCreate(&start);
    hipEventCreate(&stop);
    hipEventRecord(start);
    // Launch a kernel on the GPU with one thread for each element
    //cout<<"allocating "<<gridSize<<" blocks, "<<blockSize<<" threads per block"<<endl;
    for (int tt=0;tt<t;tt++){
        addKernel<<<gridSize, blockSize>>>(dev_c, size, n,dev_Along,t);

        if (cudaStatus != hipSuccess) {
            cout<<"hipMemcpy failed!"<<endl;
            goto Error;
        }
        cudaStatus = hipMemcpy(dev_Along, dev_c, size * sizeof(double), hipMemcpyDeviceToDevice);
    }
    
    
    
      // Check for any errors launching the kernel
    cudaStatus = hipGetLastError();
    if (cudaStatus != hipSuccess) {
        cout<<"addKernel failed: "<<hipGetErrorString(cudaStatus)<<endl;
        goto Error;
    }
    // cudaDeviceSynchronize waits for the kernel to finish, and returns
    // any errors encountered during the launch.
    cudaStatus = hipDeviceSynchronize();
    if (cudaStatus != hipSuccess) {
        cout<<"hipDeviceSynchronize failed: "<<hipGetErrorString(cudaStatus)<<endl;
        goto Error;
    }
    // Close timing
    hipEventRecord(stop);
    hipEventSynchronize(stop);
    hipEventElapsedTime(&gpu_time, start, stop);
    cout<<"Time spent: "<<gpu_time<<"ms"<<endl;
    hipEventDestroy(start);
    hipEventDestroy(stop);

    // Copy memory from devide to host
    cudaStatus = hipMemcpy(c, dev_c, size * sizeof(double), hipMemcpyDeviceToHost);
    if (cudaStatus != hipSuccess) {
        cout<<"hipMemcpy failed!"<<endl;
        goto Error;
    }

Error:
    hipFree(dev_c);
    hipFree(dev_Along);


    return cudaStatus;
}

	.text
	.file	"hw4_4.hip"
                                        # Start of file scope inline assembly
	.globl	_ZSt21ios_base_library_initv

                                        # End of file scope inline assembly
	.globl	_Z24__device_stub__addKernelPdiiS_i # -- Begin function _Z24__device_stub__addKernelPdiiS_i
	.type	_Z24__device_stub__addKernelPdiiS_i,@function
_Z24__device_stub__addKernelPdiiS_i:    # @_Z24__device_stub__addKernelPdiiS_i
	.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	$128, %rsp
	.cfi_def_cfa_offset 176
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	12(%rsp), %rdi
	movl	%esi, (%rdi)
	leaq	8(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	4(%rsp), %rcx
	movl	%r8d, (%rcx)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	24(%rsp), %r12
	leaq	16(%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	$_Z9addKernelPdiiS_i, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$144, %rsp
	.cfi_adjust_cfa_offset -144
	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_end0:
	.size	_Z24__device_stub__addKernelPdiiS_i, .Lfunc_end0-_Z24__device_stub__addKernelPdiiS_i
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function main
.LCPI1_0:
	.quad	0x3ff0000000000000              # double 1
	.text
	.globl	main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$40, %rsp
	.cfi_def_cfa_offset 96
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%rsi, %r14
	movq	8(%rsi), %rdi
	callq	atoi
	movl	%eax, %ebx
	movq	16(%r14), %rdi
	callq	atoi
	movl	%eax, %ebp
	movl	%ebx, %r12d
	imull	%r12d, %r12d
	leaq	(,%r12,8), %r14
	movq	%r14, %rdi
	callq	_Znam
	movq	%rax, %r13
	movq	%r14, %rdi
	callq	_Znam
	movq	%rax, %r15
	testl	%ebx, %ebx
	je	.LBB1_5
# %bb.1:                                # %.lr.ph.preheader
	movl	%ebp, 20(%rsp)                  # 4-byte Spill
	movq	%r14, 24(%rsp)                  # 8-byte Spill
	movq	%r15, 32(%rsp)                  # 8-byte Spill
	cmpl	$2, %r12d
	movl	$1, %r15d
	cmovael	%r12d, %r15d
	xorl	%ebp, %ebp
.LBB1_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movl	%ebp, %eax
	cltd
	idivl	%ebx
	movl	%edx, %r14d
	addl	%eax, %eax
	xorps	%xmm0, %xmm0
	cvtsi2sd	%eax, %xmm0
	callq	cos
	addsd	.LCPI1_0(%rip), %xmm0
	movsd	%xmm0, 8(%rsp)                  # 8-byte Spill
	xorps	%xmm0, %xmm0
	cvtsi2sd	%r14d, %xmm0
	callq	sin
	addsd	8(%rsp), %xmm0                  # 8-byte Folded Reload
	mulsd	%xmm0, %xmm0
	movsd	%xmm0, (%r13,%rbp,8)
	incq	%rbp
	cmpq	%rbp, %r15
	jne	.LBB1_2
# %bb.3:                                # %._crit_edge
	movq	32(%rsp), %r15                  # 8-byte Reload
	movq	%r15, %rdi
	movl	%r12d, %esi
	movl	%ebx, %edx
	movq	%r13, %rcx
	movl	20(%rsp), %r8d                  # 4-byte Reload
	callq	_Z11addWithCudaPdjiS_i
	testl	%eax, %eax
	jne	.LBB1_4
# %bb.7:
	testl	%ebx, %ebx
	je	.LBB1_6
# %bb.8:                                # %.lr.ph.i.preheader
	xorpd	%xmm0, %xmm0
	xorl	%eax, %eax
	movq	24(%rsp), %rcx                  # 8-byte Reload
.LBB1_9:                                # %.lr.ph.i
                                        # =>This Inner Loop Header: Depth=1
	addsd	(%r15,%rax), %xmm0
	addq	$8, %rax
	cmpq	%rax, %rcx
	jne	.LBB1_9
	jmp	.LBB1_10
.LBB1_5:                                # %._crit_edge.thread
	movq	%r15, %rdi
	movl	%r12d, %esi
	movl	%ebx, %edx
	movq	%r13, %rcx
	movl	%ebp, %r8d
	callq	_Z11addWithCudaPdjiS_i
	testl	%eax, %eax
	je	.LBB1_6
.LBB1_4:
	movl	$_ZSt4cout, %r14d
	movl	$_ZSt4cout, %edi
	movl	$.L.str, %esi
	movl	$19, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$-1, %ebx
	jmp	.LBB1_11
.LBB1_6:
	xorpd	%xmm0, %xmm0
.LBB1_10:                               # %_ZSt10accumulateIPddET0_T_S2_S1_.exit
	movsd	%xmm0, 8(%rsp)                  # 8-byte Spill
	movl	$_ZSt4cout, %edi
	movl	$.L.str.1, %esi
	movl	$5, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movsd	8(%rsp), %xmm0                  # 8-byte Reload
                                        # xmm0 = mem[0],zero
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r14
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r14, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r14, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$_ZSt4cout, %edi
	movl	$.L.str.2, %esi
	movl	$10, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	leal	(%rbx,%rbx,8), %eax
	leal	(%rbx,%rax,4), %eax
	cltq
	movsd	376(%r15,%rax,8), %xmm0         # xmm0 = mem[0],zero
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r14
	xorl	%ebx, %ebx
.LBB1_11:
	movq	(%r14), %rax
	movq	-24(%rax), %rdi
	addq	%r14, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r14, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	%ebx, %eax
	addq	$40, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	main, .Lfunc_end1-main
	.cfi_endproc
                                        # -- End function
	.globl	_Z11addWithCudaPdjiS_i          # -- Begin function _Z11addWithCudaPdjiS_i
	.type	_Z11addWithCudaPdjiS_i,@function
_Z11addWithCudaPdjiS_i:                 # @_Z11addWithCudaPdjiS_i
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$56, %rsp
	.cfi_def_cfa_offset 112
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	%r8d, 4(%rsp)                   # 4-byte Spill
	movq	%rcx, %rbp
	movl	%edx, %r12d
	movl	%esi, %r13d
	movq	%rdi, 48(%rsp)                  # 8-byte Spill
	movq	$0, 16(%rsp)
	movl	%esi, %r14d
	shlq	$3, %r14
	movq	%r14, %rdi
	callq	_Znam
	movq	%rax, 8(%rsp)
	testl	%r13d, %r13d
	je	.LBB2_2
# %bb.1:                                # %.lr.ph.preheader
	movq	%rax, %rdi
	xorl	%esi, %esi
	movq	%r14, %rdx
	callq	memset@PLT
.LBB2_2:                                # %._crit_edge
	movl	$0, 36(%rsp)
	leal	31(%r12), %r15d
	testl	%r12d, %r12d
	cmovnsl	%r12d, %r15d
	xorl	%edi, %edi
	callq	hipSetDevice
	testl	%eax, %eax
	je	.LBB2_4
# %bb.3:
	movl	%eax, %ebx
	movl	$_ZSt4cout, %r14d
	movl	$_ZSt4cout, %edi
	movl	$.L.str.3, %esi
	movl	$20, %edx
	jmp	.LBB2_8
.LBB2_4:
	leaq	16(%rsp), %rdi
	movq	%r14, %rsi
	callq	hipMalloc
	testl	%eax, %eax
	jne	.LBB2_6
# %bb.5:
	leaq	8(%rsp), %rdi
	movq	%r14, %rsi
	callq	hipMalloc
	testl	%eax, %eax
	je	.LBB2_10
.LBB2_6:
	movl	%eax, %ebx
	movl	$_ZSt4cout, %r14d
	movl	$_ZSt4cout, %edi
	movl	$.L.str.4, %esi
.LBB2_7:
	movl	$17, %edx
.LBB2_8:
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%r14, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
.LBB2_9:
	movq	8(%rsp), %rdi
	callq	hipFree
	movq	16(%rsp), %rdi
	callq	hipFree
	movl	%ebx, %eax
	addq	$56, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.LBB2_10:
	.cfi_def_cfa_offset 112
	movq	16(%rsp), %rdi
	movq	%rbp, %rsi
	movq	%r14, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	testl	%eax, %eax
	je	.LBB2_12
# %bb.11:
	movl	%eax, %ebx
	movl	$_ZSt4cout, %r14d
	movl	$_ZSt4cout, %edi
	movl	$.L.str.5, %esi
	jmp	.LBB2_7
.LBB2_12:
	leaq	40(%rsp), %rbx
	movq	%rbx, %rdi
	callq	hipEventCreate
	leaq	24(%rsp), %rdi
	callq	hipEventCreate
	movq	(%rbx), %rdi
	xorl	%ebx, %ebx
	xorl	%esi, %esi
	callq	hipEventRecord
	cmpl	$0, 4(%rsp)                     # 4-byte Folded Reload
	jle	.LBB2_18
# %bb.13:                               # %.lr.ph71.preheader
	sarl	$5, %r15d
	incl	%r15d
	movq	%r15, %rbp
	shlq	$32, %rbp
	orq	%r15, %rbp
	movl	4(%rsp), %r15d                  # 4-byte Reload
.LBB2_14:                               # %.lr.ph71
                                        # =>This Inner Loop Header: Depth=1
	movq	%rbp, %rdi
	movl	$1, %esi
	movabsq	$137438953504, %rdx             # imm = 0x2000000020
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB2_16
# %bb.15:                               #   in Loop: Header=BB2_14 Depth=1
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rcx
	movl	%r13d, %esi
	movl	%r12d, %edx
	movl	4(%rsp), %r8d                   # 4-byte Reload
	callq	_Z24__device_stub__addKernelPdiiS_i
.LBB2_16:                               #   in Loop: Header=BB2_14 Depth=1
	testl	%ebx, %ebx
	jne	.LBB2_20
# %bb.17:                               #   in Loop: Header=BB2_14 Depth=1
	movq	16(%rsp), %rdi
	movq	8(%rsp), %rsi
	movq	%r14, %rdx
	movl	$3, %ecx
	callq	hipMemcpy
	movl	%eax, %ebx
	decl	%r15d
	jne	.LBB2_14
.LBB2_18:                               # %._crit_edge72
	callq	hipGetLastError
	testl	%eax, %eax
	je	.LBB2_22
# %bb.19:
	movl	%eax, %ebx
	movl	$_ZSt4cout, %edi
	movl	$.L.str.6, %esi
	movl	$18, %edx
	jmp	.LBB2_24
.LBB2_20:
	movl	$_ZSt4cout, %edi
	movl	$.L.str.5, %esi
	movl	$17, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	jmp	.LBB2_21
.LBB2_22:
	callq	hipDeviceSynchronize
	testl	%eax, %eax
	je	.LBB2_25
# %bb.23:
	movl	%eax, %ebx
	movl	$_ZSt4cout, %edi
	movl	$.L.str.7, %esi
	movl	$29, %edx
.LBB2_24:
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	%ebx, %edi
	callq	hipGetErrorString
	movl	$_ZSt4cout, %edi
	movq	%rax, %rsi
	callq	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc
	movq	%rax, %rdi
.LBB2_21:
	callq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_
	jmp	.LBB2_9
.LBB2_25:
	movq	24(%rsp), %rdi
	xorl	%ebx, %ebx
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	24(%rsp), %rdi
	callq	hipEventSynchronize
	movq	40(%rsp), %rsi
	movq	24(%rsp), %rdx
	leaq	36(%rsp), %r15
	movq	%r15, %rdi
	callq	hipEventElapsedTime
	movl	$_ZSt4cout, %edi
	movl	$.L.str.8, %esi
	movl	$12, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	cvtss2sd	(%r15), %xmm0
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r15
	movl	$.L.str.9, %esi
	movl	$2, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	%r15, %rdi
	callq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_
	movq	40(%rsp), %rdi
	callq	hipEventDestroy
	movq	24(%rsp), %rdi
	callq	hipEventDestroy
	movq	8(%rsp), %rsi
	movq	48(%rsp), %rdi                  # 8-byte Reload
	movq	%r14, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	testl	%eax, %eax
	je	.LBB2_9
# %bb.26:
	movl	%eax, %ebp
	movl	$_ZSt4cout, %edi
	movl	$.L.str.5, %esi
	movl	$17, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	callq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_
	movl	%ebp, %ebx
	jmp	.LBB2_9
.Lfunc_end2:
	.size	_Z11addWithCudaPdjiS_i, .Lfunc_end2-_Z11addWithCudaPdjiS_i
	.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	$_Z9addKernelPdiiS_i, %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	_Z9addKernelPdiiS_i,@object     # @_Z9addKernelPdiiS_i
	.section	.rodata,"a",@progbits
	.globl	_Z9addKernelPdiiS_i
	.p2align	3, 0x0
_Z9addKernelPdiiS_i:
	.quad	_Z24__device_stub__addKernelPdiiS_i
	.size	_Z9addKernelPdiiS_i, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"addWithCuda failed!"
	.size	.L.str, 20

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"Sum: "
	.size	.L.str.1, 6

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"A(37,47): "
	.size	.L.str.2, 11

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"hipSetDevice failed!"
	.size	.L.str.3, 21

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"hipMalloc failed!"
	.size	.L.str.4, 18

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"hipMemcpy failed!"
	.size	.L.str.5, 18

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"addKernel failed: "
	.size	.L.str.6, 19

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	"hipDeviceSynchronize failed: "
	.size	.L.str.7, 30

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	"Time spent: "
	.size	.L.str.8, 13

	.type	.L.str.9,@object                # @.str.9
.L.str.9:
	.asciz	"ms"
	.size	.L.str.9, 3

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z9addKernelPdiiS_i"
	.size	.L__unnamed_1, 20

	.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 _Z24__device_stub__addKernelPdiiS_i
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z9addKernelPdiiS_i
	.addrsig_sym _ZSt4cout
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z9addKernelPdiiS_i     ; -- Begin function _Z9addKernelPdiiS_i
	.globl	_Z9addKernelPdiiS_i
	.p2align	8
	.type	_Z9addKernelPdiiS_i,@function
_Z9addKernelPdiiS_i:                    ; @_Z9addKernelPdiiS_i
; %bb.0:
	s_clause 0x2
	s_load_b32 s2, s[0:1], 0x2c
	s_load_b32 s9, s[0:1], 0xc
	s_load_b64 s[10:11], s[0:1], 0x10
	v_and_b32_e32 v3, 0x3ff, v0
	v_bfe_u32 v0, v0, 10, 10
	s_mov_b32 s4, exec_lo
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	s_lshr_b32 s3, s2, 16
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s14, s2, v[3:4]
	v_mad_u64_u32 v[2:3], null, s15, s3, v[0:1]
	v_mul_lo_u32 v5, v1, s9
	s_mov_b32 s3, 0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_min_i32_e32 v0, v1, v2
	v_max_i32_e32 v10, v1, v2
	v_cmp_gt_i32_e64 s2, 1, v0
	v_cmpx_lt_i32_e32 0, v0
; %bb.1:
	s_add_i32 s5, s9, -1
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_3) | instid1(SALU_CYCLE_1)
	s_and_not1_b32 s2, s2, exec_lo
	v_cmp_le_i32_e32 vcc_lo, s5, v10
	s_mov_b32 s3, exec_lo
	s_and_b32 s5, vcc_lo, exec_lo
	s_or_b32 s2, s2, s5
; %bb.2:                                ; %Flow96
	s_or_b32 exec_lo, exec_lo, s4
	v_add_nc_u32_e32 v0, v5, v2
                                        ; implicit-def: $vgpr3_vgpr4
	s_and_saveexec_b32 s4, s2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 s4, exec_lo, s4
	s_cbranch_execz .LBB0_8
; %bb.3:
	v_cmp_eq_u32_e32 vcc_lo, 0, v1
	v_cmp_eq_u32_e64 s2, 0, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 s5, vcc_lo, s2
	s_xor_b32 s2, s5, -1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_saveexec_b32 s6, s2
; %bb.4:
	s_add_i32 s2, s9, -1
	s_and_not1_b32 s5, s5, exec_lo
	v_cmp_eq_u32_e32 vcc_lo, s2, v1
	v_cmp_eq_u32_e64 s2, s2, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 s2, vcc_lo, s2
	s_and_b32 s2, s2, exec_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 s5, s5, s2
; %bb.5:                                ; %Flow
	s_or_b32 exec_lo, exec_lo, s6
                                        ; implicit-def: $vgpr3_vgpr4
	s_and_saveexec_b32 s2, s5
	s_cbranch_execz .LBB0_7
; %bb.6:
	v_ashrrev_i32_e32 v1, 31, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[3:4], 3, v[0:1]
	v_add_co_u32 v3, vcc_lo, s10, v3
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v4, vcc_lo, s11, v4, vcc_lo
	global_load_b64 v[3:4], v[3:4], off
.LBB0_7:                                ; %Flow93
	s_or_b32 exec_lo, exec_lo, s2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 s3, s3, exec_lo
.LBB0_8:                                ; %Flow97
	s_or_b32 exec_lo, exec_lo, s4
	s_and_saveexec_b32 s16, s3
	s_cbranch_execz .LBB0_15
; %bb.9:
	v_add_nc_u32_e32 v1, s9, v5
	s_lshl_b32 s2, s9, 1
	v_ashrrev_i32_e32 v5, 31, v2
	s_mov_b32 s17, 0
	s_mov_b64 s[12:13], 1
	s_waitcnt vmcnt(0)
	v_subrev_nc_u32_e32 v3, s2, v1
	v_ashrrev_i32_e32 v6, 31, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v4, 31, v3
	v_add_co_u32 v3, vcc_lo, v3, v2
	v_add_co_ci_u32_e32 v4, vcc_lo, v4, v5, vcc_lo
	v_add_co_u32 v1, vcc_lo, v1, v2
	v_add_co_ci_u32_e32 v2, vcc_lo, v6, v5, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[3:4], 3, v[3:4]
	v_lshlrev_b64 v[1:2], 3, v[1:2]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v6, vcc_lo, s10, v3
	v_add_co_ci_u32_e32 v7, vcc_lo, s11, v4, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v1, vcc_lo, s10, v1
	v_add_co_ci_u32_e32 v2, vcc_lo, s11, v2, vcc_lo
	s_clause 0x3
	global_load_b64 v[4:5], v[6:7], off offset:8
	global_load_b64 v[6:7], v[6:7], off offset:-8
	global_load_b64 v[8:9], v[1:2], off offset:8
	global_load_b64 v[2:3], v[1:2], off offset:-8
.LBB0_10:                               ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_12 Depth 2
	s_cmp_gt_u32 s17, 2
	s_cbranch_scc1 .LBB0_13
; %bb.11:                               ; %.lr.ph
                                        ;   in Loop: Header=BB0_10 Depth=1
	s_mov_b64 s[14:15], s[12:13]
.LBB0_12:                               ;   Parent Loop BB0_10 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_cmp_eq_u32 s17, 1
	s_cselect_b32 vcc_lo, -1, 0
	s_cmp_eq_u32 s17, 2
	s_waitcnt vmcnt(0)
	v_cndmask_b32_e32 v1, v9, v3, vcc_lo
	s_cselect_b32 s2, -1, 0
	s_cmp_eq_u32 s17, 3
	s_cselect_b32 s3, -1, 0
	s_delay_alu instid0(VALU_DEP_1)
	v_cndmask_b32_e64 v1, v1, v5, s2
	s_cmp_eq_u32 s14, 1
	s_cselect_b32 s4, -1, 0
	s_cmp_eq_u32 s14, 2
	v_cndmask_b32_e64 v13, v9, v3, s4
	v_cndmask_b32_e64 v14, v8, v2, s4
	v_cndmask_b32_e64 v12, v1, v7, s3
	v_cndmask_b32_e32 v11, v8, v2, vcc_lo
	s_cselect_b32 s5, -1, 0
	s_cmp_eq_u32 s14, 3
	v_cndmask_b32_e64 v1, v13, v5, s5
	v_cndmask_b32_e64 v13, v14, v4, s5
	v_cndmask_b32_e64 v11, v11, v4, s2
	s_cselect_b32 s6, -1, 0
	s_cmp_eq_u32 s17, 0
	v_cndmask_b32_e64 v14, v1, v7, s6
	v_cndmask_b32_e64 v13, v13, v6, s6
	v_cndmask_b32_e64 v11, v11, v6, s3
	s_cselect_b32 s8, -1, 0
	s_cmp_eq_u32 s14, 0
	v_cndmask_b32_e64 v1, v7, v14, s3
	v_cndmask_b32_e64 v15, v5, v14, s2
	v_cmp_gt_f64_e64 s7, v[11:12], v[13:14]
	v_cndmask_b32_e32 v16, v3, v14, vcc_lo
	v_cndmask_b32_e64 v14, v9, v14, s8
	v_cndmask_b32_e64 v17, v6, v13, s3
	v_cndmask_b32_e64 v18, v4, v13, s2
	v_cndmask_b32_e32 v19, v2, v13, vcc_lo
	v_cndmask_b32_e64 v13, v8, v13, s8
	s_cselect_b32 vcc_lo, -1, 0
	v_cndmask_b32_e64 v1, v1, v12, s6
	v_cndmask_b32_e64 v16, v16, v12, s4
	v_cndmask_b32_e64 v15, v15, v12, s5
	v_dual_cndmask_b32 v12, v14, v12 :: v_dual_cndmask_b32 v13, v13, v11
	v_cndmask_b32_e64 v17, v17, v11, s6
	v_cndmask_b32_e64 v14, v18, v11, s5
	v_cndmask_b32_e64 v11, v19, v11, s4
	s_add_u32 s2, s14, 1
	s_addc_u32 s3, s15, 0
	s_cmp_gt_u32 s14, 2
	s_mov_b64 s[14:15], s[2:3]
	v_cndmask_b32_e64 v7, v7, v1, s7
	v_cndmask_b32_e64 v6, v6, v17, s7
	v_cndmask_b32_e64 v5, v5, v15, s7
	v_cndmask_b32_e64 v4, v4, v14, s7
	v_cndmask_b32_e64 v3, v3, v16, s7
	v_cndmask_b32_e64 v2, v2, v11, s7
	v_cndmask_b32_e64 v9, v9, v12, s7
	v_cndmask_b32_e64 v8, v8, v13, s7
	s_cbranch_scc0 .LBB0_12
.LBB0_13:                               ; %.loopexit
                                        ;   in Loop: Header=BB0_10 Depth=1
	s_add_i32 s17, s17, 1
	s_add_u32 s12, s12, 1
	s_addc_u32 s13, s13, 0
	s_cmp_lg_u32 s17, 4
	s_cbranch_scc1 .LBB0_10
; %bb.14:
	v_ashrrev_i32_e32 v1, 31, v0
	s_waitcnt vmcnt(3)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[4:5], 3, v[0:1]
	v_add_co_u32 v4, vcc_lo, s10, v4
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v5, vcc_lo, s11, v5, vcc_lo
	global_load_b64 v[4:5], v[4:5], off
	s_waitcnt vmcnt(0)
	v_add_f64 v[3:4], v[2:3], v[4:5]
.LBB0_15:                               ; %Flow98
	s_or_b32 exec_lo, exec_lo, s16
	s_delay_alu instid0(SALU_CYCLE_1)
	s_mov_b32 s2, exec_lo
	v_cmpx_gt_i32_e64 s9, v10
	s_cbranch_execz .LBB0_17
; %bb.16:
	s_load_b64 s[0:1], s[0:1], 0x0
	v_ashrrev_i32_e32 v1, 31, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 3, v[0:1]
	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
	s_waitcnt vmcnt(0)
	global_store_b64 v[0:1], v[3:4], off
.LBB0_17:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z9addKernelPdiiS_i
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 288
		.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 20
		.amdhsa_next_free_sgpr 18
		.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	_Z9addKernelPdiiS_i, .Lfunc_end0-_Z9addKernelPdiiS_i
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 1028
; NumSgprs: 20
; NumVgprs: 20
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 2
; NumSGPRsForWavesPerEU: 20
; NumVGPRsForWavesPerEU: 20
; Occupancy: 16
; WaveLimiterHint : 0
; 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	__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
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         36
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         44
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         46
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         48
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         50
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         52
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         54
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         96
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 288
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z9addKernelPdiiS_i
    .private_segment_fixed_size: 0
    .sgpr_count:     20
    .sgpr_spill_count: 0
    .symbol:         _Z9addKernelPdiiS_i.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     20
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
