# include <cuda.h>
# include <cuda_runtime.h>
# include <vector_types.h>
# include <device_launch_parameters.h>
# include <cstdio>
# ifdef WIN32
# include <time.h>
# else
# include <sys/time.h>
# endif


__global__ void iterMandel(int niterMax, int n, int* mandel )
{
  int   t_x, t_y, iter;
  float step, zr;
  float2 z0,c0;
  iter = 0;
  t_x = threadIdx.x+blockIdx.x*blockDim.x;
  t_y = threadIdx.y+blockIdx.y*blockDim.y;
  if ((t_x<n) && (t_y<n)) {
    step = 2.5f/(n-1.f);
    c0.x = -2.00 + step*t_x;
    c0.y = -1.25 + step*t_y;
    z0 = c0;
    while ((z0.x*z0.x+z0.y*z0.y < 4.f)&&(iter<niterMax)) {
      iter += 1;
      zr = z0.x*z0.x-z0.y*z0.y+c0.x;
      z0.y = 2.f*z0.x*z0.y+c0.y;
      z0.x = zr;
    }
    mandel[t_x+t_y*n] = iter;
  }
}

/* Petite routine renvoyant un instant
   en seconde. En faisant la différence
   entre deux instants, on peut mesure le temps
   pris par une routine.
   La précision est de l'ordre de la micro-seconde.
*/
float topChrono()
{
# ifdef WIN32
  clock_t chrono;
  chrono = clock();
  return ((float)chrono)/CLOCKS_PER_SEC;
# else
  struct timeval tv;
  gettimeofday(&tv,NULL);
  return tv.tv_sec+1.E-6*tv.tv_usec;
# endif
}

/* Calcul les coordonnees des sommets de la grille */
void compGrid(int n, float*& xh, float*& yh)
{
  xh = new float[n];
  yh = new float[n];
  float step = 2.5/(n-1.);
  for (int i = 0; i < n; i++) {
    xh[i] = -2.+i*step;
    yh[i] = -1.25+i*step;
  }
}

/* Sauve l'ensemble de mandelbrot au format Gnuplot
 */
void saveMandel(int n, const float* xh, const float* yh,
		const int* mandel)
{
  FILE* fich = fopen("Mandel.dat", "w");
  for (int j = 0; j < n; j++) {
    for (int i  = 0; i < n; i++) {
      fprintf(fich,"%f\t%f\t%d\n",yh[i], xh[j], mandel[i*n+j]);
    }
    fprintf(fich, "\n");
  }
  fclose(fich);
}

int main(int nargc, char* argv[])
{
  float t1,t2;
  float *xh, *yh;
  int n = 1280;
  int maxIter = 1000;
  int* gmandel, *cmandel;
  cudaEvent_t start, stop; 
  float time;
  if (nargc > 1) n = atoi(argv[1]);
  if (nargc > 2) maxIter = atoi(argv[2]);

  compGrid(n,xh,yh);

  cudaMalloc((void**)&gmandel, n*n*sizeof(int));
  cmandel = new int[n*n];
  
  cudaEventCreate(&start); cudaEventCreate(&stop);
  dim3 blockDim(16,16);
  dim3 gridDim((n+15)/16, (n+15)/16);
  cudaEventRecord( start, 0 );
  iterMandel<<<gridDim,blockDim>>>(maxIter, n, gmandel);
  cudaMemcpy(cmandel, gmandel, n*n*sizeof(int), cudaMemcpyDeviceToHost);
  cudaEventRecord( stop, 0 ); cudaEventSynchronize( stop );
  cudaFree(gmandel);
  cudaEventElapsedTime( &time, start, stop );
  cudaEventDestroy( start ); cudaEventDestroy( stop );
  printf("Temps de calcul du Bouddha : %f secondes\n",time/1000.);
  saveMandel(n,xh,yh,cmandel);
  delete [] cmandel;
  delete [] xh;
  delete [] yh;
  return EXIT_SUCCESS;
};

	.file	"tmpxft_002e29bb_00000000-6_MandelGPU.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2033:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2033:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z9topChronov
	.type	_Z9topChronov, @function
_Z9topChronov:
.LFB2027:
	.cfi_startproc
	endbr64
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	xorl	%esi, %esi
	movq	%fs:40, %rax
	movq	%rax, 24(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rdi
	call	gettimeofday@PLT
	cvtsi2sdq	16(%rsp), %xmm0
	mulsd	.LC0(%rip), %xmm0
	cvtsi2sdq	8(%rsp), %xmm1
	addsd	%xmm1, %xmm0
	movq	24(%rsp), %rax
	subq	%fs:40, %rax
	je	.L3
	call	__stack_chk_fail@PLT
.L3:
	cvtsd2ss	%xmm0, %xmm0
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2027:
	.size	_Z9topChronov, .-_Z9topChronov
	.globl	_Z8compGridiRPfS0_
	.type	_Z8compGridiRPfS0_, @function
_Z8compGridiRPfS0_:
.LFB2028:
	.cfi_startproc
	endbr64
	movabsq	$2305843009213693950, %rax
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	movslq	%edi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	pushq	%rcx
	.cfi_def_cfa_offset 48
	cmpq	%rbp, %rax
	jb	.L7
	movq	%rbp, %rbx
	salq	$2, %rbp
	movq	%rdx, %r13
	movq	%rsi, %r12
	movq	%rbp, %rdi
	call	_Znam@PLT
	movq	%rbp, %rdi
	movq	%rax, (%r12)
	call	_Znam@PLT
	cvtsi2sdl	%ebx, %xmm0
	xorl	%edx, %edx
	subsd	.LC1(%rip), %xmm0
	movsd	.LC2(%rip), %xmm1
	movq	%rax, 0(%r13)
	movss	.LC3(%rip), %xmm2
	movss	.LC4(%rip), %xmm3
	divsd	%xmm0, %xmm1
	cvtsd2ss	%xmm1, %xmm1
	jmp	.L8
.L7:
	call	__cxa_throw_bad_array_new_length@PLT
.L8:
	cmpl	%edx, %ebx
	jle	.L11
	cvtsi2ssl	%edx, %xmm0
	movq	(%r12), %rcx
	mulss	%xmm1, %xmm0
	movaps	%xmm0, %xmm4
	subss	%xmm3, %xmm0
	subss	%xmm2, %xmm4
	movss	%xmm4, (%rcx,%rdx,4)
	movss	%xmm0, (%rax,%rdx,4)
	incq	%rdx
	jmp	.L8
.L11:
	popq	%rax
	.cfi_def_cfa_offset 40
	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	_Z8compGridiRPfS0_, .-_Z8compGridiRPfS0_
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC5:
	.string	"w"
.LC6:
	.string	"Mandel.dat"
.LC7:
	.string	"%f\t%f\t%d\n"
.LC8:
	.string	"\n"
	.text
	.globl	_Z10saveMandeliPKfS0_PKi
	.type	_Z10saveMandeliPKfS0_PKi, @function
_Z10saveMandeliPKfS0_PKi:
.LFB2029:
	.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
	movq	%rcx, %r14
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movq	%rsi, %r13
	leaq	.LC5(%rip), %rsi
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%rdx, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	xorl	%ebx, %ebx
	subq	$40, %rsp
	.cfi_def_cfa_offset 96
	movl	%edi, 12(%rsp)
	leaq	.LC6(%rip), %rdi
	call	fopen@PLT
	movslq	12(%rsp), %rbp
	movq	%rax, (%rsp)
	salq	$2, %rbp
	leaq	(%r12,%rbp), %r15
.L13:
	cmpl	%ebx, 12(%rsp)
	jle	.L18
	leaq	(%r14,%rbx,4), %r10
	movq	%r12, %r9
.L14:
	movl	(%r10), %ecx
	movq	(%rsp), %rdi
	leaq	.LC7(%rip), %rdx
	movb	$2, %al
	movl	$2, %esi
	movq	%r9, 24(%rsp)
	cvtss2sd	(%r9), %xmm0
	cvtss2sd	0(%r13,%rbx,4), %xmm1
	movq	%r10, 16(%rsp)
	call	__fprintf_chk@PLT
	movq	24(%rsp), %r9
	movq	16(%rsp), %r10
	addq	$4, %r9
	addq	%rbp, %r10
	cmpq	%r15, %r9
	jne	.L14
	movq	(%rsp), %rdi
	leaq	.LC8(%rip), %rdx
	xorl	%eax, %eax
	incq	%rbx
	movl	$2, %esi
	call	__fprintf_chk@PLT
	jmp	.L13
.L18:
	movq	(%rsp), %rdi
	addq	$40, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	jmp	fclose@PLT
	.cfi_endproc
.LFE2029:
	.size	_Z10saveMandeliPKfS0_PKi, .-_Z10saveMandeliPKfS0_PKi
	.globl	_Z32__device_stub__Z10iterMandeliiPiiiPi
	.type	_Z32__device_stub__Z10iterMandeliiPiiiPi, @function
_Z32__device_stub__Z10iterMandeliiPiiiPi:
.LFB2055:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movl	%edi, 12(%rsp)
	leaq	24(%rsp), %rcx
	leaq	32(%rsp), %rdi
	movl	%esi, 8(%rsp)
	leaq	44(%rsp), %rsi
	movq	%rdx, (%rsp)
	leaq	16(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	12(%rsp), %rax
	movl	$1, 40(%rsp)
	movq	%rax, 80(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 32(%rsp)
	movq	%rax, 44(%rsp)
	movl	$1, 52(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L19
	pushq	24(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z10iterMandeliiPi(%rip), %rdi
	pushq	24(%rsp)
	.cfi_def_cfa_offset 144
	movq	60(%rsp), %rcx
	movl	68(%rsp), %r8d
	movq	48(%rsp), %rsi
	movl	56(%rsp), %edx
	leaq	96(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L19:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L21
	call	__stack_chk_fail@PLT
.L21:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2055:
	.size	_Z32__device_stub__Z10iterMandeliiPiiiPi, .-_Z32__device_stub__Z10iterMandeliiPiiiPi
	.globl	_Z10iterMandeliiPi
	.type	_Z10iterMandeliiPi, @function
_Z10iterMandeliiPi:
.LFB2056:
	.cfi_startproc
	endbr64
	jmp	_Z32__device_stub__Z10iterMandeliiPiiiPi
	.cfi_endproc
.LFE2056:
	.size	_Z10iterMandeliiPi, .-_Z10iterMandeliiPi
	.section	.rodata.str1.1
.LC10:
	.string	"Temps de calcul du Bouddha : %f secondes\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2030:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	subq	$88, %rsp
	.cfi_def_cfa_offset 128
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	xorl	%eax, %eax
	cmpl	$1, %edi
	jle	.L30
	movl	%edi, %r12d
	movq	8(%rsi), %rdi
	movq	%rsi, %rbp
	movl	$1000, %r13d
	call	atoi@PLT
	movl	%eax, %ebx
	cmpl	$2, %r12d
	je	.L25
	movq	16(%rbp), %rdi
	call	atoi@PLT
	movl	%eax, %r13d
	jmp	.L25
.L30:
	movl	$1280, %ebx
	movl	$1000, %r13d
.L25:
	movl	%ebx, %ebp
	leaq	16(%rsp), %rdx
	leaq	8(%rsp), %rsi
	movl	%ebx, %edi
	call	_Z8compGridiRPfS0_
	imull	%ebx, %ebp
	leaq	24(%rsp), %rdi
	movslq	%ebp, %rbp
	salq	$2, %rbp
	movq	%rbp, %rsi
	call	cudaMalloc@PLT
	movq	%rbp, %rdi
	call	_Znam@PLT
	leaq	32(%rsp), %rdi
	movq	%rax, %r12
	call	cudaEventCreate@PLT
	leaq	40(%rsp), %rdi
	call	cudaEventCreate@PLT
	movl	$16, %ecx
	xorl	%esi, %esi
	movabsq	$68719476752, %rax
	movq	%rax, 48(%rsp)
	leal	15(%rbx), %eax
	movq	32(%rsp), %rdi
	cltd
	movl	$1, 56(%rsp)
	idivl	%ecx
	movl	$1, 68(%rsp)
	movl	%eax, 60(%rsp)
	movl	%eax, 64(%rsp)
	call	cudaEventRecord@PLT
	movl	56(%rsp), %ecx
	movl	68(%rsp), %esi
	xorl	%r9d, %r9d
	movq	48(%rsp), %rdx
	movq	60(%rsp), %rdi
	xorl	%r8d, %r8d
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L26
	movq	24(%rsp), %rdx
	movl	%ebx, %esi
	movl	%r13d, %edi
	call	_Z32__device_stub__Z10iterMandeliiPiiiPi
.L26:
	movq	24(%rsp), %rsi
	movl	$2, %ecx
	movq	%rbp, %rdx
	movq	%r12, %rdi
	call	cudaMemcpy@PLT
	movq	40(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	40(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
	movq	40(%rsp), %rdx
	movq	32(%rsp), %rsi
	leaq	4(%rsp), %rdi
	call	cudaEventElapsedTime@PLT
	movq	32(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movq	40(%rsp), %rdi
	call	cudaEventDestroy@PLT
	leaq	.LC10(%rip), %rsi
	movl	$2, %edi
	movb	$1, %al
	cvtss2sd	4(%rsp), %xmm0
	divsd	.LC9(%rip), %xmm0
	call	__printf_chk@PLT
	movl	%ebx, %edi
	movq	8(%rsp), %r13
	movq	16(%rsp), %rbp
	movq	%r12, %rcx
	movq	%rbp, %rdx
	movq	%r13, %rsi
	call	_Z10saveMandeliPKfS0_PKi
	movq	%r12, %rdi
	call	_ZdaPv@PLT
	testq	%r13, %r13
	je	.L27
	movq	%r13, %rdi
	call	_ZdaPv@PLT
.L27:
	testq	%rbp, %rbp
	je	.L28
	movq	%rbp, %rdi
	call	_ZdaPv@PLT
.L28:
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L29
	call	__stack_chk_fail@PLT
.L29:
	addq	$88, %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
.LFE2030:
	.size	main, .-main
	.section	.rodata.str1.1
.LC11:
	.string	"_Z10iterMandeliiPi"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2058:
	.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	.LC11(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z10iterMandeliiPi(%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
.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
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC0:
	.long	-1598689907
	.long	1051772663
	.align 8
.LC1:
	.long	0
	.long	1072693248
	.align 8
.LC2:
	.long	0
	.long	1074003968
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC3:
	.long	1073741824
	.align 4
.LC4:
	.long	1067450368
	.section	.rodata.cst8
	.align 8
.LC9:
	.long	0
	.long	1083129856
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:

# include <hip/hip_runtime.h>

# include <hip/hip_vector_types.h>

# include <cstdio>
# ifdef WIN32
# include <time.h>
# else
# include <sys/time.h>
# endif


__global__ void iterMandel(int niterMax, int n, int* mandel )
{
  int   t_x, t_y, iter;
  float step, zr;
  float2 z0,c0;
  iter = 0;
  t_x = threadIdx.x+blockIdx.x*blockDim.x;
  t_y = threadIdx.y+blockIdx.y*blockDim.y;
  if ((t_x<n) && (t_y<n)) {
    step = 2.5f/(n-1.f);
    c0.x = -2.00 + step*t_x;
    c0.y = -1.25 + step*t_y;
    z0 = c0;
    while ((z0.x*z0.x+z0.y*z0.y < 4.f)&&(iter<niterMax)) {
      iter += 1;
      zr = z0.x*z0.x-z0.y*z0.y+c0.x;
      z0.y = 2.f*z0.x*z0.y+c0.y;
      z0.x = zr;
    }
    mandel[t_x+t_y*n] = iter;
  }
}

/* Petite routine renvoyant un instant
   en seconde. En faisant la différence
   entre deux instants, on peut mesure le temps
   pris par une routine.
   La précision est de l'ordre de la micro-seconde.
*/
float topChrono()
{
# ifdef WIN32
  clock_t chrono;
  chrono = clock();
  return ((float)chrono)/CLOCKS_PER_SEC;
# else
  struct timeval tv;
  gettimeofday(&tv,NULL);
  return tv.tv_sec+1.E-6*tv.tv_usec;
# endif
}

/* Calcul les coordonnees des sommets de la grille */
void compGrid(int n, float*& xh, float*& yh)
{
  xh = new float[n];
  yh = new float[n];
  float step = 2.5/(n-1.);
  for (int i = 0; i < n; i++) {
    xh[i] = -2.+i*step;
    yh[i] = -1.25+i*step;
  }
}

/* Sauve l'ensemble de mandelbrot au format Gnuplot
 */
void saveMandel(int n, const float* xh, const float* yh,
		const int* mandel)
{
  FILE* fich = fopen("Mandel.dat", "w");
  for (int j = 0; j < n; j++) {
    for (int i  = 0; i < n; i++) {
      fprintf(fich,"%f\t%f\t%d\n",yh[i], xh[j], mandel[i*n+j]);
    }
    fprintf(fich, "\n");
  }
  fclose(fich);
}

int main(int nargc, char* argv[])
{
  float t1,t2;
  float *xh, *yh;
  int n = 1280;
  int maxIter = 1000;
  int* gmandel, *cmandel;
  hipEvent_t start, stop; 
  float time;
  if (nargc > 1) n = atoi(argv[1]);
  if (nargc > 2) maxIter = atoi(argv[2]);

  compGrid(n,xh,yh);

  hipMalloc((void**)&gmandel, n*n*sizeof(int));
  cmandel = new int[n*n];
  
  hipEventCreate(&start); hipEventCreate(&stop);
  dim3 blockDim(16,16);
  dim3 gridDim((n+15)/16, (n+15)/16);
  hipEventRecord( start, 0 );
  iterMandel<<<gridDim,blockDim>>>(maxIter, n, gmandel);
  hipMemcpy(cmandel, gmandel, n*n*sizeof(int), hipMemcpyDeviceToHost);
  hipEventRecord( stop, 0 ); hipEventSynchronize( stop );
  hipFree(gmandel);
  hipEventElapsedTime( &time, start, stop );
  hipEventDestroy( start ); hipEventDestroy( stop );
  printf("Temps de calcul du Bouddha : %f secondes\n",time/1000.);
  saveMandel(n,xh,yh,cmandel);
  delete [] cmandel;
  delete [] xh;
  delete [] yh;
  return EXIT_SUCCESS;
};

	.text
	.file	"MandelGPU.hip"
	.globl	_Z25__device_stub__iterMandeliiPi # -- Begin function _Z25__device_stub__iterMandeliiPi
	.type	_Z25__device_stub__iterMandeliiPi,@function
_Z25__device_stub__iterMandeliiPi:      # @_Z25__device_stub__iterMandeliiPi
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$96, %rsp
	.cfi_def_cfa_offset 144
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	4(%rsp), %rax
	movl	%edi, (%rax)
	movq	%rsp, %rcx
	movl	%esi, (%rcx)
	leaq	24(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	64(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	movq	%rsi, 16(%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z10iterMandeliiPi, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$112, %rsp
	.cfi_adjust_cfa_offset -112
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z25__device_stub__iterMandeliiPi, .Lfunc_end0-_Z25__device_stub__iterMandeliiPi
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function _Z9topChronov
.LCPI1_0:
	.quad	0x3eb0c6f7a0b5ed8d              # double 9.9999999999999995E-7
	.text
	.globl	_Z9topChronov
	.type	_Z9topChronov,@function
_Z9topChronov:                          # @_Z9topChronov
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$16, %rsp
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -16
	movq	%rsp, %rbx
	movq	%rbx, %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	cvtsi2sdq	(%rbx), %xmm0
	cvtsi2sdq	8(%rbx), %xmm1
	mulsd	.LCPI1_0(%rip), %xmm1
	addsd	%xmm0, %xmm1
	xorps	%xmm0, %xmm0
	cvtsd2ss	%xmm1, %xmm0
	addq	$16, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	_Z9topChronov, .Lfunc_end1-_Z9topChronov
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function _Z8compGridiRPfS0_
.LCPI2_0:
	.quad	0xbff0000000000000              # double -1
.LCPI2_1:
	.quad	0x4004000000000000              # double 2.5
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0
.LCPI2_2:
	.long	0xc0000000                      # float -2
.LCPI2_3:
	.long	0xbfa00000                      # float -1.25
	.text
	.globl	_Z8compGridiRPfS0_
	.type	_Z8compGridiRPfS0_,@function
_Z8compGridiRPfS0_:                     # @_Z8compGridiRPfS0_
	.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	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%rdx, %r14
	movq	%rsi, %rbx
	movl	%edi, %ebp
	movl	%edi, %r12d
	leaq	(,%r12,4), %rax
	testl	%edi, %edi
	movq	$-1, %r15
	cmovnsq	%rax, %r15
	movq	%r15, %rdi
	callq	_Znam
	movq	%rax, (%rbx)
	movq	%r15, %rdi
	callq	_Znam
	movq	%rax, (%r14)
	testl	%ebp, %ebp
	jle	.LBB2_3
# %bb.1:                                # %.lr.ph
	cvtsi2sd	%ebp, %xmm0
	addsd	.LCPI2_0(%rip), %xmm0
	movsd	.LCPI2_1(%rip), %xmm1           # xmm1 = mem[0],zero
	divsd	%xmm0, %xmm1
	xorps	%xmm0, %xmm0
	cvtsd2ss	%xmm1, %xmm0
	movq	(%rbx), %rcx
	xorl	%edx, %edx
	movss	.LCPI2_2(%rip), %xmm1           # xmm1 = mem[0],zero,zero,zero
	movss	.LCPI2_3(%rip), %xmm2           # xmm2 = mem[0],zero,zero,zero
.LBB2_2:                                # =>This Inner Loop Header: Depth=1
	xorps	%xmm3, %xmm3
	cvtsi2ss	%edx, %xmm3
	mulss	%xmm0, %xmm3
	movaps	%xmm3, %xmm4
	addss	%xmm1, %xmm4
	movss	%xmm4, (%rcx,%rdx,4)
	addss	%xmm2, %xmm3
	movss	%xmm3, (%rax,%rdx,4)
	incq	%rdx
	cmpq	%rdx, %r12
	jne	.LBB2_2
.LBB2_3:                                # %._crit_edge
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end2:
	.size	_Z8compGridiRPfS0_, .Lfunc_end2-_Z8compGridiRPfS0_
	.cfi_endproc
                                        # -- End function
	.globl	_Z10saveMandeliPKfS0_PKi        # -- Begin function _Z10saveMandeliPKfS0_PKi
	.type	_Z10saveMandeliPKfS0_PKi,@function
_Z10saveMandeliPKfS0_PKi:               # @_Z10saveMandeliPKfS0_PKi
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$56, %rsp
	.cfi_def_cfa_offset 112
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%rcx, %r14
	movq	%rdx, 40(%rsp)                  # 8-byte Spill
	movq	%rsi, 32(%rsp)                  # 8-byte Spill
	movl	%edi, %ebp
	movl	$.L.str, %edi
	movl	$.L.str.1, %esi
	callq	fopen
	movq	%rax, %rbx
	testl	%ebp, %ebp
	jle	.LBB3_5
# %bb.1:                                # %.preheader.lr.ph
	movl	%ebp, %edx
	leaq	(,%rdx,4), %rcx
	movq	%rcx, 24(%rsp)                  # 8-byte Spill
	xorl	%ecx, %ecx
	movq	%rdx, 8(%rsp)                   # 8-byte Spill
	movq	40(%rsp), %r12                  # 8-byte Reload
	movq	32(%rsp), %rbp                  # 8-byte Reload
	movq	24(%rsp), %r13                  # 8-byte Reload
.LBB3_2:                                # %.lr.ph
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB3_3 Depth 2
	movq	%rcx, 16(%rsp)                  # 8-byte Spill
	movq	%r14, 48(%rsp)                  # 8-byte Spill
	xorl	%r15d, %r15d
.LBB3_3:                                #   Parent Loop BB3_2 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	xorps	%xmm0, %xmm0
	cvtss2sd	(%r12,%r15,4), %xmm0
	movq	16(%rsp), %rax                  # 8-byte Reload
	xorps	%xmm1, %xmm1
	cvtss2sd	(%rbp,%rax,4), %xmm1
	movl	(%r14), %edx
	movl	$.L.str.2, %esi
	movq	%rbx, %rdi
	movb	$2, %al
	callq	fprintf
	incq	%r15
	addq	%r13, %r14
	cmpq	%r15, 8(%rsp)                   # 8-byte Folded Reload
	jne	.LBB3_3
# %bb.4:                                # %._crit_edge
                                        #   in Loop: Header=BB3_2 Depth=1
	movl	$10, %edi
	movq	%rbx, %rsi
	callq	fputc@PLT
	movq	16(%rsp), %rsi                  # 8-byte Reload
	incq	%rsi
	movq	48(%rsp), %r14                  # 8-byte Reload
	addq	$4, %r14
	movq	8(%rsp), %rcx                   # 8-byte Reload
	cmpq	%rcx, %rsi
	movq	%rsi, %rcx
	jne	.LBB3_2
.LBB3_5:                                # %._crit_edge20
	movq	%rbx, %rdi
	addq	$56, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	jmp	fclose                          # TAILCALL
.Lfunc_end3:
	.size	_Z10saveMandeliPKfS0_PKi, .Lfunc_end3-_Z10saveMandeliPKfS0_PKi
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function main
.LCPI4_0:
	.quad	0x408f400000000000              # double 1000
	.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	$56, %rsp
	.cfi_def_cfa_offset 112
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	$1000, %ebp                     # imm = 0x3E8
	movl	$1280, %ebx                     # imm = 0x500
	cmpl	$2, %edi
	jl	.LBB4_3
# %bb.1:
	movq	%rsi, %r14
	movl	%edi, %r15d
	movq	8(%rsi), %rdi
	callq	atoi
	movl	%eax, %ebx
	cmpl	$2, %r15d
	je	.LBB4_3
# %bb.2:
	movq	16(%r14), %rdi
	callq	atoi
	movl	%eax, %ebp
.LBB4_3:                                # %.thread
	leaq	40(%rsp), %rsi
	leaq	32(%rsp), %rdx
	movl	%ebx, %edi
	callq	_Z8compGridiRPfS0_
	movl	%ebx, %r15d
	imull	%r15d, %r15d
	shlq	$2, %r15
	leaq	16(%rsp), %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	movq	%r15, %rdi
	callq	_Znam
	movq	%rax, %r14
	leaq	24(%rsp), %r12
	movq	%r12, %rdi
	callq	hipEventCreate
	leaq	8(%rsp), %rdi
	callq	hipEventCreate
	leal	15(%rbx), %eax
	leal	30(%rbx), %ecx
	testl	%eax, %eax
	cmovnsl	%eax, %ecx
	sarl	$4, %ecx
	movq	%rcx, %r13
	shlq	$32, %r13
	orq	%rcx, %r13
	movq	(%r12), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movabsq	$68719476752, %rdx              # imm = 0x1000000010
	movq	%r13, %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB4_5
# %bb.4:
	movq	16(%rsp), %rdx
	movl	%ebp, %edi
	movl	%ebx, %esi
	callq	_Z25__device_stub__iterMandeliiPi
.LBB4_5:
	movq	16(%rsp), %rsi
	movq	%r14, %rdi
	movq	%r15, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movq	8(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	8(%rsp), %rdi
	callq	hipEventSynchronize
	movq	16(%rsp), %rdi
	callq	hipFree
	movq	24(%rsp), %rsi
	movq	8(%rsp), %rdx
	leaq	52(%rsp), %r15
	movq	%r15, %rdi
	callq	hipEventElapsedTime
	movq	24(%rsp), %rdi
	callq	hipEventDestroy
	movq	8(%rsp), %rdi
	callq	hipEventDestroy
	cvtss2sd	(%r15), %xmm0
	divsd	.LCPI4_0(%rip), %xmm0
	movl	$.L.str.4, %edi
	movb	$1, %al
	callq	printf
	movq	40(%rsp), %r12
	movq	32(%rsp), %r15
	movl	%ebx, %edi
	movq	%r12, %rsi
	movq	%r15, %rdx
	movq	%r14, %rcx
	callq	_Z10saveMandeliPKfS0_PKi
	movq	%r14, %rdi
	callq	_ZdaPv
	testq	%r12, %r12
	je	.LBB4_7
# %bb.6:
	movq	%r12, %rdi
	callq	_ZdaPv
.LBB4_7:
	testq	%r15, %r15
	je	.LBB4_9
# %bb.8:
	movq	%r15, %rdi
	callq	_ZdaPv
.LBB4_9:
	xorl	%eax, %eax
	addq	$56, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end4:
	.size	main, .Lfunc_end4-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	.LBB5_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB5_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z10iterMandeliiPi, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end5:
	.size	__hip_module_ctor, .Lfunc_end5-__hip_module_ctor
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_dtor
	.type	__hip_module_dtor,@function
__hip_module_dtor:                      # @__hip_module_dtor
	.cfi_startproc
# %bb.0:
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB6_2
# %bb.1:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__hipUnregisterFatBinary
	movq	$0, __hip_gpubin_handle(%rip)
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB6_2:
	retq
.Lfunc_end6:
	.size	__hip_module_dtor, .Lfunc_end6-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z10iterMandeliiPi,@object      # @_Z10iterMandeliiPi
	.section	.rodata,"a",@progbits
	.globl	_Z10iterMandeliiPi
	.p2align	3, 0x0
_Z10iterMandeliiPi:
	.quad	_Z25__device_stub__iterMandeliiPi
	.size	_Z10iterMandeliiPi, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"Mandel.dat"
	.size	.L.str, 11

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"w"
	.size	.L.str.1, 2

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"%f\t%f\t%d\n"
	.size	.L.str.2, 10

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"Temps de calcul du Bouddha : %f secondes\n"
	.size	.L.str.4, 42

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z10iterMandeliiPi"
	.size	.L__unnamed_1, 19

	.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 _Z25__device_stub__iterMandeliiPi
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z10iterMandeliiPi
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z10iterMandeliiPi      ; -- Begin function _Z10iterMandeliiPi
	.globl	_Z10iterMandeliiPi
	.p2align	8
	.type	_Z10iterMandeliiPi,@function
_Z10iterMandeliiPi:                     ; @_Z10iterMandeliiPi
; %bb.0:
	s_clause 0x1
	s_load_b32 s4, s[0:1], 0x1c
	s_load_b64 s[2:3], s[0:1], 0x0
	v_and_b32_e32 v2, 0x3ff, v0
	v_bfe_u32 v3, v0, 10, 10
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s5, s4, 16
	s_and_b32 s4, s4, 0xffff
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_mad_u64_u32 v[0:1], null, s14, s4, v[2:3]
	v_mad_u64_u32 v[1:2], null, s15, s5, v[3:4]
	s_mov_b32 s4, exec_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_max_i32_e32 v2, v0, v1
	v_cmpx_gt_i32_e64 s3, v2
	s_cbranch_execz .LBB0_6
; %bb.1:
	v_cvt_f32_i32_e32 v2, s3
	s_cmp_gt_i32 s2, 0
	s_mov_b32 s5, 0
	s_cselect_b32 s4, -1, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v2, -1.0, v2
	v_div_scale_f32 v3, null, v2, v2, 0x40200000
	v_div_scale_f32 v6, vcc_lo, 0x40200000, v2, 0x40200000
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f32_e32 v4, v3
	s_waitcnt_depctr 0xfff
	v_fma_f32 v5, -v3, v4, 1.0
	v_fmac_f32_e32 v4, v5, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v5, v6, v4
	v_fma_f32 v7, -v3, v5, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v5, v7, v4
	v_fma_f32 v3, -v3, v5, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_div_fmas_f32 v3, v3, v4, v5
	v_cvt_f32_i32_e32 v4, v0
	v_div_fixup_f32 v3, v3, v2, 0x40200000
	v_cvt_f32_i32_e32 v2, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_fmaak_f32 v2, v3, v2, 0xbfa00000
	v_fma_f32 v3, v3, v4, -2.0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v4, v2, v2
	v_fma_f32 v5, v3, v3, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(SALU_CYCLE_1)
	v_cmp_gt_f32_e32 vcc_lo, 4.0, v5
	v_mov_b32_e32 v5, 0
	s_and_b32 s6, vcc_lo, s4
	s_and_saveexec_b32 s4, s6
	s_cbranch_execz .LBB0_5
; %bb.2:                                ; %.lr.ph.preheader
	v_dual_mul_f32 v7, v3, v3 :: v_dual_mov_b32 v6, v2
	v_mov_b32_e32 v8, v3
	s_mov_b32 s6, 0
.LBB0_3:                                ; %.lr.ph
                                        ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(VALU_DEP_1)
	v_add_f32_e32 v5, v8, v8
	s_add_i32 s6, s6, 1
	v_sub_f32_e32 v4, v7, v4
	s_cmp_ge_i32 s6, s2
	s_cselect_b32 s7, -1, 0
	v_fma_f32 v6, v5, v6, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v8, v3, v4
	v_dual_mul_f32 v4, v6, v6 :: v_dual_mul_f32 v7, v8, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v5, v8, v8, v4
	v_cmp_ngt_f32_e32 vcc_lo, 4.0, v5
	v_mov_b32_e32 v5, s6
	s_or_b32 s7, vcc_lo, s7
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s7, exec_lo, s7
	s_or_b32 s5, s7, s5
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s5
	s_cbranch_execnz .LBB0_3
; %bb.4:                                ; %Flow
	s_or_b32 exec_lo, exec_lo, s5
.LBB0_5:                                ; %Flow43
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	s_or_b32 exec_lo, exec_lo, s4
	s_load_b64 s[0:1], s[0:1], 0x8
	v_mad_u64_u32 v[2:3], null, v1, s3, v[0:1]
	v_ashrrev_i32_e32 v3, 31, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 2, v[2:3]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v0, vcc_lo, s0, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
	global_store_b32 v[0:1], v5, off
.LBB0_6:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z10iterMandeliiPi
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 272
		.amdhsa_user_sgpr_count 14
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 1
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 1
		.amdhsa_next_free_vgpr 9
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end0:
	.size	_Z10iterMandeliiPi, .Lfunc_end0-_Z10iterMandeliiPi
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 488
; NumSgprs: 18
; NumVgprs: 9
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 9
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 14
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 1
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .offset:         0
        .size:           4
        .value_kind:     by_value
      - .offset:         4
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .offset:         16
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         20
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         24
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         28
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         30
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         32
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         34
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         36
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         38
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         56
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         80
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 272
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z10iterMandeliiPi
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z10iterMandeliiPi.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     9
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
