/*
******************************************************
This file is the single GPU version of 2D Heat Equation 
using CUDA programming model. This implementation is based
on the CPU version from
http://www.many-core.group.cam.ac.uk/archive/CUDAcourse09/

Permission to use, copy, distribute and modify this software for any 
purpose with or without fee is hereby granted. This software is        
provided "as is" without express or implied warranty. 

Send comments or suggestions for this OpenACC version to
            rxu6@uh.edu, schandra@udel.edu

Authors: Rengan Xu, Sunita Chandrasekaran

May 26th, 2016
******************************************************
*/

#include <stdlib.h>
#include <stdio.h>
#include <string.h>
#include <sys/time.h>
#include <fcntl.h>
#include <unistd.h>
#include <cuda_runtime.h>

// I2D to index into a linear memory space from a 2D array index pair
#define I2D(ni, i, j) ((i) + (ni)*(j))
#define THREADS 128

// kernel to update temperatures - CPU version
__global__ void step_kernel_cpu(int ni, 
                     int nj,
                     double tfac, 
                     double *temp_in,
                     double *temp_out) {
    int i, j, i00, im10, ip10, i0m1, i0p1;
    double d2tdx2, d2tdy2;

    j = blockIdx.y + 1;
    i = threadIdx.x + blockIdx.x*blockDim.x + 1;

    // loop over all points in domain (not boundary points)
    for (j=j; j < nj-1; j+=gridDim.y) {
        for (i=i; i < ni-1; i+=blockDim.x*gridDim.x) {
	    // find indices into linear memory for central point and neighbours
            i00 = I2D(ni, i, j);
            im10 = I2D(ni, i-1, j);
            ip10 = I2D(ni, i+1, j);
            i0m1 = I2D(ni, i, j-1);
            i0p1 = I2D(ni, i, j+1);

	    // evaluate derivatives
            d2tdx2 = temp_in[im10] - 2*temp_in[i00] + temp_in[ip10];
            d2tdy2 = temp_in[i0m1] - 2*temp_in[i00] + temp_in[i0p1];

	    // update temperatures
            temp_out[i00] = temp_in[i00] + tfac*(d2tdx2 + d2tdy2);
        }
    }

}


int main(int argc, char *argv[]) 
{
    
    if(argc < 5)
    {
        printf("Usage: %s <ni> <ni> <nstep> <output file>\n", argv[0]);
        exit(0);
    }

    int ni, nj, nstep;
    double tfac, *temp1_h, *temp2_h, *temp_tmp, *temp1_d, *temp2_d;
    int i, j, i2d, istep;
    double temp_bl, temp_br, temp_tl, temp_tr;
	struct timeval tim;
	double start, end;
	int fd;
   
    // domain size and number of timesteps (iterations)
    ni = atoi(argv[1]);
    nj = atoi(argv[2]);
    nstep = atoi(argv[3]);
    
    // allocate temperature array on host
    temp1_h = (double *)malloc(sizeof(double)*(ni+2)*(nj+2));
    temp2_h = (double *)malloc(sizeof(double)*(ni+2)*(nj+2));

    // initial temperature in interior
    for (j=1; j < nj+1; j++) {
        for (i=1; i < ni+1; i++) {
            i2d = i + (ni+2)*j;
            temp1_h[i2d] = 0.0;
        }
    }
    
    // initial temperature on boundaries - set corners
    temp_bl = 200.0f;
    temp_br = 300.0f;
    temp_tl = 200.0f;
    temp_tr = 300.0f;

    // set edges by linear interpolation from corners
    for (i=0; i < ni+2; i++) {
        // bottom
        j = 0;
        i2d = i + (ni+2)*j;
        temp1_h[i2d] = temp_bl + (temp_br-temp_bl)*(double)i/(double)(ni+1);

        // top
        j = nj+1;
        i2d = i + (ni+2)*j;
        temp1_h[i2d] = temp_tl + (temp_tr-temp_tl)*(double)i/(double)(ni+1);
    }

    for (j=0; j < nj+2; j++) {
        // left
        i = 0;
        i2d = i + (ni+2)*j;
        temp1_h[i2d] = temp_bl + (temp_tl-temp_bl)*(double)j/(double)(nj+1);

        // right
        i = ni+1;
        i2d = i + (ni+2)*j;
        temp1_h[i2d] = temp_br + (temp_tr-temp_br)*(double)j/(double)(nj+1);
    }

    // duplicate temeperature array on host
    memcpy(temp2_h, temp1_h, sizeof(double)*(ni+2)*(nj+2));
    
	gettimeofday(&tim, NULL);
	start = tim.tv_sec + (tim.tv_usec/1000000.0);
	
    // allocate temperature arrays on device
    cudaMalloc((void **)&temp1_d, sizeof(double)*(ni+2)*(nj+2));
    cudaMalloc((void **)&temp2_d, sizeof(double)*(ni+2)*(nj+2));
    
    // transfer temperature array from host to device
    cudaMemcpy((void *)temp1_d, (void *)temp1_h, sizeof(double)*(ni+2)*(nj+2),
               cudaMemcpyHostToDevice);
    cudaMemcpy((void *)temp2_d, (void *)temp1_h, sizeof(double)*(ni+2)*(nj+2),
               cudaMemcpyHostToDevice);

    tfac = 0.2;
    
    // main iteration loop  

    for (istep=0; istep < nstep; istep++) {
            dim3 grid_dim((ni+THREADS-1)/THREADS, nj, 1); 
            dim3 block_dim(THREADS, 1, 1); 
            step_kernel_cpu<<<grid_dim, block_dim>>>(ni+2, nj+2, tfac, temp1_d, temp2_d);

	    // swap the temp pointers
            temp_tmp = temp1_d;
            temp1_d = temp2_d;
            temp2_d = temp_tmp;
            
    } 
    
    cudaMemcpy((void *)temp1_h, (void *)temp1_d, sizeof(double)*(ni+2)*(nj+2),
                   cudaMemcpyDeviceToHost);
	
    gettimeofday(&tim, NULL);
    end = tim.tv_sec + (tim.tv_usec/1000000.0);
    printf("Time for computing: %.2f s\n",end-start);


    // output temp1 to a file
    
	fd = creat(argv[4], 00666);
	fd = open(argv[4], O_WRONLY);
	write(fd, temp1_h, (size_t)(ni+2)*(nj+2)*sizeof(double));
	close(fd);
    
/*
    FILE *fp;
    fp = fopen(filename, "w");
    fprintf(fp, "%d %d\n", ni, nj);
    for (j=0; j < nj; j++) {
        for (i=0; i < ni; i++) {
            fprintf(fp, "%.4f\n", j, i, temp1_h[i + ni*j]);
        }
    }
    fclose(fp);
*/
}



	.file	"tmpxft_002a0b7b_00000000-6_cuda_heat_single.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2047:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2047:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z40__device_stub__Z15step_kernel_cpuiidPdS_iidPdS_
	.type	_Z40__device_stub__Z15step_kernel_cpuiidPdS_iidPdS_, @function
_Z40__device_stub__Z15step_kernel_cpuiidPdS_iidPdS_:
.LFB2069:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movl	%edi, 28(%rsp)
	leaq	48(%rsp), %rdi
	movl	%esi, 24(%rsp)
	leaq	60(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%rsp)
	leaq	40(%rsp), %rcx
	movsd	%xmm0, 16(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	28(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	movq	%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	.L2
	pushq	40(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z15step_kernel_cpuiidPdS_(%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
.L2:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2069:
	.size	_Z40__device_stub__Z15step_kernel_cpuiidPdS_iidPdS_, .-_Z40__device_stub__Z15step_kernel_cpuiidPdS_iidPdS_
	.globl	_Z15step_kernel_cpuiidPdS_
	.type	_Z15step_kernel_cpuiidPdS_, @function
_Z15step_kernel_cpuiidPdS_:
.LFB2070:
	.cfi_startproc
	endbr64
	jmp	_Z40__device_stub__Z15step_kernel_cpuiidPdS_iidPdS_
	.cfi_endproc
.LFE2070:
	.size	_Z15step_kernel_cpuiidPdS_, .-_Z15step_kernel_cpuiidPdS_
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"Usage: %s <ni> <ni> <nstep> <output file>\n"
.LC7:
	.string	"Time for computing: %.2f s\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2044:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	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
	movq	%rsi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	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
	cmpl	$4, %edi
	jg	.L9
	movq	(%rsi), %rdx
	movl	$2, %edi
	leaq	.LC0(%rip), %rsi
	call	__printf_chk@PLT
	xorl	%edi, %edi
	call	exit@PLT
.L9:
	movq	8(%rsi), %rdi
	call	atoi@PLT
	movq	16(%r12), %rdi
	movl	%eax, %ebp
	call	atoi@PLT
	movq	24(%r12), %rdi
	leal	2(%rbp), %r14d
	movl	%eax, %r13d
	call	atoi@PLT
	movslq	%r14d, %rcx
	movl	%eax, %r15d
	leal	2(%r13), %eax
	movq	%rcx, 8(%rsp)
	movl	%eax, 24(%rsp)
	cltq
	imulq	%rcx, %rax
	salq	$3, %rax
	movq	%rax, %rdi
	movq	%rax, (%rsp)
	call	malloc@PLT
	movq	(%rsp), %rdi
	movq	%rax, %rbx
	call	malloc@PLT
	movq	8(%rsp), %rcx
	movl	%r14d, %edx
	movl	$1, %esi
	movq	%rax, %rdi
.L10:
	cmpl	%r13d, %esi
	jg	.L12
	movslq	%edx, %rax
	leaq	(%rbx,%rax,8), %r8
	xorl	%eax, %eax
.L13:
	incq	%rax
	cmpl	%eax, %ebp
	jl	.L23
	movq	$0x000000000, (%r8,%rax,8)
	jmp	.L13
.L23:
	incl	%esi
	addl	%r14d, %edx
	jmp	.L10
.L12:
	leal	1(%r13), %edx
	movsd	.LC2(%rip), %xmm1
	movsd	.LC3(%rip), %xmm2
	leal	1(%rbp), %esi
	movl	%edx, %eax
	imull	%r14d, %eax
	cltq
	leaq	(%rbx,%rax,8), %r8
	xorl	%eax, %eax
.L14:
	cmpl	%eax, %esi
	jl	.L24
	cvtsi2sdl	%eax, %xmm0
	cvtsi2sdl	%esi, %xmm3
	mulsd	%xmm1, %xmm0
	divsd	%xmm3, %xmm0
	addsd	%xmm2, %xmm0
	movsd	%xmm0, (%rbx,%rax,8)
	movsd	%xmm0, (%r8,%rax,8)
	incq	%rax
	jmp	.L14
.L24:
	movsd	.LC3(%rip), %xmm2
	movq	%rbx, %rsi
	xorl	%eax, %eax
	xorps	%xmm1, %xmm1
	movsd	.LC4(%rip), %xmm3
	salq	$3, %rcx
	movslq	%ebp, %r8
.L16:
	cmpl	%eax, %edx
	jl	.L25
	cvtsi2sdl	%eax, %xmm0
	incl	%eax
	cvtsi2sdl	%edx, %xmm4
	mulsd	%xmm1, %xmm0
	divsd	%xmm4, %xmm0
	movaps	%xmm0, %xmm4
	addsd	%xmm3, %xmm0
	addsd	%xmm2, %xmm4
	movsd	%xmm4, (%rsi)
	movsd	%xmm0, 8(%rsi,%r8,8)
	addq	%rcx, %rsi
	jmp	.L16
.L25:
	movq	(%rsp), %rcx
	movq	%rbx, %rsi
	movq	%rcx, %rdx
	call	__memcpy_chk@PLT
	leaq	72(%rsp), %rax
	xorl	%esi, %esi
	movq	%rax, %rdi
	movq	%rax, 8(%rsp)
	call	gettimeofday@PLT
	cvtsi2sdq	72(%rsp), %xmm1
	movq	(%rsp), %rsi
	leaq	32(%rsp), %rdi
	cvtsi2sdq	80(%rsp), %xmm0
	divsd	.LC5(%rip), %xmm0
	addsd	%xmm1, %xmm0
	movsd	%xmm0, 16(%rsp)
	call	cudaMalloc@PLT
	movq	(%rsp), %rsi
	leaq	40(%rsp), %rdi
	call	cudaMalloc@PLT
	movq	(%rsp), %rdx
	movq	32(%rsp), %rdi
	movq	%rbx, %rsi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movq	(%rsp), %rdx
	movq	40(%rsp), %rdi
	movq	%rbx, %rsi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	leal	127(%rbp), %eax
	movl	$128, %ecx
	xorl	%ebp, %ebp
	cltd
	idivl	%ecx
	movl	%eax, 28(%rsp)
.L18:
	cmpl	%r15d, %ebp
	jge	.L26
	movl	28(%rsp), %eax
	movl	%r13d, 52(%rsp)
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$33554433, %edx
	movl	$1, %ecx
	movl	$1, %esi
	movl	%eax, 48(%rsp)
	movq	48(%rsp), %rdi
	salq	$7, %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L19
	movq	40(%rsp), %rcx
	movq	32(%rsp), %rdx
	movl	%r14d, %edi
	movsd	.LC6(%rip), %xmm0
	movl	24(%rsp), %esi
	call	_Z40__device_stub__Z15step_kernel_cpuiidPdS_iidPdS_
.L19:
	movq	32(%rsp), %rax
	movq	40(%rsp), %rdx
	incl	%ebp
	movq	%rdx, 32(%rsp)
	movq	%rax, 40(%rsp)
	jmp	.L18
.L26:
	movq	(%rsp), %rdx
	movq	32(%rsp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdi
	call	cudaMemcpy@PLT
	movq	8(%rsp), %rdi
	xorl	%esi, %esi
	call	gettimeofday@PLT
	cvtsi2sdq	72(%rsp), %xmm1
	movl	$2, %edi
	movb	$1, %al
	cvtsi2sdq	80(%rsp), %xmm0
	divsd	.LC5(%rip), %xmm0
	leaq	.LC7(%rip), %rsi
	addsd	%xmm1, %xmm0
	subsd	16(%rsp), %xmm0
	call	__printf_chk@PLT
	movq	32(%r12), %rdi
	movl	$438, %esi
	call	creat@PLT
	movq	32(%r12), %rdi
	movl	$1, %esi
	call	__open_2@PLT
	movq	(%rsp), %rdx
	movq	%rbx, %rsi
	movl	%eax, %ebp
	movl	%eax, %edi
	call	write@PLT
	movl	%ebp, %edi
	call	close@PLT
	movq	88(%rsp), %rax
	subq	%fs:40, %rax
	je	.L21
	call	__stack_chk_fail@PLT
.L21:
	addq	$104, %rsp
	.cfi_def_cfa_offset 56
	xorl	%eax, %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
.LFE2044:
	.size	main, .-main
	.section	.rodata.str1.1
.LC8:
	.string	"_Z15step_kernel_cpuiidPdS_"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2072:
	.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	.LC8(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z15step_kernel_cpuiidPdS_(%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
.LFE2072:
	.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	1079574528
	.align 8
.LC3:
	.long	0
	.long	1080623104
	.align 8
.LC4:
	.long	0
	.long	1081262080
	.align 8
.LC5:
	.long	0
	.long	1093567616
	.align 8
.LC6:
	.long	-1717986918
	.long	1070176665
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:

/*
******************************************************
This file is the single GPU version of 2D Heat Equation 
using CUDA programming model. This implementation is based
on the CPU version from
http://www.many-core.group.cam.ac.uk/archive/CUDAcourse09/

Permission to use, copy, distribute and modify this software for any 
purpose with or without fee is hereby granted. This software is        
provided "as is" without express or implied warranty. 

Send comments or suggestions for this OpenACC version to
            rxu6@uh.edu, schandra@udel.edu

Authors: Rengan Xu, Sunita Chandrasekaran

May 26th, 2016
******************************************************
*/

#include <stdlib.h>
#include <stdio.h>
#include <string.h>
#include <sys/time.h>
#include <fcntl.h>
#include <unistd.h>
#include <hip/hip_runtime.h>

// I2D to index into a linear memory space from a 2D array index pair
#define I2D(ni, i, j) ((i) + (ni)*(j))
#define THREADS 128

// kernel to update temperatures - CPU version
__global__ void step_kernel_cpu(int ni, 
                     int nj,
                     double tfac, 
                     double *temp_in,
                     double *temp_out) {
    int i, j, i00, im10, ip10, i0m1, i0p1;
    double d2tdx2, d2tdy2;

    j = blockIdx.y + 1;
    i = threadIdx.x + blockIdx.x*blockDim.x + 1;

    // loop over all points in domain (not boundary points)
    for (j=j; j < nj-1; j+=gridDim.y) {
        for (i=i; i < ni-1; i+=blockDim.x*gridDim.x) {
	    // find indices into linear memory for central point and neighbours
            i00 = I2D(ni, i, j);
            im10 = I2D(ni, i-1, j);
            ip10 = I2D(ni, i+1, j);
            i0m1 = I2D(ni, i, j-1);
            i0p1 = I2D(ni, i, j+1);

	    // evaluate derivatives
            d2tdx2 = temp_in[im10] - 2*temp_in[i00] + temp_in[ip10];
            d2tdy2 = temp_in[i0m1] - 2*temp_in[i00] + temp_in[i0p1];

	    // update temperatures
            temp_out[i00] = temp_in[i00] + tfac*(d2tdx2 + d2tdy2);
        }
    }

}


int main(int argc, char *argv[]) 
{
    
    if(argc < 5)
    {
        printf("Usage: %s <ni> <ni> <nstep> <output file>\n", argv[0]);
        exit(0);
    }

    int ni, nj, nstep;
    double tfac, *temp1_h, *temp2_h, *temp_tmp, *temp1_d, *temp2_d;
    int i, j, i2d, istep;
    double temp_bl, temp_br, temp_tl, temp_tr;
	struct timeval tim;
	double start, end;
	int fd;
   
    // domain size and number of timesteps (iterations)
    ni = atoi(argv[1]);
    nj = atoi(argv[2]);
    nstep = atoi(argv[3]);
    
    // allocate temperature array on host
    temp1_h = (double *)malloc(sizeof(double)*(ni+2)*(nj+2));
    temp2_h = (double *)malloc(sizeof(double)*(ni+2)*(nj+2));

    // initial temperature in interior
    for (j=1; j < nj+1; j++) {
        for (i=1; i < ni+1; i++) {
            i2d = i + (ni+2)*j;
            temp1_h[i2d] = 0.0;
        }
    }
    
    // initial temperature on boundaries - set corners
    temp_bl = 200.0f;
    temp_br = 300.0f;
    temp_tl = 200.0f;
    temp_tr = 300.0f;

    // set edges by linear interpolation from corners
    for (i=0; i < ni+2; i++) {
        // bottom
        j = 0;
        i2d = i + (ni+2)*j;
        temp1_h[i2d] = temp_bl + (temp_br-temp_bl)*(double)i/(double)(ni+1);

        // top
        j = nj+1;
        i2d = i + (ni+2)*j;
        temp1_h[i2d] = temp_tl + (temp_tr-temp_tl)*(double)i/(double)(ni+1);
    }

    for (j=0; j < nj+2; j++) {
        // left
        i = 0;
        i2d = i + (ni+2)*j;
        temp1_h[i2d] = temp_bl + (temp_tl-temp_bl)*(double)j/(double)(nj+1);

        // right
        i = ni+1;
        i2d = i + (ni+2)*j;
        temp1_h[i2d] = temp_br + (temp_tr-temp_br)*(double)j/(double)(nj+1);
    }

    // duplicate temeperature array on host
    memcpy(temp2_h, temp1_h, sizeof(double)*(ni+2)*(nj+2));
    
	gettimeofday(&tim, NULL);
	start = tim.tv_sec + (tim.tv_usec/1000000.0);
	
    // allocate temperature arrays on device
    hipMalloc((void **)&temp1_d, sizeof(double)*(ni+2)*(nj+2));
    hipMalloc((void **)&temp2_d, sizeof(double)*(ni+2)*(nj+2));
    
    // transfer temperature array from host to device
    hipMemcpy((void *)temp1_d, (void *)temp1_h, sizeof(double)*(ni+2)*(nj+2),
               hipMemcpyHostToDevice);
    hipMemcpy((void *)temp2_d, (void *)temp1_h, sizeof(double)*(ni+2)*(nj+2),
               hipMemcpyHostToDevice);

    tfac = 0.2;
    
    // main iteration loop  

    for (istep=0; istep < nstep; istep++) {
            dim3 grid_dim((ni+THREADS-1)/THREADS, nj, 1); 
            dim3 block_dim(THREADS, 1, 1); 
            step_kernel_cpu<<<grid_dim, block_dim>>>(ni+2, nj+2, tfac, temp1_d, temp2_d);

	    // swap the temp pointers
            temp_tmp = temp1_d;
            temp1_d = temp2_d;
            temp2_d = temp_tmp;
            
    } 
    
    hipMemcpy((void *)temp1_h, (void *)temp1_d, sizeof(double)*(ni+2)*(nj+2),
                   hipMemcpyDeviceToHost);
	
    gettimeofday(&tim, NULL);
    end = tim.tv_sec + (tim.tv_usec/1000000.0);
    printf("Time for computing: %.2f s\n",end-start);


    // output temp1 to a file
    
	fd = creat(argv[4], 00666);
	fd = open(argv[4], O_WRONLY);
	write(fd, temp1_h, (size_t)(ni+2)*(nj+2)*sizeof(double));
	close(fd);
    
/*
    FILE *fp;
    fp = fopen(filename, "w");
    fprintf(fp, "%d %d\n", ni, nj);
    for (j=0; j < nj; j++) {
        for (i=0; i < ni; i++) {
            fprintf(fp, "%.4f\n", j, i, temp1_h[i + ni*j]);
        }
    }
    fclose(fp);
*/
}



	.text
	.file	"cuda_heat_single.hip"
	.globl	_Z30__device_stub__step_kernel_cpuiidPdS_ # -- Begin function _Z30__device_stub__step_kernel_cpuiidPdS_
	.type	_Z30__device_stub__step_kernel_cpuiidPdS_,@function
_Z30__device_stub__step_kernel_cpuiidPdS_: # @_Z30__device_stub__step_kernel_cpuiidPdS_
	.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	4(%rsp), %rax
	movl	%edi, (%rax)
	movq	%rsp, %rdi
	movl	%esi, (%rdi)
	leaq	40(%rsp), %rsi
	movsd	%xmm0, (%rsi)
	leaq	32(%rsp), %r8
	movq	%rdx, (%r8)
	leaq	24(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%r8, 24(%rbx)
	movq	%rdx, 32(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%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	$_Z15step_kernel_cpuiidPdS_, %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	_Z30__device_stub__step_kernel_cpuiidPdS_, .Lfunc_end0-_Z30__device_stub__step_kernel_cpuiidPdS_
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function main
.LCPI1_0:
	.quad	0x4059000000000000              # double 100
.LCPI1_1:
	.quad	0x4069000000000000              # double 200
.LCPI1_2:
	.quad	0x4072c00000000000              # double 300
.LCPI1_3:
	.quad	0x3fc999999999999a              # double 0.20000000000000001
.LCPI1_4:
	.quad	0x412e848000000000              # double 1.0E+6
	.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	$88, %rsp
	.cfi_def_cfa_offset 144
	.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, %rbx
	cmpl	$4, %edi
	jle	.LBB1_19
# %bb.1:
	movq	8(%rbx), %rdi
	callq	atoi
	movl	%eax, %r14d
	movq	16(%rbx), %rdi
	callq	atoi
	movl	%eax, %r15d
	movq	%rbx, 64(%rsp)                  # 8-byte Spill
	movq	24(%rbx), %rdi
	callq	atoi
	movl	%eax, %ebp
	leal	2(%r14), %eax
	movl	%eax, (%rsp)                    # 4-byte Spill
	cltq
	leal	2(%r15), %ecx
	movl	%ecx, 12(%rsp)                  # 4-byte Spill
	movslq	%ecx, %rdi
	imulq	%rax, %rdi
	shlq	$3, %rdi
	movq	%rdi, 56(%rsp)                  # 8-byte Spill
	callq	malloc
	movq	%rax, %r13
	movq	%r15, 32(%rsp)                  # 8-byte Spill
	testl	%r15d, %r15d
	jle	.LBB1_6
# %bb.2:                                # %.preheader121.lr.ph
	leal	3(%r14), %r12d
	movl	%r14d, %ebx
	shlq	$3, %rbx
	movl	32(%rsp), %r15d                 # 4-byte Reload
.LBB1_3:                                # %.preheader121
                                        # =>This Inner Loop Header: Depth=1
	testl	%r14d, %r14d
	jle	.LBB1_5
# %bb.4:                                # %.lr.ph
                                        #   in Loop: Header=BB1_3 Depth=1
	movl	%r12d, %eax
	leaq	(,%rax,8), %rdi
	addq	%r13, %rdi
	xorl	%esi, %esi
	movq	%rbx, %rdx
	callq	memset@PLT
.LBB1_5:                                # %._crit_edge
                                        #   in Loop: Header=BB1_3 Depth=1
	addl	(%rsp), %r12d                   # 4-byte Folded Reload
	decq	%r15
	jne	.LBB1_3
.LBB1_6:                                # %.preheader120
	movq	32(%rsp), %rax                  # 8-byte Reload
	incl	%eax
	cmpl	$-1, %r14d
	movq	%r13, %r12
	jl	.LBB1_9
# %bb.7:                                # %.lr.ph127
	leal	1(%r14), %ecx
	cvtsi2sd	%ecx, %xmm0
	movl	%eax, %ecx
	movl	(%rsp), %esi                    # 4-byte Reload
	imull	%esi, %ecx
	movslq	%ecx, %rdx
	cmpl	$2, %esi
	movl	$1, %ecx
	cmovgel	%esi, %ecx
	leaq	(%r12,%rdx,8), %rdx
	xorl	%esi, %esi
	movsd	.LCPI1_0(%rip), %xmm1           # xmm1 = mem[0],zero
	movsd	.LCPI1_1(%rip), %xmm2           # xmm2 = mem[0],zero
.LBB1_8:                                # =>This Inner Loop Header: Depth=1
	xorps	%xmm3, %xmm3
	cvtsi2sd	%esi, %xmm3
	mulsd	%xmm1, %xmm3
	divsd	%xmm0, %xmm3
	addsd	%xmm2, %xmm3
	movsd	%xmm3, (%r12,%rsi,8)
	movsd	%xmm3, (%rdx,%rsi,8)
	incq	%rsi
	cmpq	%rsi, %rcx
	jne	.LBB1_8
.LBB1_9:                                # %.preheader
	movq	32(%rsp), %r13                  # 8-byte Reload
	cmpl	$-1, %r13d
	movq	56(%rsp), %r15                  # 8-byte Reload
	jl	.LBB1_12
# %bb.10:                               # %.lr.ph129
	xorps	%xmm0, %xmm0
	cvtsi2sd	%eax, %xmm0
	movl	12(%rsp), %ecx                  # 4-byte Reload
	cmpl	$2, %ecx
	movl	$1, %eax
	cmovgel	%ecx, %eax
	movslq	%r14d, %rcx
	leaq	16(,%rcx,8), %rdx
	shlq	$3, %rcx
	xorl	%esi, %esi
	xorpd	%xmm1, %xmm1
	movsd	.LCPI1_1(%rip), %xmm2           # xmm2 = mem[0],zero
	movsd	.LCPI1_2(%rip), %xmm3           # xmm3 = mem[0],zero
	movq	%r12, %rdi
.LBB1_11:                               # =>This Inner Loop Header: Depth=1
	xorps	%xmm4, %xmm4
	cvtsi2sd	%esi, %xmm4
	mulsd	%xmm1, %xmm4
	divsd	%xmm0, %xmm4
	movapd	%xmm4, %xmm5
	addsd	%xmm2, %xmm5
	movsd	%xmm5, (%rdi)
	addsd	%xmm3, %xmm4
	movsd	%xmm4, 8(%rdi,%rcx)
	incq	%rsi
	addq	%rdx, %rdi
	cmpq	%rsi, %rax
	jne	.LBB1_11
.LBB1_12:                               # %._crit_edge130
	leaq	72(%rsp), %rbx
	movq	%rbx, %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	movq	(%rbx), %rax
	movq	%rax, 48(%rsp)                  # 8-byte Spill
	movq	8(%rbx), %rax
	movq	%rax, 40(%rsp)                  # 8-byte Spill
	leaq	24(%rsp), %rbx
	movq	%rbx, %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	leaq	16(%rsp), %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	movq	(%rbx), %rdi
	movq	%r12, %rsi
	movq	%r15, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	16(%rsp), %rax
	movq	(%rax), %rdi
	movq	%r12, %rsi
	movq	%r15, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	testl	%ebp, %ebp
	jle	.LBB1_13
# %bb.14:                               # %.lr.ph133
	leal	127(%r14), %eax
	addl	$254, %r14d
	testl	%eax, %eax
	cmovnsl	%eax, %r14d
	sarl	$7, %r14d
	shlq	$32, %r13
	orq	%r14, %r13
	movabsq	$4294967424, %r14               # imm = 0x100000080
	movl	12(%rsp), %ebx                  # 4-byte Reload
.LBB1_15:                               # =>This Inner Loop Header: Depth=1
	movq	%r13, %rdi
	movl	$1, %esi
	movq	%r14, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB1_17
# %bb.16:                               #   in Loop: Header=BB1_15 Depth=1
	movq	24(%rsp), %rdx
	movq	16(%rsp), %rcx
	movl	(%rsp), %edi                    # 4-byte Reload
	movl	%ebx, %esi
	movsd	.LCPI1_3(%rip), %xmm0           # xmm0 = mem[0],zero
	callq	_Z30__device_stub__step_kernel_cpuiidPdS_
.LBB1_17:                               #   in Loop: Header=BB1_15 Depth=1
	movq	24(%rsp), %rax
	movq	16(%rsp), %rsi
	movq	%rsi, 24(%rsp)
	movq	%rax, 16(%rsp)
	decl	%ebp
	jne	.LBB1_15
	jmp	.LBB1_18
.LBB1_13:                               # %._crit_edge130.._crit_edge134_crit_edge
	movq	24(%rsp), %rsi
.LBB1_18:                               # %._crit_edge134
	xorps	%xmm0, %xmm0
	cvtsi2sdq	48(%rsp), %xmm0         # 8-byte Folded Reload
	xorps	%xmm1, %xmm1
	cvtsi2sdq	40(%rsp), %xmm1         # 8-byte Folded Reload
	divsd	.LCPI1_4(%rip), %xmm1
	addsd	%xmm0, %xmm1
	movsd	%xmm1, (%rsp)                   # 8-byte Spill
	movq	%r12, %rdi
	movq	%r15, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	leaq	72(%rsp), %rbx
	movq	%rbx, %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	xorps	%xmm1, %xmm1
	cvtsi2sdq	(%rbx), %xmm1
	xorps	%xmm0, %xmm0
	cvtsi2sdq	8(%rbx), %xmm0
	divsd	.LCPI1_4(%rip), %xmm0
	addsd	%xmm1, %xmm0
	subsd	(%rsp), %xmm0                   # 8-byte Folded Reload
	movl	$.L.str.1, %edi
	movb	$1, %al
	callq	printf
	movq	64(%rsp), %rbx                  # 8-byte Reload
	movq	32(%rbx), %rdi
	movl	$438, %esi                      # imm = 0x1B6
	callq	creat
	movq	32(%rbx), %rdi
	movl	$1, %esi
	xorl	%eax, %eax
	callq	open
	movl	%eax, %ebx
	movl	%eax, %edi
	movq	%r12, %rsi
	movq	%r15, %rdx
	callq	write
	movl	%ebx, %edi
	callq	close
	xorl	%eax, %eax
	addq	$88, %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
.LBB1_19:
	.cfi_def_cfa_offset 144
	movq	(%rbx), %rsi
	movl	$.L.str, %edi
	xorl	%eax, %eax
	callq	printf
	xorl	%edi, %edi
	callq	exit
.Lfunc_end1:
	.size	main, .Lfunc_end1-main
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	jne	.LBB2_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB2_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z15step_kernel_cpuiidPdS_, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end2:
	.size	__hip_module_ctor, .Lfunc_end2-__hip_module_ctor
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_dtor
	.type	__hip_module_dtor,@function
__hip_module_dtor:                      # @__hip_module_dtor
	.cfi_startproc
# %bb.0:
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB3_2
# %bb.1:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__hipUnregisterFatBinary
	movq	$0, __hip_gpubin_handle(%rip)
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB3_2:
	retq
.Lfunc_end3:
	.size	__hip_module_dtor, .Lfunc_end3-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z15step_kernel_cpuiidPdS_,@object # @_Z15step_kernel_cpuiidPdS_
	.section	.rodata,"a",@progbits
	.globl	_Z15step_kernel_cpuiidPdS_
	.p2align	3, 0x0
_Z15step_kernel_cpuiidPdS_:
	.quad	_Z30__device_stub__step_kernel_cpuiidPdS_
	.size	_Z15step_kernel_cpuiidPdS_, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"Usage: %s <ni> <ni> <nstep> <output file>\n"
	.size	.L.str, 43

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"Time for computing: %.2f s\n"
	.size	.L.str.1, 28

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z15step_kernel_cpuiidPdS_"
	.size	.L__unnamed_1, 27

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

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

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

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z15step_kernel_cpuiidPdS_ ; -- Begin function _Z15step_kernel_cpuiidPdS_
	.globl	_Z15step_kernel_cpuiidPdS_
	.p2align	8
	.type	_Z15step_kernel_cpuiidPdS_,@function
_Z15step_kernel_cpuiidPdS_:             ; @_Z15step_kernel_cpuiidPdS_
; %bb.0:
	s_load_b64 s[2:3], s[0:1], 0x0
	s_add_i32 s12, s15, 1
	s_waitcnt lgkmcnt(0)
	s_add_i32 s3, s3, -1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_ge_i32 s12, s3
	s_cbranch_scc1 .LBB0_7
; %bb.1:                                ; %.preheader.lr.ph
	s_clause 0x1
	s_load_b32 s13, s[0:1], 0x2c
	s_load_b256 s[4:11], s[0:1], 0x8
	s_add_i32 s1, s2, -1
	s_waitcnt lgkmcnt(0)
	s_and_b32 s0, s13, 0xffff
	s_add_i32 s13, s15, 2
	s_mul_i32 s14, s14, s0
	s_mul_i32 s10, s10, s0
	v_add3_u32 v0, v0, s14, 1
	s_mul_i32 s13, s2, s13
	s_mul_i32 s14, s11, s2
	s_mul_i32 s15, s15, s2
	s_mul_i32 s2, s2, s12
.LBB0_2:                                ; %.preheader
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_4 Depth 2
	s_mov_b32 s16, exec_lo
	v_cmpx_gt_i32_e64 s1, v0
	s_cbranch_execz .LBB0_6
; %bb.3:                                ; %.lr.ph
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_mov_b32 s17, 0
.LBB0_4:                                ;   Parent Loop BB0_2 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	v_add_nc_u32_e32 v1, s2, v0
	v_add_nc_u32_e32 v3, s15, v0
	v_add_nc_u32_e32 v11, s13, v0
	v_add_nc_u32_e32 v0, s10, v0
	s_delay_alu instid0(VALU_DEP_4)
	v_ashrrev_i32_e32 v2, 31, v1
	v_add_nc_u32_e32 v5, -1, v1
	v_ashrrev_i32_e32 v4, 31, v3
	v_add_nc_u32_e32 v7, 1, v1
	v_ashrrev_i32_e32 v12, 31, v11
	v_lshlrev_b64 v[1:2], 3, v[1:2]
	v_ashrrev_i32_e32 v6, 31, v5
	v_lshlrev_b64 v[3:4], 3, v[3:4]
	v_ashrrev_i32_e32 v8, 31, v7
	v_lshlrev_b64 v[11:12], 3, v[11:12]
	s_delay_alu instid0(VALU_DEP_4)
	v_lshlrev_b64 v[5:6], 3, v[5:6]
	v_add_co_u32 v9, vcc_lo, s6, v1
	v_add_co_ci_u32_e32 v10, vcc_lo, s7, v2, vcc_lo
	v_lshlrev_b64 v[7:8], 3, v[7:8]
	v_add_co_u32 v3, vcc_lo, s6, v3
	v_add_co_ci_u32_e32 v4, vcc_lo, s7, v4, vcc_lo
	v_add_co_u32 v5, vcc_lo, s6, v5
	v_add_co_ci_u32_e32 v6, vcc_lo, s7, v6, vcc_lo
	v_add_co_u32 v7, vcc_lo, s6, v7
	v_add_co_ci_u32_e32 v8, vcc_lo, s7, v8, vcc_lo
	s_clause 0x3
	global_load_b64 v[5:6], v[5:6], off
	global_load_b64 v[9:10], v[9:10], off
	global_load_b64 v[3:4], v[3:4], off
	global_load_b64 v[7:8], v[7:8], off
	v_add_co_u32 v11, vcc_lo, s6, v11
	v_add_co_ci_u32_e32 v12, vcc_lo, s7, v12, vcc_lo
	v_cmp_le_i32_e32 vcc_lo, s1, v0
	v_add_co_u32 v1, s0, s8, v1
	global_load_b64 v[11:12], v[11:12], off
	v_add_co_ci_u32_e64 v2, s0, s9, v2, s0
	s_or_b32 s17, vcc_lo, s17
	s_waitcnt vmcnt(3)
	v_fma_f64 v[5:6], v[9:10], -2.0, v[5:6]
	s_waitcnt vmcnt(2)
	v_fma_f64 v[3:4], v[9:10], -2.0, v[3:4]
	s_waitcnt vmcnt(1)
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_f64 v[5:6], v[7:8], v[5:6]
	s_waitcnt vmcnt(0)
	v_add_f64 v[3:4], v[3:4], v[11:12]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[3:4], v[5:6], v[3:4]
	v_fma_f64 v[3:4], v[3:4], s[4:5], v[9:10]
	global_store_b64 v[1:2], v[3:4], off
	s_and_not1_b32 exec_lo, exec_lo, s17
	s_cbranch_execnz .LBB0_4
; %bb.5:                                ; %Flow
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s17
.LBB0_6:                                ; %Flow68
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s16
	s_add_i32 s12, s11, s12
	s_add_i32 s13, s13, s14
	s_add_i32 s15, s15, s14
	s_add_i32 s2, s2, s14
	s_cmp_lt_i32 s12, s3
	s_cbranch_scc1 .LBB0_2
.LBB0_7:                                ; %._crit_edge49
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z15step_kernel_cpuiidPdS_
		.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 0
		.amdhsa_next_free_vgpr 13
		.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	_Z15step_kernel_cpuiidPdS_, .Lfunc_end0-_Z15step_kernel_cpuiidPdS_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 480
; NumSgprs: 20
; NumVgprs: 13
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 20
; NumVGPRsForWavesPerEU: 13
; 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: 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:
      - .offset:         0
        .size:           4
        .value_kind:     by_value
      - .offset:         4
        .size:           4
        .value_kind:     by_value
      - .offset:         8
        .size:           8
        .value_kind:     by_value
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .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:           _Z15step_kernel_cpuiidPdS_
    .private_segment_fixed_size: 0
    .sgpr_count:     20
    .sgpr_spill_count: 0
    .symbol:         _Z15step_kernel_cpuiidPdS_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     13
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
