//This is a generated CUDA code
#include<stdio.h>
#include<stdlib.h>
#include<time.h>

//f_alu = 1
//f_mad =8
//f_sfu =6
//d_alu =8
//d_mad =6
//b_alu =5
__constant__ float kconst[115] = {
2640.27049419,496.788317279,755.85277182,4070.19741521,6510.34703498,2039.14289025,3704.61925152,7755.66914948,
1861.26002473,1253.44798987,2629.46951817,7940.16158665,7767.42297282,5605.65608719,5373.03014677,3891.71871656,
366.965906878,8379.0600811,3930.37788456,8751.34651216,6043.60764794,9682.30745623,6178.41261724,4564.3931025,
374.255925925,4577.84509961,1671.08086172,9345.83987891,2901.2040663,4154.98833481,8208.38540292,4352.1328191,
9573.36924331,3580.73886635,269.083439352,352.007680208,5305.30563288,3596.20743502,438.919179395,9456.31492571,
8206.88128285,7265.8087816,8790.45564504,2464.18248902,4064.78800258,6832.31591731,9795.16429237,5786.85994295,
2010.51460232,4034.59881526,6585.45210739,8957.74158847,6403.29279037,3070.06110398,6175.24851691,1461.27925187,
3336.50691137,6910.80070577,9967.11557409,3549.38634683,2606.03554605,1575.68922804,4906.98800552,679.32387231,
5310.51577039,6347.74789192,4980.14693502,7079.55612995,2304.94310865,253.435257548,6597.60226307,9160.92905888,
2656.51142121,1532.34098062,9800.24073027,4697.5946278,4549.56735803,895.369414414,81.2862597505,954.653076872,
9802.40869968,3068.94073212,4343.28386194,4288.43198471,9454.29190141,687.820715136,8555.10227894,1020.62527779,
8887.51891817,1622.30149751,5517.20514654,2660.16395818,9482.94902673,8842.14459278,3834.4145655,700.939818034,
8816.58299456,4351.4402913,3601.03531056,2832.71861204,5178.80878256,3393.39850047,3436.36898163,2336.76169531,
3094.84512625,9409.78159435,1337.62467857,5439.82313727,7485.84090711,3687.98225492,3607.41760755,4976.65341527,
7191.02440951,4980.12665621,3064.46805377
};


//this is an unused kernel
__global__ void kk(float * d)
{
	d[0]=kconst[0];}


__global__ void kernel(float*,float*,float*,int*,int*,int *,int,float *,int *);
float uniform(float start,float end)
{
	return start+(end-start)*rand()/RAND_MAX;
}

int main(int argc, char* argv[])
{
	int gridx = 179;
	int gridy = 1;
	int gridz = 1;
	int bx = 115;
	int by = 1;
	int bz = 1;

	dim3 gridDim(gridx,gridy,gridz);
	dim3 blockDim(bx,by,bz);

	//srand( (unsigned)time( NULL ) );
	srand(4);
	//declare variables
	int N=64,size=0,loop=0,globalN=1024*1024*128,globalspace;

	//initialize the variables
	N=gridx*gridy*gridz*bx*by*bz;
	loop=16;

	N=gridx*gridy*gridz*bx*by*bz*17;
	size=(N)*sizeof(float);
	globalspace=(globalN)*sizeof(float);
	//declare the float pointer variables
	float * f0,*fd0,*f1,*fd1,*f2,*fd2,*gf,*gf0;
	//parameter f0
	f0=(float*)malloc(size);
	memset(f0,0,size);
	cudaMalloc((void**)&fd0,size);
	for(int i=0;i<N;i++)
	f0[i]=uniform(1,10000);
	cudaMemcpy(fd0,f0,size,cudaMemcpyHostToDevice);

	//parameter f1
	f1=(float*)malloc(size);
	memset(f1,0,size);
	cudaMalloc((void**)&fd1,size);
	for(int i=0;i<N;i++)
	f1[i]=uniform(1,10000);
	cudaMemcpy(fd1,f1,size,cudaMemcpyHostToDevice);

	//parameter f2
	f2=(float*)malloc(size);
	memset(f2,0,size);
	cudaMalloc((void**)&fd2,size);
	for(int i=0;i<N;i++)
	f2[i]=uniform(1,10000);
	cudaMemcpy(fd2,f2,size,cudaMemcpyHostToDevice);

	//global memory declare(float)
	gf=(float*)malloc(globalspace);
	memset(gf,0,globalspace);
	cudaMalloc((void**)&gf0,globalspace);
	for(int i=0;i<globalN;i++)
		gf[i]=uniform(1,10000);
	cudaMemcpy(gf0,gf,globalspace,cudaMemcpyHostToDevice);

	//declare the int pointer variables
	int * d0,*dd0,*d1,*dd1,*d2,*dd2,*gd,*gd0;
	size=(N)*sizeof(int);

	globalspace=(globalN)*sizeof(int);
	//parameter d0
	d0=(int*)malloc(size);
	memset(d0,0,size);
	cudaMalloc((void**)&dd0,size);
	for(int i=0;i<N;i++)
	d0[i]=uniform(1,10000);
	cudaMemcpy(dd0,d0,size,cudaMemcpyHostToDevice);

	//parameter d1
	d1=(int*)malloc(size);
	memset(d1,0,size);
	cudaMalloc((void**)&dd1,size);
	for(int i=0;i<N;i++)
	d1[i]=uniform(1,10000);
	cudaMemcpy(dd1,d1,size,cudaMemcpyHostToDevice);

	//parameter d2
	d2=(int*)malloc(size);
	memset(d2,0,size);
	cudaMalloc((void**)&dd2,size);
	for(int i=0;i<N;i++)
	d2[i]=uniform(1,10000);
	cudaMemcpy(dd2,d2,size,cudaMemcpyHostToDevice);

	//global memory declare(int)
	gd=(int*)malloc(globalspace);
	memset(gd,0,globalspace);
	cudaMalloc((void**)&gd0,globalspace);
	for(int i=0;i<globalN;i++)
		gd[i]=uniform(1,10000);
	cudaMemcpy(gd0,gd,globalspace,cudaMemcpyHostToDevice);

	kernel<<<gridx*gridy,bx*by*bz>>>(fd0,fd1,fd2,dd0,dd1,dd2,loop,gf0,gd0);

	cudaMemcpy(d2,dd2, size, cudaMemcpyDeviceToHost);
	/*for(int i=0;i<N;i++)
		printf("%d	",d2[i]);
	printf("\n");*/
	cudaMemcpy(f2,fd2, size, cudaMemcpyDeviceToHost);
	/*for(int i=0;i<N;i++)
		printf("%f	",f2[i]);
	printf("\n");*/
}

__global__ void kernel(float* f0,float* f1,float* f2,int *d0,int *d1,int *d2,int loop,float* gf,int* gd)
{
	//declear the regs
	asm volatile(".reg .u32	%rt<4>;\n\
	.reg .u32 %rnt<4>;\n\
	.reg .u32 %rc<4>;\n\
	.reg .u32 %rnc<4>;\n\
	.reg .u32 %rg<14>;\n\
	.reg .u32 %rgtid;\n\
	.reg .u32 %rbtid;");
	//compute global threadID
	asm volatile("mov.u32	%rt1,%tid.x;\n\
	mov.u32	%rt2,%tid.y;\n\
	mov.u32	%rt3,%tid.z;\n\
	mov.u32	%rnt1,%ntid.x;\n\
	mov.u32	%rnt2,%ntid.y;\n\
	mov.u32	%rnt3,%ntid.z;\n\
	mov.u32	%rc1,%ctaid.x;\n\
	mov.u32	%rc2,%ctaid.y;\n\
	mov.u32	%rc3,%ctaid.z;\n\
	mov.u32	%rnc1,%nctaid.x;\n\
	mov.u32	%rnc2,%nctaid.y;\n\
	mov.u32	%rnc3,%nctaid.z;");	//compute threadID in a block
	asm volatile("mul.lo.u32	%rg1,%rnt1,%rnt2;\n\
	mul.lo.u32	%rg2,%rg1,%rt3;\n\
	mul.lo.u32	%rg3,%rt2,%rnt1;\n\
	add.u32	%rg4,%rt1,%rg3;\n\
	add.u32	%rg5,%rg4,%rg2;\n\
	mov.u32 %rbtid,%rg5;");
	//compute blockid in a grid
	asm volatile("mul.lo.u32	%rg6,%rnc1,%rnc2;\n\
	mul.lo.u32	%rg7,%rg6,%rc3;\n\
	mul.lo.u32	%rg8,%rc2,%rnc1;\n\
	add.u32	%rg9,%rc1,%rg8;\n\
	add.u32	%rg10,%rg9,%rg7;");
	//compute blocksize
	asm volatile("mul.lo.u32	%rg11,%rnt1,%rnt2;\n\
	mul.lo.u32	%rg12,%rg11,%rnt3;\n\
	mul.lo.u32	%rg13,%rg10,%rg12;\n\
	add.u32	%rgtid,%rg13,%rbtid;");
	asm volatile(".reg .u64	%rdf<3>;\n\
	.reg .u64	%rpf<3>;\n\
	.reg .u64	%rdd<3>;\n\
	.reg .u64	%rpd<3>;\n\
	.reg .pred	%p_<10>;\n\
	.reg .f32	%f_<3>;\n\
	.reg .s32	%d_<3>;\n\
	.reg .u64	%offset;\n\
	.reg .u64	%g_distance;\n\
	.reg .u32	%loop;\n\
	.reg .u32	%pass;\n\
	.reg .u64	%distance;\n\
	.reg .u64	%distance_mem;\n\
	.reg .u64	%gfp;\n\
	.reg .u64	%gdp;\n\
	.reg .u32	%counter;");
	asm volatile(".reg .u32	%rd<15>;\n\
	.reg .f32	%rf<16>;\n\
	.reg .f32	%rgf<32>;\n\
	.reg .u32	%rgd<32>;\n\
	.reg .b32	%rb<5>;\n\
	.reg .f64    %rdouble<8>;\n\
	.reg .f64	%rmadf<3>;\n\
	.reg .s32	%ropd3;\n\
	.reg .b32	%ropb<2>;\n\
	mov .b32	%ropb0,652;\n\
	mov .b32	%ropb1,256;\n\
	mov .u64	%distance,20585;\n\
	mov .u64	%distance_mem,262144;\n\
	mov .u64	%g_distance,65536;\n\
	mov .s32	%ropd3,5786;\n\
	mov .f64	%rmadf1,19.5923428225;\n\
	mov .f64     %rmadf2,23.9162632576;\n\
	mov .f64     %rmadf0,86.2923868728;");

	asm volatile(".reg .u64    %c_base;\n\
	.reg .u64    %c_offset;\n\
	.reg .u64    %c_mem;\n\
	mov.u64  %c_base,kconst;\n\
	mul.wide.u32   %c_offset,%rbtid,4;\n\
	add.u64  %c_mem,%c_base,%c_offset;");

	//get the loop number
	asm volatile("ld.param.u32   %loop, [_Z6kernelPfS_S_PiS0_S0_iS_S0__param_6];");//loop

	asm volatile("mul.wide.u32	%offset,%rgtid,4;");//get the address offset
//d0 data

	asm volatile("ld.param.u64	%rpd0, [_Z6kernelPfS_S_PiS0_S0_iS_S0__param_0];\n\
	add.u64	%rdd0,%rpd0,%offset;");
//d1 data

	asm volatile("ld.param.u64	%rpd1, [_Z6kernelPfS_S_PiS0_S0_iS_S0__param_1];\n\
	add.u64	%rdd1,%rpd1,%offset;");
//d2 data

	asm volatile("ld.param.u64	%rpd2, [_Z6kernelPfS_S_PiS0_S0_iS_S0__param_2];\n\
	add.u64	%rdd2,%rpd2,%offset;");
//f0 data

	asm volatile("ld.param.u64	%rpf0, [_Z6kernelPfS_S_PiS0_S0_iS_S0__param_3];\n\
	add.u64	%rdf0,%rpf0,%offset;");
//f1 data

	asm volatile("ld.param.u64	%rpf1, [_Z6kernelPfS_S_PiS0_S0_iS_S0__param_4];\n\
	add.u64	%rdf1,%rpf1,%offset;");
//f2 data

	asm volatile("ld.param.u64	%rpf2, [_Z6kernelPfS_S_PiS0_S0_iS_S0__param_5];\n\
	add.u64	%rdf2,%rpf2,%offset;");

	asm volatile("ld.param.u64	%gfp, [_Z6kernelPfS_S_PiS0_S0_iS_S0__param_7];\n\
	add.u64	%gfp,%gfp,%offset;");

	asm volatile("ld.param.u64	%gdp, [_Z6kernelPfS_S_PiS0_S0_iS_S0__param_8];\n\
	add.u64	%gdp,%gdp,%offset;");
	
//load the values of parameters
	asm volatile("ld.global.s32	%d_0,[%rdd0+0];");
	asm volatile("ld.global.s32	%d_1,[%rdd1+0];");
	asm volatile("ld.global.f32	%f_0,[%rdf0+0];");
	asm volatile("ld.global.f32	%f_1,[%rdf1+0];");


	asm volatile("mul.wide.u32	%distance,20585,4;\n\
	add.u64	%rdd0,%rdd0,%distance;\n\
	add.u64	%rdd1,%rdd1,%distance;\n\
	add.u64	%rdf0,%rdf0,%distance;\n\
	add.u64	%rdf1,%rdf1,%distance;");
	asm volatile("mov.u32	%counter,0;");
	asm volatile("$BB_LABEL0:");
	// instruction mix
	asm volatile("mov.u32 %pass,0;\n\
	fma.rn.f64	%rdouble0, %rmadf0, %rmadf1, %rmadf2;\n\
	cvt.rn.f32.f64     %rf0,%rdouble0;\n\
	lg2.approx.f32	%rf1, %f_1;\n\
	mad.lo.s32	%rd0, %d_0, %d_1, %ropd3;\n\
	mad.rn.f64	%rdouble1, %rmadf0, %rmadf1, %rmadf2;\n\
	cvt.rn.f32.f64     %rf2,%rdouble1;\n\
	lg2.approx.f32	%rf3, %f_0;\n\
	sqrt.rn.f32	%rf4, %f_1;\n\
	st.global.f32  [%gfp+0],%rgf19;\n\
	add.u64 %gfp,%gfp,%g_distance;\n\
	mul.f32	%rf5, %rf0, %rf1;\n\
	ld.const.f32	%rf0, [%c_mem+0];\n\
	cvt.f64.f32     %rmadf0, %rf0;\n\
	cvt.f64.f32     %rmadf1, %rf1;\n\
	cvt.f64.f32     %rmadf2, %rf2;\n\
	fma.rn.f64	%rdouble2, %rmadf0, %rmadf1, %rmadf2;\n\
	cvt.rn.f32.f64     %rf6,%rdouble2;\n\
	ld.const.f32	%rf3, [%c_mem+0];\n\
	xor.b32	%rb0, %ropb0, %ropb1;\n\
	ld.global.f32	%rgf11, [%gfp+0];\n\
	add.u64 %gfp,%gfp,%g_distance;\n\
	ld.global.f32	%rgf17, [%gfp+0];\n\
	add.u64 %gfp,%gfp,%g_distance;\n\
	mad.lo.s32	%rd1, %d_0, %d_1, %ropd3;\n\
	mad.lo.s32	%rd2, %d_0, %d_1, %ropd3;\n\
	ld.global.f32	%rgf8, [%gfp+0];\n\
	add.u64 %gfp,%gfp,%g_distance;\n\
	cvt.f64.f32     %rmadf0, %rf1;\n\
	cvt.f64.f32     %rmadf1, %rf2;\n\
	cvt.f64.f32     %rmadf2, %rf3;\n\
	mad.rn.f64	%rdouble3, %rmadf0, %rmadf1, %rmadf2;\n\
	cvt.rn.f32.f64     %rf7,%rdouble3;\n\
	mul.lo.s32	%rd3, %d_0, %d_1;\n\
	ld.const.f32	%rf1, [%c_mem+0];\n\
	neg.s32	%rd4, %d_0;\n\
	not.b32	%rb1, %ropb0;\n\
	cvt.f64.f32     %rmadf0, %rf2;\n\
	cvt.f64.f32     %rmadf1, %rf3;\n\
	cvt.f64.f32     %rmadf2, %rf4;\n\
	fma.rn.f64	%rdouble4, %rmadf0, %rmadf1, %rmadf2;\n\
	cvt.rn.f32.f64     %rf8,%rdouble4;\n\
	mad.lo.s32	%rd5, %rd0, %rd1, %rd2;\n\
	ld.global.f32	%rgf2, [%gfp+0];\n\
	add.u64 %gfp,%gfp,%g_distance;\n\
	cvt.f64.f32     %rmadf0, %rf3;\n\
	cvt.f64.f32     %rmadf1, %rf4;\n\
	cvt.f64.f32     %rmadf2, %rf5;\n\
	mad.rn.f64	%rdouble5, %rmadf0, %rmadf1, %rmadf2;\n\
	cvt.rn.f32.f64     %rf9,%rdouble5;\n\
	and.b32	%rb2, %ropb0, %ropb1;\n\
	cvt.f64.f32     %rmadf0, %rf4;\n\
	cvt.f64.f32     %rmadf1, %rf5;\n\
	cvt.f64.f32     %rmadf2, %rf6;\n\
	mad.rn.f64	%rdouble6, %rmadf0, %rmadf1, %rmadf2;\n\
	cvt.rn.f32.f64     %rf10,%rdouble6;\n\
	lg2.approx.f32	%rf11, %rf6;\n\
	cos.approx.f32	%rf12, %rf7;\n\
	rem.s32	%rd6, %rd1, %d_0;\n\
	rem.s32	%rd7, %rd2, %d_0;\n\
	mad.lo.s32	%rd8, %rd3, %rd4, %rd5;\n\
	cvt.f64.f32     %rmadf0, %rf7;\n\
	cvt.f64.f32     %rmadf1, %rf8;\n\
	cvt.f64.f32     %rmadf2, %rf9;\n\
	fma.rn.f64	%rdouble7, %rmadf0, %rmadf1, %rmadf2;\n\
	cvt.rn.f32.f64     %rf13,%rdouble7;\n\
	xor.b32	%rb3, %ropb0, %ropb1;\n\
	ex2.approx.f32	%rf14, %rf9;\n\
	ld.global.f32	%rgf31, [%gfp+0];\n\
	add.u64 %gfp,%gfp,%g_distance;\n\
	ld.global.f32	%rgf25, [%gfp+0];\n\
	add.u64 %gfp,%gfp,%g_distance;\n\
	shl.b32	%rb4, %ropb0, %ropb1;\n\
	max.s32	%rd9, %rd4, %rd5;\n\
	sad.s32	%rd10, %rd5, %rd6, %rd7;\n\
	ld.global.f32	%rgf20, [%gfp+0];\n\
	add.u64 %gfp,%gfp,%g_distance;\n\
	div.s32	%rd11, %rd6, %d_0;\n\
	mad.lo.s32	%rd12, %rd7, %rd8, %rd9;\n\
	min.s32	%rd13, %rd8, %rd9;\n\
	");
	asm volatile("add.u32	%counter,%counter,1;\n\
	setp.ge.u32	 %p_0, %counter, %loop;\n\
	@!%p_0	bra	$BB_LABEL0;");

	asm volatile("st.global.f32 [%rdf2+0],%rf15;\n\
	st.global.s32 [%rdd2+0],%rd14;");
	asm volatile("add.u64	%rdf2,%rdf2,%distance;\n\
	add.u64	%rdd2,%rdd2,%distance;");

	//end basic block;
	asm volatile("$BB_LABEL1:");
	asm volatile("exit;");
}

	.file	"tmpxft_0023d880_00000000-6_2018-8-20-7-20-52-1537453252-699299-13834.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2031:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2031:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z7uniformff
	.type	_Z7uniformff, @function
_Z7uniformff:
.LFB2027:
	.cfi_startproc
	endbr64
	subss	%xmm0, %xmm1
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movss	%xmm0, 12(%rsp)
	movss	%xmm1, 8(%rsp)
	call	rand@PLT
	movss	12(%rsp), %xmm2
	cvtsi2ssl	%eax, %xmm0
	mulss	8(%rsp), %xmm0
	mulss	.LC0(%rip), %xmm0
	addq	$24, %rsp
	.cfi_def_cfa_offset 8
	addss	%xmm2, %xmm0
	ret
	.cfi_endproc
.LFE2027:
	.size	_Z7uniformff, .-_Z7uniformff
	.globl	_Z21__device_stub__Z2kkPfPf
	.type	_Z21__device_stub__Z2kkPfPf, @function
_Z21__device_stub__Z2kkPfPf:
.LFB2053:
	.cfi_startproc
	endbr64
	subq	$104, %rsp
	.cfi_def_cfa_offset 112
	movq	%rdi, 8(%rsp)
	leaq	24(%rsp), %rcx
	leaq	16(%rsp), %rdx
	leaq	44(%rsp), %rsi
	leaq	32(%rsp), %rdi
	movq	%fs:40, %rax
	movq	%rax, 88(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 40(%rsp)
	movq	%rax, 80(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 32(%rsp)
	movq	%rax, 44(%rsp)
	movl	$1, 52(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L4
	pushq	24(%rsp)
	.cfi_def_cfa_offset 120
	leaq	_Z2kkPf(%rip), %rdi
	pushq	24(%rsp)
	.cfi_def_cfa_offset 128
	movq	60(%rsp), %rcx
	movl	68(%rsp), %r8d
	movq	48(%rsp), %rsi
	movl	56(%rsp), %edx
	leaq	96(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 120
	popq	%rdx
	.cfi_def_cfa_offset 112
.L4:
	movq	88(%rsp), %rax
	subq	%fs:40, %rax
	je	.L6
	call	__stack_chk_fail@PLT
.L6:
	addq	$104, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2053:
	.size	_Z21__device_stub__Z2kkPfPf, .-_Z21__device_stub__Z2kkPfPf
	.globl	_Z2kkPf
	.type	_Z2kkPf, @function
_Z2kkPf:
.LFB2054:
	.cfi_startproc
	endbr64
	jmp	_Z21__device_stub__Z2kkPfPf
	.cfi_endproc
.LFE2054:
	.size	_Z2kkPf, .-_Z2kkPf
	.globl	_Z43__device_stub__Z6kernelPfS_S_PiS0_S0_iS_S0_PfS_S_PiS0_S0_iS_S0_
	.type	_Z43__device_stub__Z6kernelPfS_S_PiS0_S0_iS_S0_PfS_S_PiS0_S0_iS_S0_, @function
_Z43__device_stub__Z6kernelPfS_S_PiS0_S0_iS_S0_PfS_S_PiS0_S0_iS_S0_:
.LFB2055:
	.cfi_startproc
	endbr64
	subq	$216, %rsp
	.cfi_def_cfa_offset 224
	movq	232(%rsp), %rax
	movq	%rdi, 56(%rsp)
	leaq	80(%rsp), %rdi
	movq	%rsi, 48(%rsp)
	leaq	92(%rsp), %rsi
	movq	%rax, 8(%rsp)
	movq	240(%rsp), %rax
	movq	%rdx, 40(%rsp)
	leaq	64(%rsp), %rdx
	movq	%rcx, 32(%rsp)
	leaq	72(%rsp), %rcx
	movq	%r8, 24(%rsp)
	movq	%r9, 16(%rsp)
	movq	%rax, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 200(%rsp)
	xorl	%eax, %eax
	leaq	56(%rsp), %rax
	movl	$1, 88(%rsp)
	movq	%rax, 128(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	224(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 184(%rsp)
	movq	%rsp, %rax
	movq	%rax, 192(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 80(%rsp)
	movq	%rax, 92(%rsp)
	movl	$1, 100(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L10
	pushq	72(%rsp)
	.cfi_def_cfa_offset 232
	leaq	_Z6kernelPfS_S_PiS0_S0_iS_S0_(%rip), %rdi
	pushq	72(%rsp)
	.cfi_def_cfa_offset 240
	movq	108(%rsp), %rcx
	movl	116(%rsp), %r8d
	movq	96(%rsp), %rsi
	movl	104(%rsp), %edx
	leaq	144(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 232
	popq	%rdx
	.cfi_def_cfa_offset 224
.L10:
	movq	200(%rsp), %rax
	subq	%fs:40, %rax
	je	.L12
	call	__stack_chk_fail@PLT
.L12:
	addq	$216, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2055:
	.size	_Z43__device_stub__Z6kernelPfS_S_PiS0_S0_iS_S0_PfS_S_PiS0_S0_iS_S0_, .-_Z43__device_stub__Z6kernelPfS_S_PiS0_S0_iS_S0_PfS_S_PiS0_S0_iS_S0_
	.globl	_Z6kernelPfS_S_PiS0_S0_iS_S0_
	.type	_Z6kernelPfS_S_PiS0_S0_iS_S0_, @function
_Z6kernelPfS_S_PiS0_S0_iS_S0_:
.LFB2056:
	.cfi_startproc
	endbr64
	jmp	_Z43__device_stub__Z6kernelPfS_S_PiS0_S0_iS_S0_PfS_S_PiS0_S0_iS_S0_
	.cfi_endproc
.LFE2056:
	.size	_Z6kernelPfS_S_PiS0_S0_iS_S0_, .-_Z6kernelPfS_S_PiS0_S0_iS_S0_
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2028:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	movl	$4, %edi
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	xorl	%ebx, %ebx
	subq	$104, %rsp
	.cfi_def_cfa_offset 144
	movq	%fs:40, %rax
	movq	%rax, 88(%rsp)
	xorl	%eax, %eax
	call	srand@PLT
	movl	$1399780, %edi
	call	malloc@PLT
	movl	$349945, %ecx
	movl	$1399780, %esi
	movq	%rax, %rbp
	xorl	%eax, %eax
	movq	%rbp, %rdi
	rep stosl
	movq	%rsp, %rdi
	call	cudaMalloc@PLT
.L16:
	movss	.LC1(%rip), %xmm1
	movss	.LC2(%rip), %xmm0
	call	_Z7uniformff
	movss	%xmm0, 0(%rbp,%rbx,4)
	incq	%rbx
	cmpq	$349945, %rbx
	jne	.L16
	movq	(%rsp), %rdi
	movq	%rbp, %rsi
	movl	$1, %ecx
	xorl	%ebp, %ebp
	movl	$1399780, %edx
	call	cudaMemcpy@PLT
	movl	$1399780, %edi
	call	malloc@PLT
	movq	%rbx, %rcx
	movl	$1399780, %esi
	movq	%rax, %r12
	xorl	%eax, %eax
	movq	%r12, %rdi
	rep stosl
	leaq	8(%rsp), %rdi
	call	cudaMalloc@PLT
.L17:
	movss	.LC1(%rip), %xmm1
	movss	.LC2(%rip), %xmm0
	call	_Z7uniformff
	movss	%xmm0, (%r12,%rbp,4)
	incq	%rbp
	cmpq	$349945, %rbp
	jne	.L17
	movq	8(%rsp), %rdi
	movl	$1, %ecx
	movl	$1399780, %edx
	movq	%r12, %rsi
	call	cudaMemcpy@PLT
	movl	$1399780, %edi
	call	malloc@PLT
	movq	%rbp, %rcx
	movl	$1399780, %esi
	xorl	%ebp, %ebp
	movq	%rax, %rbx
	xorl	%eax, %eax
	movq	%rbx, %rdi
	rep stosl
	leaq	16(%rsp), %rdi
	call	cudaMalloc@PLT
.L18:
	movss	.LC1(%rip), %xmm1
	movss	.LC2(%rip), %xmm0
	call	_Z7uniformff
	movss	%xmm0, (%rbx,%rbp,4)
	incq	%rbp
	cmpq	$349945, %rbp
	jne	.L18
	movq	16(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbx, %rsi
	xorl	%r12d, %r12d
	movl	$1399780, %edx
	call	cudaMemcpy@PLT
	movl	$536870912, %edi
	call	malloc@PLT
	movl	$134217728, %ecx
	movl	$536870912, %esi
	movq	%rax, %rbp
	xorl	%eax, %eax
	movq	%rbp, %rdi
	rep stosl
	leaq	24(%rsp), %rdi
	call	cudaMalloc@PLT
.L19:
	movss	.LC1(%rip), %xmm1
	movss	.LC2(%rip), %xmm0
	call	_Z7uniformff
	movss	%xmm0, 0(%rbp,%r12,4)
	incq	%r12
	cmpq	$134217728, %r12
	jne	.L19
	movq	24(%rsp), %rdi
	movq	%rbp, %rsi
	movl	$1, %ecx
	xorl	%ebp, %ebp
	movl	$536870912, %edx
	call	cudaMemcpy@PLT
	movl	$1399780, %edi
	call	malloc@PLT
	movl	$349945, %ecx
	movl	$1399780, %esi
	movq	%rax, %r12
	xorl	%eax, %eax
	movq	%r12, %rdi
	rep stosl
	leaq	32(%rsp), %rdi
	call	cudaMalloc@PLT
.L20:
	movss	.LC1(%rip), %xmm1
	movss	.LC2(%rip), %xmm0
	call	_Z7uniformff
	cvttss2sil	%xmm0, %eax
	movl	%eax, (%r12,%rbp,4)
	incq	%rbp
	cmpq	$349945, %rbp
	jne	.L20
	movq	32(%rsp), %rdi
	movq	%r12, %rsi
	movl	$1, %ecx
	xorl	%r12d, %r12d
	movl	$1399780, %edx
	call	cudaMemcpy@PLT
	movl	$1399780, %edi
	call	malloc@PLT
	movq	%rbp, %rcx
	movl	$1399780, %esi
	movq	%rax, %r13
	xorl	%eax, %eax
	movq	%r13, %rdi
	rep stosl
	leaq	40(%rsp), %rdi
	call	cudaMalloc@PLT
.L21:
	movss	.LC1(%rip), %xmm1
	movss	.LC2(%rip), %xmm0
	call	_Z7uniformff
	cvttss2sil	%xmm0, %eax
	movl	%eax, 0(%r13,%r12,4)
	incq	%r12
	cmpq	$349945, %r12
	jne	.L21
	movq	40(%rsp), %rdi
	movl	$1, %ecx
	movl	$1399780, %edx
	movq	%r13, %rsi
	call	cudaMemcpy@PLT
	movl	$1399780, %edi
	call	malloc@PLT
	movq	%r12, %rcx
	movl	$1399780, %esi
	xorl	%r12d, %r12d
	movq	%rax, %rbp
	xorl	%eax, %eax
	movq	%rbp, %rdi
	rep stosl
	leaq	48(%rsp), %rdi
	call	cudaMalloc@PLT
.L22:
	movss	.LC1(%rip), %xmm1
	movss	.LC2(%rip), %xmm0
	call	_Z7uniformff
	cvttss2sil	%xmm0, %eax
	movl	%eax, 0(%rbp,%r12,4)
	incq	%r12
	cmpq	$349945, %r12
	jne	.L22
	movq	48(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbp, %rsi
	xorl	%r13d, %r13d
	movl	$1399780, %edx
	call	cudaMemcpy@PLT
	movl	$536870912, %edi
	call	malloc@PLT
	movl	$134217728, %ecx
	movl	$536870912, %esi
	movq	%rax, %r12
	xorl	%eax, %eax
	movq	%r12, %rdi
	rep stosl
	leaq	56(%rsp), %rdi
	call	cudaMalloc@PLT
.L23:
	movss	.LC1(%rip), %xmm1
	movss	.LC2(%rip), %xmm0
	call	_Z7uniformff
	cvttss2sil	%xmm0, %eax
	movl	%eax, (%r12,%r13,4)
	incq	%r13
	cmpq	$134217728, %r13
	jne	.L23
	movq	56(%rsp), %rdi
	movl	$1, %ecx
	movl	$536870912, %edx
	movq	%r12, %rsi
	call	cudaMemcpy@PLT
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	movabsq	$4294967411, %rdx
	movl	$1, %esi
	movabsq	$4294967475, %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L24
	pushq	%rax
	.cfi_def_cfa_offset 152
	pushq	64(%rsp)
	.cfi_def_cfa_offset 160
	pushq	40(%rsp)
	.cfi_def_cfa_offset 168
	pushq	$16
	.cfi_def_cfa_offset 176
	movq	80(%rsp), %r9
	movq	72(%rsp), %r8
	movq	64(%rsp), %rcx
	movq	48(%rsp), %rdx
	movq	40(%rsp), %rsi
	movq	32(%rsp), %rdi
	call	_Z43__device_stub__Z6kernelPfS_S_PiS0_S0_iS_S0_PfS_S_PiS0_S0_iS_S0_
	addq	$32, %rsp
	.cfi_def_cfa_offset 144
.L24:
	movq	48(%rsp), %rsi
	movl	$2, %ecx
	movl	$1399780, %edx
	movq	%rbp, %rdi
	call	cudaMemcpy@PLT
	movq	16(%rsp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdi
	movl	$1399780, %edx
	call	cudaMemcpy@PLT
	movq	88(%rsp), %rax
	subq	%fs:40, %rax
	je	.L25
	call	__stack_chk_fail@PLT
.L25:
	addq	$104, %rsp
	.cfi_def_cfa_offset 40
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2028:
	.size	main, .-main
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC3:
	.string	"_Z6kernelPfS_S_PiS0_S0_iS_S0_"
.LC4:
	.string	"_Z2kkPf"
.LC5:
	.string	"kconst"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2058:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -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	.LC3(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z6kernelPfS_S_PiS0_S0_iS_S0_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC4(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z2kkPf(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC5(%rip), %rdx
	movl	$460, %r9d
	leaq	_ZL6kconst(%rip), %rsi
	pushq	$1
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	call	__cudaRegisterVar@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rdx
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2058:
	.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
	.local	_ZL6kconst
	.comm	_ZL6kconst,460,32
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC0:
	.long	805306368
	.align 4
.LC1:
	.long	1176256512
	.align 4
.LC2:
	.long	1065353216
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:

//This is a generated CUDA code

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

//f_alu = 1
//f_mad =8
//f_sfu =6
//d_alu =8
//d_mad =6
//b_alu =5
__constant__ float kconst[115] = {
2640.27049419,496.788317279,755.85277182,4070.19741521,6510.34703498,2039.14289025,3704.61925152,7755.66914948,
1861.26002473,1253.44798987,2629.46951817,7940.16158665,7767.42297282,5605.65608719,5373.03014677,3891.71871656,
366.965906878,8379.0600811,3930.37788456,8751.34651216,6043.60764794,9682.30745623,6178.41261724,4564.3931025,
374.255925925,4577.84509961,1671.08086172,9345.83987891,2901.2040663,4154.98833481,8208.38540292,4352.1328191,
9573.36924331,3580.73886635,269.083439352,352.007680208,5305.30563288,3596.20743502,438.919179395,9456.31492571,
8206.88128285,7265.8087816,8790.45564504,2464.18248902,4064.78800258,6832.31591731,9795.16429237,5786.85994295,
2010.51460232,4034.59881526,6585.45210739,8957.74158847,6403.29279037,3070.06110398,6175.24851691,1461.27925187,
3336.50691137,6910.80070577,9967.11557409,3549.38634683,2606.03554605,1575.68922804,4906.98800552,679.32387231,
5310.51577039,6347.74789192,4980.14693502,7079.55612995,2304.94310865,253.435257548,6597.60226307,9160.92905888,
2656.51142121,1532.34098062,9800.24073027,4697.5946278,4549.56735803,895.369414414,81.2862597505,954.653076872,
9802.40869968,3068.94073212,4343.28386194,4288.43198471,9454.29190141,687.820715136,8555.10227894,1020.62527779,
8887.51891817,1622.30149751,5517.20514654,2660.16395818,9482.94902673,8842.14459278,3834.4145655,700.939818034,
8816.58299456,4351.4402913,3601.03531056,2832.71861204,5178.80878256,3393.39850047,3436.36898163,2336.76169531,
3094.84512625,9409.78159435,1337.62467857,5439.82313727,7485.84090711,3687.98225492,3607.41760755,4976.65341527,
7191.02440951,4980.12665621,3064.46805377
};


//this is an unused kernel
__global__ void kk(float * d)
{
	d[0]=kconst[0];}


__global__ void kernel(float*,float*,float*,int*,int*,int *,int,float *,int *);
float uniform(float start,float end)
{
	return start+(end-start)*rand()/RAND_MAX;
}

int main(int argc, char* argv[])
{
	int gridx = 179;
	int gridy = 1;
	int gridz = 1;
	int bx = 115;
	int by = 1;
	int bz = 1;

	dim3 gridDim(gridx,gridy,gridz);
	dim3 blockDim(bx,by,bz);

	//srand( (unsigned)time( NULL ) );
	srand(4);
	//declare variables
	int N=64,size=0,loop=0,globalN=1024*1024*128,globalspace;

	//initialize the variables
	N=gridx*gridy*gridz*bx*by*bz;
	loop=16;

	N=gridx*gridy*gridz*bx*by*bz*17;
	size=(N)*sizeof(float);
	globalspace=(globalN)*sizeof(float);
	//declare the float pointer variables
	float * f0,*fd0,*f1,*fd1,*f2,*fd2,*gf,*gf0;
	//parameter f0
	f0=(float*)malloc(size);
	memset(f0,0,size);
	hipMalloc((void**)&fd0,size);
	for(int i=0;i<N;i++)
	f0[i]=uniform(1,10000);
	hipMemcpy(fd0,f0,size,hipMemcpyHostToDevice);

	//parameter f1
	f1=(float*)malloc(size);
	memset(f1,0,size);
	hipMalloc((void**)&fd1,size);
	for(int i=0;i<N;i++)
	f1[i]=uniform(1,10000);
	hipMemcpy(fd1,f1,size,hipMemcpyHostToDevice);

	//parameter f2
	f2=(float*)malloc(size);
	memset(f2,0,size);
	hipMalloc((void**)&fd2,size);
	for(int i=0;i<N;i++)
	f2[i]=uniform(1,10000);
	hipMemcpy(fd2,f2,size,hipMemcpyHostToDevice);

	//global memory declare(float)
	gf=(float*)malloc(globalspace);
	memset(gf,0,globalspace);
	hipMalloc((void**)&gf0,globalspace);
	for(int i=0;i<globalN;i++)
		gf[i]=uniform(1,10000);
	hipMemcpy(gf0,gf,globalspace,hipMemcpyHostToDevice);

	//declare the int pointer variables
	int * d0,*dd0,*d1,*dd1,*d2,*dd2,*gd,*gd0;
	size=(N)*sizeof(int);

	globalspace=(globalN)*sizeof(int);
	//parameter d0
	d0=(int*)malloc(size);
	memset(d0,0,size);
	hipMalloc((void**)&dd0,size);
	for(int i=0;i<N;i++)
	d0[i]=uniform(1,10000);
	hipMemcpy(dd0,d0,size,hipMemcpyHostToDevice);

	//parameter d1
	d1=(int*)malloc(size);
	memset(d1,0,size);
	hipMalloc((void**)&dd1,size);
	for(int i=0;i<N;i++)
	d1[i]=uniform(1,10000);
	hipMemcpy(dd1,d1,size,hipMemcpyHostToDevice);

	//parameter d2
	d2=(int*)malloc(size);
	memset(d2,0,size);
	hipMalloc((void**)&dd2,size);
	for(int i=0;i<N;i++)
	d2[i]=uniform(1,10000);
	hipMemcpy(dd2,d2,size,hipMemcpyHostToDevice);

	//global memory declare(int)
	gd=(int*)malloc(globalspace);
	memset(gd,0,globalspace);
	hipMalloc((void**)&gd0,globalspace);
	for(int i=0;i<globalN;i++)
		gd[i]=uniform(1,10000);
	hipMemcpy(gd0,gd,globalspace,hipMemcpyHostToDevice);

	kernel<<<gridx*gridy,bx*by*bz>>>(fd0,fd1,fd2,dd0,dd1,dd2,loop,gf0,gd0);

	hipMemcpy(d2,dd2, size, hipMemcpyDeviceToHost);
	/*for(int i=0;i<N;i++)
		printf("%d	",d2[i]);
	printf("\n");*/
	hipMemcpy(f2,fd2, size, hipMemcpyDeviceToHost);
	/*for(int i=0;i<N;i++)
		printf("%f	",f2[i]);
	printf("\n");*/
}

__global__ void kernel(float* f0,float* f1,float* f2,int *d0,int *d1,int *d2,int loop,float* gf,int* gd)
{
	//declear the regs
	asm volatile(".reg .u32	%rt<4>;\n\
	.reg .u32 %rnt<4>;\n\
	.reg .u32 %rc<4>;\n\
	.reg .u32 %rnc<4>;\n\
	.reg .u32 %rg<14>;\n\
	.reg .u32 %rgtid;\n\
	.reg .u32 %rbtid;");
	//compute global threadID
	asm volatile("mov.u32	%rt1,%tid.x;\n\
	mov.u32	%rt2,%tid.y;\n\
	mov.u32	%rt3,%tid.z;\n\
	mov.u32	%rnt1,%ntid.x;\n\
	mov.u32	%rnt2,%ntid.y;\n\
	mov.u32	%rnt3,%ntid.z;\n\
	mov.u32	%rc1,%ctaid.x;\n\
	mov.u32	%rc2,%ctaid.y;\n\
	mov.u32	%rc3,%ctaid.z;\n\
	mov.u32	%rnc1,%nctaid.x;\n\
	mov.u32	%rnc2,%nctaid.y;\n\
	mov.u32	%rnc3,%nctaid.z;");	//compute threadID in a block
	asm volatile("mul.lo.u32	%rg1,%rnt1,%rnt2;\n\
	mul.lo.u32	%rg2,%rg1,%rt3;\n\
	mul.lo.u32	%rg3,%rt2,%rnt1;\n\
	add.u32	%rg4,%rt1,%rg3;\n\
	add.u32	%rg5,%rg4,%rg2;\n\
	mov.u32 %rbtid,%rg5;");
	//compute blockid in a grid
	asm volatile("mul.lo.u32	%rg6,%rnc1,%rnc2;\n\
	mul.lo.u32	%rg7,%rg6,%rc3;\n\
	mul.lo.u32	%rg8,%rc2,%rnc1;\n\
	add.u32	%rg9,%rc1,%rg8;\n\
	add.u32	%rg10,%rg9,%rg7;");
	//compute blocksize
	asm volatile("mul.lo.u32	%rg11,%rnt1,%rnt2;\n\
	mul.lo.u32	%rg12,%rg11,%rnt3;\n\
	mul.lo.u32	%rg13,%rg10,%rg12;\n\
	add.u32	%rgtid,%rg13,%rbtid;");
	asm volatile(".reg .u64	%rdf<3>;\n\
	.reg .u64	%rpf<3>;\n\
	.reg .u64	%rdd<3>;\n\
	.reg .u64	%rpd<3>;\n\
	.reg .pred	%p_<10>;\n\
	.reg .f32	%f_<3>;\n\
	.reg .s32	%d_<3>;\n\
	.reg .u64	%offset;\n\
	.reg .u64	%g_distance;\n\
	.reg .u32	%loop;\n\
	.reg .u32	%pass;\n\
	.reg .u64	%distance;\n\
	.reg .u64	%distance_mem;\n\
	.reg .u64	%gfp;\n\
	.reg .u64	%gdp;\n\
	.reg .u32	%counter;");
	asm volatile(".reg .u32	%rd<15>;\n\
	.reg .f32	%rf<16>;\n\
	.reg .f32	%rgf<32>;\n\
	.reg .u32	%rgd<32>;\n\
	.reg .b32	%rb<5>;\n\
	.reg .f64    %rdouble<8>;\n\
	.reg .f64	%rmadf<3>;\n\
	.reg .s32	%ropd3;\n\
	.reg .b32	%ropb<2>;\n\
	mov .b32	%ropb0,652;\n\
	mov .b32	%ropb1,256;\n\
	mov .u64	%distance,20585;\n\
	mov .u64	%distance_mem,262144;\n\
	mov .u64	%g_distance,65536;\n\
	mov .s32	%ropd3,5786;\n\
	mov .f64	%rmadf1,19.5923428225;\n\
	mov .f64     %rmadf2,23.9162632576;\n\
	mov .f64     %rmadf0,86.2923868728;");

	asm volatile(".reg .u64    %c_base;\n\
	.reg .u64    %c_offset;\n\
	.reg .u64    %c_mem;\n\
	mov.u64  %c_base,kconst;\n\
	mul.wide.u32   %c_offset,%rbtid,4;\n\
	add.u64  %c_mem,%c_base,%c_offset;");

	//get the loop number
	asm volatile("ld.param.u32   %loop, [_Z6kernelPfS_S_PiS0_S0_iS_S0__param_6];");//loop

	asm volatile("mul.wide.u32	%offset,%rgtid,4;");//get the address offset
//d0 data

	asm volatile("ld.param.u64	%rpd0, [_Z6kernelPfS_S_PiS0_S0_iS_S0__param_0];\n\
	add.u64	%rdd0,%rpd0,%offset;");
//d1 data

	asm volatile("ld.param.u64	%rpd1, [_Z6kernelPfS_S_PiS0_S0_iS_S0__param_1];\n\
	add.u64	%rdd1,%rpd1,%offset;");
//d2 data

	asm volatile("ld.param.u64	%rpd2, [_Z6kernelPfS_S_PiS0_S0_iS_S0__param_2];\n\
	add.u64	%rdd2,%rpd2,%offset;");
//f0 data

	asm volatile("ld.param.u64	%rpf0, [_Z6kernelPfS_S_PiS0_S0_iS_S0__param_3];\n\
	add.u64	%rdf0,%rpf0,%offset;");
//f1 data

	asm volatile("ld.param.u64	%rpf1, [_Z6kernelPfS_S_PiS0_S0_iS_S0__param_4];\n\
	add.u64	%rdf1,%rpf1,%offset;");
//f2 data

	asm volatile("ld.param.u64	%rpf2, [_Z6kernelPfS_S_PiS0_S0_iS_S0__param_5];\n\
	add.u64	%rdf2,%rpf2,%offset;");

	asm volatile("ld.param.u64	%gfp, [_Z6kernelPfS_S_PiS0_S0_iS_S0__param_7];\n\
	add.u64	%gfp,%gfp,%offset;");

	asm volatile("ld.param.u64	%gdp, [_Z6kernelPfS_S_PiS0_S0_iS_S0__param_8];\n\
	add.u64	%gdp,%gdp,%offset;");
	
//load the values of parameters
	asm volatile("ld.global.s32	%d_0,[%rdd0+0];");
	asm volatile("ld.global.s32	%d_1,[%rdd1+0];");
	asm volatile("ld.global.f32	%f_0,[%rdf0+0];");
	asm volatile("ld.global.f32	%f_1,[%rdf1+0];");


	asm volatile("mul.wide.u32	%distance,20585,4;\n\
	add.u64	%rdd0,%rdd0,%distance;\n\
	add.u64	%rdd1,%rdd1,%distance;\n\
	add.u64	%rdf0,%rdf0,%distance;\n\
	add.u64	%rdf1,%rdf1,%distance;");
	asm volatile("mov.u32	%counter,0;");
	asm volatile("$BB_LABEL0:");
	// instruction mix
	asm volatile("mov.u32 %pass,0;\n\
	fma.rn.f64	%rdouble0, %rmadf0, %rmadf1, %rmadf2;\n\
	cvt.rn.f32.f64     %rf0,%rdouble0;\n\
	lg2.approx.f32	%rf1, %f_1;\n\
	mad.lo.s32	%rd0, %d_0, %d_1, %ropd3;\n\
	mad.rn.f64	%rdouble1, %rmadf0, %rmadf1, %rmadf2;\n\
	cvt.rn.f32.f64     %rf2,%rdouble1;\n\
	lg2.approx.f32	%rf3, %f_0;\n\
	sqrt.rn.f32	%rf4, %f_1;\n\
	st.global.f32  [%gfp+0],%rgf19;\n\
	add.u64 %gfp,%gfp,%g_distance;\n\
	mul.f32	%rf5, %rf0, %rf1;\n\
	ld.const.f32	%rf0, [%c_mem+0];\n\
	cvt.f64.f32     %rmadf0, %rf0;\n\
	cvt.f64.f32     %rmadf1, %rf1;\n\
	cvt.f64.f32     %rmadf2, %rf2;\n\
	fma.rn.f64	%rdouble2, %rmadf0, %rmadf1, %rmadf2;\n\
	cvt.rn.f32.f64     %rf6,%rdouble2;\n\
	ld.const.f32	%rf3, [%c_mem+0];\n\
	xor.b32	%rb0, %ropb0, %ropb1;\n\
	ld.global.f32	%rgf11, [%gfp+0];\n\
	add.u64 %gfp,%gfp,%g_distance;\n\
	ld.global.f32	%rgf17, [%gfp+0];\n\
	add.u64 %gfp,%gfp,%g_distance;\n\
	mad.lo.s32	%rd1, %d_0, %d_1, %ropd3;\n\
	mad.lo.s32	%rd2, %d_0, %d_1, %ropd3;\n\
	ld.global.f32	%rgf8, [%gfp+0];\n\
	add.u64 %gfp,%gfp,%g_distance;\n\
	cvt.f64.f32     %rmadf0, %rf1;\n\
	cvt.f64.f32     %rmadf1, %rf2;\n\
	cvt.f64.f32     %rmadf2, %rf3;\n\
	mad.rn.f64	%rdouble3, %rmadf0, %rmadf1, %rmadf2;\n\
	cvt.rn.f32.f64     %rf7,%rdouble3;\n\
	mul.lo.s32	%rd3, %d_0, %d_1;\n\
	ld.const.f32	%rf1, [%c_mem+0];\n\
	neg.s32	%rd4, %d_0;\n\
	not.b32	%rb1, %ropb0;\n\
	cvt.f64.f32     %rmadf0, %rf2;\n\
	cvt.f64.f32     %rmadf1, %rf3;\n\
	cvt.f64.f32     %rmadf2, %rf4;\n\
	fma.rn.f64	%rdouble4, %rmadf0, %rmadf1, %rmadf2;\n\
	cvt.rn.f32.f64     %rf8,%rdouble4;\n\
	mad.lo.s32	%rd5, %rd0, %rd1, %rd2;\n\
	ld.global.f32	%rgf2, [%gfp+0];\n\
	add.u64 %gfp,%gfp,%g_distance;\n\
	cvt.f64.f32     %rmadf0, %rf3;\n\
	cvt.f64.f32     %rmadf1, %rf4;\n\
	cvt.f64.f32     %rmadf2, %rf5;\n\
	mad.rn.f64	%rdouble5, %rmadf0, %rmadf1, %rmadf2;\n\
	cvt.rn.f32.f64     %rf9,%rdouble5;\n\
	and.b32	%rb2, %ropb0, %ropb1;\n\
	cvt.f64.f32     %rmadf0, %rf4;\n\
	cvt.f64.f32     %rmadf1, %rf5;\n\
	cvt.f64.f32     %rmadf2, %rf6;\n\
	mad.rn.f64	%rdouble6, %rmadf0, %rmadf1, %rmadf2;\n\
	cvt.rn.f32.f64     %rf10,%rdouble6;\n\
	lg2.approx.f32	%rf11, %rf6;\n\
	cos.approx.f32	%rf12, %rf7;\n\
	rem.s32	%rd6, %rd1, %d_0;\n\
	rem.s32	%rd7, %rd2, %d_0;\n\
	mad.lo.s32	%rd8, %rd3, %rd4, %rd5;\n\
	cvt.f64.f32     %rmadf0, %rf7;\n\
	cvt.f64.f32     %rmadf1, %rf8;\n\
	cvt.f64.f32     %rmadf2, %rf9;\n\
	fma.rn.f64	%rdouble7, %rmadf0, %rmadf1, %rmadf2;\n\
	cvt.rn.f32.f64     %rf13,%rdouble7;\n\
	xor.b32	%rb3, %ropb0, %ropb1;\n\
	ex2.approx.f32	%rf14, %rf9;\n\
	ld.global.f32	%rgf31, [%gfp+0];\n\
	add.u64 %gfp,%gfp,%g_distance;\n\
	ld.global.f32	%rgf25, [%gfp+0];\n\
	add.u64 %gfp,%gfp,%g_distance;\n\
	shl.b32	%rb4, %ropb0, %ropb1;\n\
	max.s32	%rd9, %rd4, %rd5;\n\
	sad.s32	%rd10, %rd5, %rd6, %rd7;\n\
	ld.global.f32	%rgf20, [%gfp+0];\n\
	add.u64 %gfp,%gfp,%g_distance;\n\
	div.s32	%rd11, %rd6, %d_0;\n\
	mad.lo.s32	%rd12, %rd7, %rd8, %rd9;\n\
	min.s32	%rd13, %rd8, %rd9;\n\
	");
	asm volatile("add.u32	%counter,%counter,1;\n\
	setp.ge.u32	 %p_0, %counter, %loop;\n\
	@!%p_0	bra	$BB_LABEL0;");

	asm volatile("st.global.f32 [%rdf2+0],%rf15;\n\
	st.global.s32 [%rdd2+0],%rd14;");
	asm volatile("add.u64	%rdf2,%rdf2,%distance;\n\
	add.u64	%rdd2,%rdd2,%distance;");

	//end basic block;
	asm volatile("$BB_LABEL1:");
	asm volatile("exit;");
}

	.text
	.file	"2018-8-20-7-20-52-1537453252-699299-13834.hip"
	.globl	_Z17__device_stub__kkPf         # -- Begin function _Z17__device_stub__kkPf
	.type	_Z17__device_stub__kkPf,@function
_Z17__device_stub__kkPf:                # @_Z17__device_stub__kkPf
	.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	$64, %rsp
	.cfi_def_cfa_offset 112
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	24(%rsp), %rax
	movq	%rdi, (%rax)
	movq	%rsp, %rbx
	movq	%rax, (%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z2kkPf, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$80, %rsp
	.cfi_adjust_cfa_offset -80
	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	_Z17__device_stub__kkPf, .Lfunc_end0-_Z17__device_stub__kkPf
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function _Z7uniformff
.LCPI1_0:
	.long	0x30000000                      # float 4.65661287E-10
	.text
	.globl	_Z7uniformff
	.type	_Z7uniformff,@function
_Z7uniformff:                           # @_Z7uniformff
	.cfi_startproc
# %bb.0:
	pushq	%rax
	.cfi_def_cfa_offset 16
	movss	%xmm0, 4(%rsp)                  # 4-byte Spill
	subss	%xmm0, %xmm1
	movss	%xmm1, (%rsp)                   # 4-byte Spill
	callq	rand
	xorps	%xmm0, %xmm0
	cvtsi2ss	%eax, %xmm0
	mulss	(%rsp), %xmm0                   # 4-byte Folded Reload
	mulss	.LCPI1_0(%rip), %xmm0
	addss	4(%rsp), %xmm0                  # 4-byte Folded Reload
	popq	%rax
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	_Z7uniformff, .Lfunc_end1-_Z7uniformff
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function main
.LCPI2_0:
	.long	0x461c3c00                      # float 9999
.LCPI2_1:
	.long	0x30000000                      # float 4.65661287E-10
.LCPI2_2:
	.long	0x3f800000                      # float 1
	.text
	.globl	main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r12
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	subq	$72, %rsp
	.cfi_def_cfa_offset 112
	.cfi_offset %rbx, -40
	.cfi_offset %r12, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movl	$4, %edi
	callq	srand
	movl	$1, %edi
	movl	$1399780, %esi                  # imm = 0x155BE4
	callq	calloc@PLT
	movq	%rax, %rbx
	leaq	64(%rsp), %rdi
	movl	$1399780, %esi                  # imm = 0x155BE4
	callq	hipMalloc
	xorl	%r14d, %r14d
.LBB2_1:                                # =>This Inner Loop Header: Depth=1
	callq	rand
	xorps	%xmm0, %xmm0
	cvtsi2ss	%eax, %xmm0
	mulss	.LCPI2_0(%rip), %xmm0
	mulss	.LCPI2_1(%rip), %xmm0
	addss	.LCPI2_2(%rip), %xmm0
	movss	%xmm0, (%rbx,%r14,4)
	incq	%r14
	cmpq	$349945, %r14                   # imm = 0x556F9
	jne	.LBB2_1
# %bb.2:
	movq	64(%rsp), %rdi
	movl	$1399780, %edx                  # imm = 0x155BE4
	movq	%rbx, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movl	$1, %edi
	movl	$1399780, %esi                  # imm = 0x155BE4
	callq	calloc@PLT
	movq	%rax, %rbx
	leaq	56(%rsp), %rdi
	movl	$1399780, %esi                  # imm = 0x155BE4
	callq	hipMalloc
	xorl	%r14d, %r14d
.LBB2_3:                                # =>This Inner Loop Header: Depth=1
	callq	rand
	xorps	%xmm0, %xmm0
	cvtsi2ss	%eax, %xmm0
	mulss	.LCPI2_0(%rip), %xmm0
	mulss	.LCPI2_1(%rip), %xmm0
	addss	.LCPI2_2(%rip), %xmm0
	movss	%xmm0, (%rbx,%r14,4)
	incq	%r14
	cmpq	$349945, %r14                   # imm = 0x556F9
	jne	.LBB2_3
# %bb.4:
	movq	56(%rsp), %rdi
	movl	$1399780, %edx                  # imm = 0x155BE4
	movq	%rbx, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movl	$1, %edi
	movl	$1399780, %esi                  # imm = 0x155BE4
	callq	calloc@PLT
	movq	%rax, %rbx
	leaq	16(%rsp), %rdi
	movl	$1399780, %esi                  # imm = 0x155BE4
	callq	hipMalloc
	xorl	%r14d, %r14d
.LBB2_5:                                # =>This Inner Loop Header: Depth=1
	callq	rand
	xorps	%xmm0, %xmm0
	cvtsi2ss	%eax, %xmm0
	mulss	.LCPI2_0(%rip), %xmm0
	mulss	.LCPI2_1(%rip), %xmm0
	addss	.LCPI2_2(%rip), %xmm0
	movss	%xmm0, (%rbx,%r14,4)
	incq	%r14
	cmpq	$349945, %r14                   # imm = 0x556F9
	jne	.LBB2_5
# %bb.6:
	movq	16(%rsp), %rdi
	movl	$1399780, %edx                  # imm = 0x155BE4
	movq	%rbx, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movl	$1, %edi
	movl	$536870912, %esi                # imm = 0x20000000
	callq	calloc@PLT
	movq	%rax, %r14
	leaq	48(%rsp), %rdi
	movl	$536870912, %esi                # imm = 0x20000000
	callq	hipMalloc
	xorl	%r15d, %r15d
.LBB2_7:                                # =>This Inner Loop Header: Depth=1
	callq	rand
	xorps	%xmm0, %xmm0
	cvtsi2ss	%eax, %xmm0
	mulss	.LCPI2_0(%rip), %xmm0
	mulss	.LCPI2_1(%rip), %xmm0
	addss	.LCPI2_2(%rip), %xmm0
	movss	%xmm0, (%r14,%r15,4)
	incq	%r15
	cmpq	$134217728, %r15                # imm = 0x8000000
	jne	.LBB2_7
# %bb.8:
	movq	48(%rsp), %rdi
	movl	$536870912, %edx                # imm = 0x20000000
	movq	%r14, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movl	$1, %edi
	movl	$1399780, %esi                  # imm = 0x155BE4
	callq	calloc@PLT
	movq	%rax, %r14
	leaq	40(%rsp), %rdi
	movl	$1399780, %esi                  # imm = 0x155BE4
	callq	hipMalloc
	xorl	%r15d, %r15d
.LBB2_9:                                # =>This Inner Loop Header: Depth=1
	callq	rand
	xorps	%xmm0, %xmm0
	cvtsi2ss	%eax, %xmm0
	mulss	.LCPI2_0(%rip), %xmm0
	mulss	.LCPI2_1(%rip), %xmm0
	addss	.LCPI2_2(%rip), %xmm0
	cvttss2si	%xmm0, %eax
	movl	%eax, (%r14,%r15,4)
	incq	%r15
	cmpq	$349945, %r15                   # imm = 0x556F9
	jne	.LBB2_9
# %bb.10:
	movq	40(%rsp), %rdi
	movl	$1399780, %edx                  # imm = 0x155BE4
	movq	%r14, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movl	$1, %edi
	movl	$1399780, %esi                  # imm = 0x155BE4
	callq	calloc@PLT
	movq	%rax, %r14
	leaq	32(%rsp), %rdi
	movl	$1399780, %esi                  # imm = 0x155BE4
	callq	hipMalloc
	xorl	%r15d, %r15d
.LBB2_11:                               # =>This Inner Loop Header: Depth=1
	callq	rand
	xorps	%xmm0, %xmm0
	cvtsi2ss	%eax, %xmm0
	mulss	.LCPI2_0(%rip), %xmm0
	mulss	.LCPI2_1(%rip), %xmm0
	addss	.LCPI2_2(%rip), %xmm0
	cvttss2si	%xmm0, %eax
	movl	%eax, (%r14,%r15,4)
	incq	%r15
	cmpq	$349945, %r15                   # imm = 0x556F9
	jne	.LBB2_11
# %bb.12:
	movq	32(%rsp), %rdi
	movl	$1399780, %edx                  # imm = 0x155BE4
	movq	%r14, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movl	$1, %edi
	movl	$1399780, %esi                  # imm = 0x155BE4
	callq	calloc@PLT
	movq	%rax, %r14
	leaq	8(%rsp), %rdi
	movl	$1399780, %esi                  # imm = 0x155BE4
	callq	hipMalloc
	xorl	%r15d, %r15d
.LBB2_13:                               # =>This Inner Loop Header: Depth=1
	callq	rand
	xorps	%xmm0, %xmm0
	cvtsi2ss	%eax, %xmm0
	mulss	.LCPI2_0(%rip), %xmm0
	mulss	.LCPI2_1(%rip), %xmm0
	addss	.LCPI2_2(%rip), %xmm0
	cvttss2si	%xmm0, %eax
	movl	%eax, (%r14,%r15,4)
	incq	%r15
	cmpq	$349945, %r15                   # imm = 0x556F9
	jne	.LBB2_13
# %bb.14:
	movq	8(%rsp), %rdi
	movl	$1399780, %edx                  # imm = 0x155BE4
	movq	%r14, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movl	$1, %edi
	movl	$536870912, %esi                # imm = 0x20000000
	callq	calloc@PLT
	movq	%rax, %r15
	leaq	24(%rsp), %rdi
	movl	$536870912, %esi                # imm = 0x20000000
	callq	hipMalloc
	xorl	%r12d, %r12d
.LBB2_15:                               # =>This Inner Loop Header: Depth=1
	callq	rand
	xorps	%xmm0, %xmm0
	cvtsi2ss	%eax, %xmm0
	mulss	.LCPI2_0(%rip), %xmm0
	mulss	.LCPI2_1(%rip), %xmm0
	addss	.LCPI2_2(%rip), %xmm0
	cvttss2si	%xmm0, %eax
	movl	%eax, (%r15,%r12,4)
	incq	%r12
	cmpq	$134217728, %r12                # imm = 0x8000000
	jne	.LBB2_15
# %bb.16:
	movq	24(%rsp), %rdi
	movl	$536870912, %edx                # imm = 0x20000000
	movq	%r15, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movabsq	$4294967411, %rdx               # imm = 0x100000073
	leaq	64(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB2_18
# %bb.17:
	movq	64(%rsp), %rdi
	movq	56(%rsp), %rsi
	movq	16(%rsp), %rdx
	movq	40(%rsp), %rcx
	movq	32(%rsp), %r8
	movq	8(%rsp), %r9
	subq	$8, %rsp
	.cfi_adjust_cfa_offset 8
	pushq	32(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	64(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	$16
	.cfi_adjust_cfa_offset 8
	callq	_Z21__device_stub__kernelPfS_S_PiS0_S0_iS_S0_
	addq	$32, %rsp
	.cfi_adjust_cfa_offset -32
.LBB2_18:
	movq	8(%rsp), %rsi
	movl	$1399780, %edx                  # imm = 0x155BE4
	movq	%r14, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movq	16(%rsp), %rsi
	movl	$1399780, %edx                  # imm = 0x155BE4
	movq	%rbx, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	xorl	%eax, %eax
	addq	$72, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end2:
	.size	main, .Lfunc_end2-main
	.cfi_endproc
                                        # -- End function
	.globl	_Z21__device_stub__kernelPfS_S_PiS0_S0_iS_S0_ # -- Begin function _Z21__device_stub__kernelPfS_S_PiS0_S0_iS_S0_
	.type	_Z21__device_stub__kernelPfS_S_PiS0_S0_iS_S0_,@function
_Z21__device_stub__kernelPfS_S_PiS0_S0_iS_S0_: # @_Z21__device_stub__kernelPfS_S_PiS0_S0_iS_S0_
	.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	$176, %rsp
	.cfi_def_cfa_offset 224
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	56(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	48(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	40(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	24(%rsp), %rcx
	movq	%r8, (%rcx)
	leaq	16(%rsp), %r8
	movq	%r9, (%r8)
	leaq	96(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r8, 40(%rbx)
	leaq	224(%rsp), %rax
	movq	%rax, 48(%rbx)
	leaq	232(%rsp), %rax
	movq	%rax, 56(%rbx)
	leaq	240(%rsp), %rax
	movq	%rax, 64(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%rsp), %r15
	leaq	8(%rsp), %r12
	movq	%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	$_Z6kernelPfS_S_PiS0_S0_iS_S0_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$192, %rsp
	.cfi_adjust_cfa_offset -192
	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_end3:
	.size	_Z21__device_stub__kernelPfS_S_PiS0_S0_iS_S0_, .Lfunc_end3-_Z21__device_stub__kernelPfS_S_PiS0_S0_iS_S0_
	.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:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -16
	movq	__hip_gpubin_handle(%rip), %rbx
	testq	%rbx, %rbx
	jne	.LBB4_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB4_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z2kkPf, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z6kernelPfS_S_PiS0_S0_iS_S0_, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$0, 8(%rsp)
	movl	$1, (%rsp)
	movl	$kconst, %esi
	movl	$.L__unnamed_3, %edx
	movl	$.L__unnamed_3, %ecx
	movl	$460, %r9d                      # imm = 0x1CC
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	callq	__hipRegisterVar
	movl	$__hip_module_dtor, %edi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end4:
	.size	__hip_module_ctor, .Lfunc_end4-__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	.LBB5_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
.LBB5_2:
	retq
.Lfunc_end5:
	.size	__hip_module_dtor, .Lfunc_end5-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	kconst,@object                  # @kconst
	.local	kconst
	.comm	kconst,460,16
	.type	_Z2kkPf,@object                 # @_Z2kkPf
	.section	.rodata,"a",@progbits
	.globl	_Z2kkPf
	.p2align	3, 0x0
_Z2kkPf:
	.quad	_Z17__device_stub__kkPf
	.size	_Z2kkPf, 8

	.type	_Z6kernelPfS_S_PiS0_S0_iS_S0_,@object # @_Z6kernelPfS_S_PiS0_S0_iS_S0_
	.globl	_Z6kernelPfS_S_PiS0_S0_iS_S0_
	.p2align	3, 0x0
_Z6kernelPfS_S_PiS0_S0_iS_S0_:
	.quad	_Z21__device_stub__kernelPfS_S_PiS0_S0_iS_S0_
	.size	_Z6kernelPfS_S_PiS0_S0_iS_S0_, 8

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

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z6kernelPfS_S_PiS0_S0_iS_S0_"
	.size	.L__unnamed_2, 30

	.type	.L__unnamed_3,@object           # @2
.L__unnamed_3:
	.asciz	"kconst"
	.size	.L__unnamed_3, 7

	.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 _Z17__device_stub__kkPf
	.addrsig_sym _Z21__device_stub__kernelPfS_S_PiS0_S0_iS_S0_
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym kconst
	.addrsig_sym _Z2kkPf
	.addrsig_sym _Z6kernelPfS_S_PiS0_S0_iS_S0_
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z2kkPf                 ; -- Begin function _Z2kkPf
	.globl	_Z2kkPf
	.p2align	8
	.type	_Z2kkPf,@function
_Z2kkPf:                                ; @_Z2kkPf
; %bb.0:
	s_getpc_b64 s[2:3]
	s_add_u32 s2, s2, kconst@rel32@lo+4
	s_addc_u32 s3, s3, kconst@rel32@hi+12
	s_load_b32 s2, s[2:3], 0x0
	s_load_b64 s[0:1], s[0:1], 0x0
	s_waitcnt lgkmcnt(0)
	v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s2
	global_store_b32 v0, v1, s[0:1]
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z2kkPf
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 8
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 2
		.amdhsa_next_free_sgpr 4
		.amdhsa_reserve_vcc 0
		.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	_Z2kkPf, .Lfunc_end0-_Z2kkPf
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 68
; NumSgprs: 4
; NumVgprs: 2
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 0
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 4
; NumVGPRsForWavesPerEU: 2
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.protected	_Z6kernelPfS_S_PiS0_S0_iS_S0_ ; -- Begin function _Z6kernelPfS_S_PiS0_S0_iS_S0_
	.globl	_Z6kernelPfS_S_PiS0_S0_iS_S0_
	.p2align	8
	.type	_Z6kernelPfS_S_PiS0_S0_iS_S0_,@function
_Z6kernelPfS_S_PiS0_S0_iS_S0_:          ; @_Z6kernelPfS_S_PiS0_S0_iS_S0_
; %bb.0:
	;;#ASMSTART
	.reg .u32	%rt<4>;
	.reg .u32 %rnt<4>;
	.reg .u32 %rc<4>;
	.reg .u32 %rnc<4>;
	.reg .u32 %rg<14>;
	.reg .u32 %rgtid;
	.reg .u32 %rbtid;
	;;#ASMEND
	;;#ASMSTART
	mov.u32	%rt1,%tid.x;
	mov.u32	%rt2,%tid.y;
	mov.u32	%rt3,%tid.z;
	mov.u32	%rnt1,%ntid.x;
	mov.u32	%rnt2,%ntid.y;
	mov.u32	%rnt3,%ntid.z;
	mov.u32	%rc1,%ctaid.x;
	mov.u32	%rc2,%ctaid.y;
	mov.u32	%rc3,%ctaid.z;
	mov.u32	%rnc1,%nctaid.x;
	mov.u32	%rnc2,%nctaid.y;
	mov.u32	%rnc3,%nctaid.z;
	;;#ASMEND
	;;#ASMSTART
	mul.lo.u32	%rg1,%rnt1,%rnt2;
	mul.lo.u32	%rg2,%rg1,%rt3;
	mul.lo.u32	%rg3,%rt2,%rnt1;
	add.u32	%rg4,%rt1,%rg3;
	add.u32	%rg5,%rg4,%rg2;
	mov.u32 %rbtid,%rg5;
	;;#ASMEND
	;;#ASMSTART
	mul.lo.u32	%rg6,%rnc1,%rnc2;
	mul.lo.u32	%rg7,%rg6,%rc3;
	mul.lo.u32	%rg8,%rc2,%rnc1;
	add.u32	%rg9,%rc1,%rg8;
	add.u32	%rg10,%rg9,%rg7;
	;;#ASMEND
	;;#ASMSTART
	mul.lo.u32	%rg11,%rnt1,%rnt2;
	mul.lo.u32	%rg12,%rg11,%rnt3;
	mul.lo.u32	%rg13,%rg10,%rg12;
	add.u32	%rgtid,%rg13,%rbtid;
	;;#ASMEND
	;;#ASMSTART
	.reg .u64	%rdf<3>;
	.reg .u64	%rpf<3>;
	.reg .u64	%rdd<3>;
	.reg .u64	%rpd<3>;
	.reg .pred	%p_<10>;
	.reg .f32	%f_<3>;
	.reg .s32	%d_<3>;
	.reg .u64	%offset;
	.reg .u64	%g_distance;
	.reg .u32	%loop;
	.reg .u32	%pass;
	.reg .u64	%distance;
	.reg .u64	%distance_mem;
	.reg .u64	%gfp;
	.reg .u64	%gdp;
	.reg .u32	%counter;
	;;#ASMEND
	;;#ASMSTART
	.reg .u32	%rd<15>;
	.reg .f32	%rf<16>;
	.reg .f32	%rgf<32>;
	.reg .u32	%rgd<32>;
	.reg .b32	%rb<5>;
	.reg .f64    %rdouble<8>;
	.reg .f64	%rmadf<3>;
	.reg .s32	%ropd3;
	.reg .b32	%ropb<2>;
	mov .b32	%ropb0,652;
	mov .b32	%ropb1,256;
	mov .u64	%distance,20585;
	mov .u64	%distance_mem,262144;
	mov .u64	%g_distance,65536;
	mov .s32	%ropd3,5786;
	mov .f64	%rmadf1,19.5923428225;
	mov .f64     %rmadf2,23.9162632576;
	mov .f64     %rmadf0,86.2923868728;
	;;#ASMEND
	;;#ASMSTART
	.reg .u64    %c_base;
	.reg .u64    %c_offset;
	.reg .u64    %c_mem;
	mov.u64  %c_base,kconst;
	mul.wide.u32   %c_offset,%rbtid,4;
	add.u64  %c_mem,%c_base,%c_offset;
	;;#ASMEND
	;;#ASMSTART
	ld.param.u32   %loop, [_Z6kernelPfS_S_PiS0_S0_iS_S0__param_6];
	;;#ASMEND
	;;#ASMSTART
	mul.wide.u32	%offset,%rgtid,4;
	;;#ASMEND
	;;#ASMSTART
	ld.param.u64	%rpd0, [_Z6kernelPfS_S_PiS0_S0_iS_S0__param_0];
	add.u64	%rdd0,%rpd0,%offset;
	;;#ASMEND
	;;#ASMSTART
	ld.param.u64	%rpd1, [_Z6kernelPfS_S_PiS0_S0_iS_S0__param_1];
	add.u64	%rdd1,%rpd1,%offset;
	;;#ASMEND
	;;#ASMSTART
	ld.param.u64	%rpd2, [_Z6kernelPfS_S_PiS0_S0_iS_S0__param_2];
	add.u64	%rdd2,%rpd2,%offset;
	;;#ASMEND
	;;#ASMSTART
	ld.param.u64	%rpf0, [_Z6kernelPfS_S_PiS0_S0_iS_S0__param_3];
	add.u64	%rdf0,%rpf0,%offset;
	;;#ASMEND
	;;#ASMSTART
	ld.param.u64	%rpf1, [_Z6kernelPfS_S_PiS0_S0_iS_S0__param_4];
	add.u64	%rdf1,%rpf1,%offset;
	;;#ASMEND
	;;#ASMSTART
	ld.param.u64	%rpf2, [_Z6kernelPfS_S_PiS0_S0_iS_S0__param_5];
	add.u64	%rdf2,%rpf2,%offset;
	;;#ASMEND
	;;#ASMSTART
	ld.param.u64	%gfp, [_Z6kernelPfS_S_PiS0_S0_iS_S0__param_7];
	add.u64	%gfp,%gfp,%offset;
	;;#ASMEND
	;;#ASMSTART
	ld.param.u64	%gdp, [_Z6kernelPfS_S_PiS0_S0_iS_S0__param_8];
	add.u64	%gdp,%gdp,%offset;
	;;#ASMEND
	;;#ASMSTART
	ld.global.s32	%d_0,[%rdd0+0];
	;;#ASMEND
	;;#ASMSTART
	ld.global.s32	%d_1,[%rdd1+0];
	;;#ASMEND
	;;#ASMSTART
	ld.global.f32	%f_0,[%rdf0+0];
	;;#ASMEND
	;;#ASMSTART
	ld.global.f32	%f_1,[%rdf1+0];
	;;#ASMEND
	;;#ASMSTART
	mul.wide.u32	%distance,20585,4;
	add.u64	%rdd0,%rdd0,%distance;
	add.u64	%rdd1,%rdd1,%distance;
	add.u64	%rdf0,%rdf0,%distance;
	add.u64	%rdf1,%rdf1,%distance;
	;;#ASMEND
	;;#ASMSTART
	mov.u32	%counter,0;
	;;#ASMEND
	;;#ASMSTART
	$BB_LABEL0:
	;;#ASMEND
	;;#ASMSTART
	mov.u32 %pass,0;
	fma.rn.f64	%rdouble0, %rmadf0, %rmadf1, %rmadf2;
	cvt.rn.f32.f64     %rf0,%rdouble0;
	lg2.approx.f32	%rf1, %f_1;
	mad.lo.s32	%rd0, %d_0, %d_1, %ropd3;
	mad.rn.f64	%rdouble1, %rmadf0, %rmadf1, %rmadf2;
	cvt.rn.f32.f64     %rf2,%rdouble1;
	lg2.approx.f32	%rf3, %f_0;
	sqrt.rn.f32	%rf4, %f_1;
	st.global.f32  [%gfp+0],%rgf19;
	add.u64 %gfp,%gfp,%g_distance;
	mul.f32	%rf5, %rf0, %rf1;
	ld.const.f32	%rf0, [%c_mem+0];
	cvt.f64.f32     %rmadf0, %rf0;
	cvt.f64.f32     %rmadf1, %rf1;
	cvt.f64.f32     %rmadf2, %rf2;
	fma.rn.f64	%rdouble2, %rmadf0, %rmadf1, %rmadf2;
	cvt.rn.f32.f64     %rf6,%rdouble2;
	ld.const.f32	%rf3, [%c_mem+0];
	xor.b32	%rb0, %ropb0, %ropb1;
	ld.global.f32	%rgf11, [%gfp+0];
	add.u64 %gfp,%gfp,%g_distance;
	ld.global.f32	%rgf17, [%gfp+0];
	add.u64 %gfp,%gfp,%g_distance;
	mad.lo.s32	%rd1, %d_0, %d_1, %ropd3;
	mad.lo.s32	%rd2, %d_0, %d_1, %ropd3;
	ld.global.f32	%rgf8, [%gfp+0];
	add.u64 %gfp,%gfp,%g_distance;
	cvt.f64.f32     %rmadf0, %rf1;
	cvt.f64.f32     %rmadf1, %rf2;
	cvt.f64.f32     %rmadf2, %rf3;
	mad.rn.f64	%rdouble3, %rmadf0, %rmadf1, %rmadf2;
	cvt.rn.f32.f64     %rf7,%rdouble3;
	mul.lo.s32	%rd3, %d_0, %d_1;
	ld.const.f32	%rf1, [%c_mem+0];
	neg.s32	%rd4, %d_0;
	not.b32	%rb1, %ropb0;
	cvt.f64.f32     %rmadf0, %rf2;
	cvt.f64.f32     %rmadf1, %rf3;
	cvt.f64.f32     %rmadf2, %rf4;
	fma.rn.f64	%rdouble4, %rmadf0, %rmadf1, %rmadf2;
	cvt.rn.f32.f64     %rf8,%rdouble4;
	mad.lo.s32	%rd5, %rd0, %rd1, %rd2;
	ld.global.f32	%rgf2, [%gfp+0];
	add.u64 %gfp,%gfp,%g_distance;
	cvt.f64.f32     %rmadf0, %rf3;
	cvt.f64.f32     %rmadf1, %rf4;
	cvt.f64.f32     %rmadf2, %rf5;
	mad.rn.f64	%rdouble5, %rmadf0, %rmadf1, %rmadf2;
	cvt.rn.f32.f64     %rf9,%rdouble5;
	and.b32	%rb2, %ropb0, %ropb1;
	cvt.f64.f32     %rmadf0, %rf4;
	cvt.f64.f32     %rmadf1, %rf5;
	cvt.f64.f32     %rmadf2, %rf6;
	mad.rn.f64	%rdouble6, %rmadf0, %rmadf1, %rmadf2;
	cvt.rn.f32.f64     %rf10,%rdouble6;
	lg2.approx.f32	%rf11, %rf6;
	cos.approx.f32	%rf12, %rf7;
	rem.s32	%rd6, %rd1, %d_0;
	rem.s32	%rd7, %rd2, %d_0;
	mad.lo.s32	%rd8, %rd3, %rd4, %rd5;
	cvt.f64.f32     %rmadf0, %rf7;
	cvt.f64.f32     %rmadf1, %rf8;
	cvt.f64.f32     %rmadf2, %rf9;
	fma.rn.f64	%rdouble7, %rmadf0, %rmadf1, %rmadf2;
	cvt.rn.f32.f64     %rf13,%rdouble7;
	xor.b32	%rb3, %ropb0, %ropb1;
	ex2.approx.f32	%rf14, %rf9;
	ld.global.f32	%rgf31, [%gfp+0];
	add.u64 %gfp,%gfp,%g_distance;
	ld.global.f32	%rgf25, [%gfp+0];
	add.u64 %gfp,%gfp,%g_distance;
	shl.b32	%rb4, %ropb0, %ropb1;
	max.s32	%rd9, %rd4, %rd5;
	sad.s32	%rd10, %rd5, %rd6, %rd7;
	ld.global.f32	%rgf20, [%gfp+0];
	add.u64 %gfp,%gfp,%g_distance;
	div.s32	%rd11, %rd6, %d_0;
	mad.lo.s32	%rd12, %rd7, %rd8, %rd9;
	min.s32	%rd13, %rd8, %rd9;
	
	;;#ASMEND
	;;#ASMSTART
	add.u32	%counter,%counter,1;
	setp.ge.u32	 %p_0, %counter, %loop;
	@!%p_0	bra	$BB_LABEL0;
	;;#ASMEND
	;;#ASMSTART
	st.global.f32 [%rdf2+0],%rf15;
	st.global.s32 [%rdd2+0],%rd14;
	;;#ASMEND
	;;#ASMSTART
	add.u64	%rdf2,%rdf2,%distance;
	add.u64	%rdd2,%rdd2,%distance;
	;;#ASMEND
	;;#ASMSTART
	$BB_LABEL1:
	;;#ASMEND
	;;#ASMSTART
	exit;
	;;#ASMEND
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z6kernelPfS_S_PiS0_S0_iS_S0_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 72
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 1
		.amdhsa_next_free_sgpr 1
		.amdhsa_reserve_vcc 0
		.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_end1:
	.size	_Z6kernelPfS_S_PiS0_S0_iS_S0_, .Lfunc_end1-_Z6kernelPfS_S_PiS0_S0_iS_S0_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 3844
; NumSgprs: 0
; NumVgprs: 0
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 0
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 1
; NumVGPRsForWavesPerEU: 1
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.protected	kconst                  ; @kconst
	.type	kconst,@object
	.data
	.globl	kconst
	.p2align	4, 0x0
kconst:
	.long	0x45250454                      ; float 2640.27051
	.long	0x43f864e8                      ; float 496.78833
	.long	0x443cf694                      ; float 755.852783
	.long	0x457e6329                      ; float 4070.19751
	.long	0x45cb72c7                      ; float 6510.34716
	.long	0x44fee493                      ; float 2039.14294
	.long	0x456789e8                      ; float 3704.61914
	.long	0x45f25d5a                      ; float 7755.66894
	.long	0x44e8a852                      ; float 1861.26001
	.long	0x449cae56                      ; float 1253.448
	.long	0x45245783                      ; float 2629.46948
	.long	0x45f8214b                      ; float 7940.16162
	.long	0x45f2bb62                      ; float 7767.42285
	.long	0x45af2d40                      ; float 5605.65625
	.long	0x45a7e83e                      ; float 5373.03027
	.long	0x45733b80                      ; float 3891.71875
	.long	0x43b77ba3                      ; float 366.965912
	.long	0x4602ec3e                      ; float 8379.06055
	.long	0x4575a60c                      ; float 3930.37793
	.long	0x4608bd63                      ; float 8751.34667
	.long	0x45bcdcdc                      ; float 6043.60742
	.long	0x4617493b                      ; float 9682.30761
	.long	0x45c1134d                      ; float 6178.41259
	.long	0x458ea325                      ; float 4564.39307
	.long	0x43bb20c2                      ; float 374.25592
	.long	0x458f0ec3                      ; float 4577.84521
	.long	0x44d0e296                      ; float 1671.08081
	.long	0x4612075c                      ; float 9345.83984
	.long	0x45355344                      ; float 2901.2041
	.long	0x4581d7e8                      ; float 4154.98828
	.long	0x4600418b                      ; float 8208.38574
	.long	0x45880110                      ; float 4352.13281
	.long	0x4615957a                      ; float 9573.36914
	.long	0x455fcbd2                      ; float 3580.73877
	.long	0x43868aae                      ; float 269.083435
	.long	0x43b000fc                      ; float 352.00769
	.long	0x45a5ca72                      ; float 5305.30566
	.long	0x4560c352                      ; float 3596.20752
	.long	0x43db75a8                      ; float 438.919189
	.long	0x4613c142                      ; float 9456.31445
	.long	0x46003b86                      ; float 8206.88086
	.long	0x45e30e78                      ; float 7265.80859
	.long	0x460959d3                      ; float 8790.45605
	.long	0x451a02eb                      ; float 2464.18237
	.long	0x457e0c9c                      ; float 4064.78809
	.long	0x45d58287                      ; float 6832.31591
	.long	0x46190ca8                      ; float 9795.16406
	.long	0x45b4d6e1                      ; float 5786.85986
	.long	0x44fb5078                      ; float 2010.51465
	.long	0x457c2995                      ; float 4034.59888
	.long	0x45cdcb9e                      ; float 6585.45215
	.long	0x460bf6f7                      ; float 8957.74121
	.long	0x45c81a58                      ; float 6403.29296
	.long	0x453fe0fa                      ; float 3070.06104
	.long	0x45c0f9fd                      ; float 6175.24853
	.long	0x44b6a8f0                      ; float 1461.2793
	.long	0x4550881c                      ; float 3336.50684
	.long	0x45d7f668                      ; float 6910.80078
	.long	0x461bbc76                      ; float 9967.11523
	.long	0x455dd62e                      ; float 3549.38623
	.long	0x4522e092                      ; float 2606.03564
	.long	0x44c4f60e                      ; float 1575.68921
	.long	0x459957e7                      ; float 4906.98779
	.long	0x4429d4ba                      ; float 679.323853
	.long	0x45a5f420                      ; float 5310.51563
	.long	0x45c65dfc                      ; float 6347.74805
	.long	0x459ba12d                      ; float 4980.14697
	.long	0x45dd3c73                      ; float 7079.55615
	.long	0x45100f17                      ; float 2304.94312
	.long	0x437d6f6d                      ; float 253.435257
	.long	0x45ce2cd1                      ; float 6597.60205
	.long	0x460f23b7                      ; float 9160.92871
	.long	0x4526082f                      ; float 2656.51147
	.long	0x44bf8ae9                      ; float 1532.34094
	.long	0x461920f7                      ; float 9800.24121
	.long	0x4592ccc2                      ; float 4697.59473
	.long	0x458e2c8a                      ; float 4549.56738
	.long	0x445fd7a4                      ; float 895.369384
	.long	0x42a29291                      ; float 81.2862625
	.long	0x446ea9cc                      ; float 954.653076
	.long	0x461929a3                      ; float 9802.40917
	.long	0x453fcf0d                      ; float 3068.94067
	.long	0x4587ba45                      ; float 4343.28369
	.long	0x45860375                      ; float 4288.43213
	.long	0x4613b92b                      ; float 9454.29199
	.long	0x442bf487                      ; float 687.82074
	.long	0x4605ac69                      ; float 8555.10253
	.long	0x447f2805                      ; float 1020.62531
	.long	0x460ade13                      ; float 8887.51855
	.long	0x44cac9a6                      ; float 1622.30151
	.long	0x45ac69a4                      ; float 5517.20508
	.long	0x452642a0                      ; float 2660.16406
	.long	0x46142bcc                      ; float 9482.94921
	.long	0x460a2894                      ; float 8842.14453
	.long	0x456fa6a2                      ; float 3834.41455
	.long	0x442f3c26                      ; float 700.939819
	.long	0x4609c255                      ; float 8816.583
	.long	0x4587fb86                      ; float 4351.44043
	.long	0x45611091                      ; float 3601.0354
	.long	0x45310b7f                      ; float 2832.71851
	.long	0x45a1d678                      ; float 5178.80859
	.long	0x45541660                      ; float 3393.39844
	.long	0x4556c5e7                      ; float 3436.3689
	.long	0x45120c30                      ; float 2336.76172
	.long	0x45416d86                      ; float 3094.84521
	.long	0x46130720                      ; float 9409.78125
	.long	0x44a733fd                      ; float 1337.62463
	.long	0x45a9fe96                      ; float 5439.82324
	.long	0x45e9eeba                      ; float 7485.84082
	.long	0x45667fb7                      ; float 3687.98218
	.long	0x456176af                      ; float 3607.41772
	.long	0x459b853a                      ; float 4976.65332
	.long	0x45e0b832                      ; float 7191.02441
	.long	0x459ba103                      ; float 4980.12646
	.long	0x453f877d                      ; float 3064.46802
	.size	kconst, 460

	.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 kconst
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 8
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z2kkPf
    .private_segment_fixed_size: 0
    .sgpr_count:     4
    .sgpr_spill_count: 0
    .symbol:         _Z2kkPf.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     2
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         32
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         40
        .size:           8
        .value_kind:     global_buffer
      - .offset:         48
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         56
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         64
        .size:           8
        .value_kind:     global_buffer
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 72
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z6kernelPfS_S_PiS0_S0_iS_S0_
    .private_segment_fixed_size: 0
    .sgpr_count:     0
    .sgpr_spill_count: 0
    .symbol:         _Z6kernelPfS_S_PiS0_S0_iS_S0_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     0
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
