/*

Calculating which points in the specified window of the complex plane
lie inside the Mandelbrot set

Can be easily extended using OpenGL to include graphics
and actually plot the results on screen

*/

#include <iostream>
#include <complex>
#include <cmath>
#include <cstdlib>

// -----------------------------
// declare some global variables

const int MAX_ITER = 200;    // maximum number of iterations

const int NX = 1000;  // number of points along X inside the window
const int NY = 1000;  // number of points along Y inside the window

// -------------------------------------------------------
// define a class for complex numbers and their operations

class dcmplx
{
public:
    double re;   // real component
    double im;   // imaginary component

// function to calculate the magnitude or absolute value of the complex number
// this function is called from and executes on the device (GPU) 

__device__
double magnitude()
{
    return pow((re*re + im*im),0.5);
}

};

// ----------------------------------------------------------------------------------
// function to check all points inside the specified window for membership in the set
// this function is called from the host (CPU) but executes on the device (GPU)

__global__
void Mandelbrot(double xmin, double xmax, double ymin, double ymax, int *dev_color)
{
    double dx = (xmax - xmin)/NX; // grid spacing along X
    double dy = (ymax - ymin)/NY; // grid spacing along Y

    int i = 10*blockIdx.x + threadIdx.x;
    int j = 10*blockIdx.y + threadIdx.y;

    double x = xmin + (double) i*dx;   // actual x coordinate (real component)
    double y = ymin + (double) j*dy;   // actual y coordinate (imaginary component)

    dcmplx c;
    c.re = x;
    c.im = y;

    dcmplx z;
    z.re = 0.0;
    z.im = 0.0;

    // ---------------
    // z <---- z*z + c

    int iter = 0;

    while(iter<MAX_ITER) {
        iter++;
        dcmplx temp = z;
        z.re = temp.re*temp.re - temp.im*temp.im  +  c.re;
        z.im = 2.0*temp.re*temp.im                +  c.im;
        
        if (z.magnitude() > 2.0) break;
    }

    // the 2D array "dev_color" stores how many iterations were required for divergence
    // for points outside the Mandelbrot set, this is typically a small number
    // points inside the set do not diverge and thus iter is a large number for such points

    dev_color[i+j*NX] = iter;

    // GRAPHICS: the value of iter can be used to pick a suitable "color" when drawing the set
}

// ---------------------------------------
// main program executes on the host (CPU)

int main(int argc, char* argv[])
{   
    // allocate memory for a 2D array on the host (CPU)
    int *color = new int[NX*NY];

    // fill the 2D array with zeroes
    for(int i=0;i<NX;i++) {
        for(int j=0;j<NY;j++) {
            color[i+j*NX] = 0;
        }
    }

    // allocate memory on the device (GPU)
    int *dev_color;
    cudaMalloc((void **) &dev_color, NX*NY*sizeof(int));

    // initialize parameters on the host
    double xmin = -2, xmax = 1, ymin = -1.5, ymax = 1.5;

    // check each pixel for membership and fill the 2D array accordingly
    // this kernel is executed in parallel on the GPU using a grid of 100 x 100 blocks
    // where each block contains 10 x 10 threads

    dim3 dimGrid (1000, 1000, 1);   // 1000 x 1000 blocks in a grid
    dim3 dimBlock(1   , 1   , 1);   // 1 x 1 x 1 threads in each block

    for(int k=0;k<10;k++) {
        Mandelbrot<<<dimGrid,dimBlock>>>(xmin,xmax,ymin,ymax,dev_color);
    }

    // copy the 2D array from device (GPU) to host (CPU)
    cudaMemcpy(color, dev_color, NX*NY*sizeof(int), cudaMemcpyDeviceToHost);
    
    return 0;
}

	.file	"tmpxft_003af1c5_00000000-6_gpu_MandelBrotSet.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB3940:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE3940:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z34__device_stub__Z10MandelbrotddddPiddddPi
	.type	_Z34__device_stub__Z10MandelbrotddddPiddddPi, @function
_Z34__device_stub__Z10MandelbrotddddPiddddPi:
.LFB3962:
	.cfi_startproc
	endbr64
	subq	$168, %rsp
	.cfi_def_cfa_offset 176
	movq	%rdi, 8(%rsp)
	leaq	56(%rsp), %rcx
	leaq	48(%rsp), %rdx
	movsd	%xmm0, 40(%rsp)
	leaq	76(%rsp), %rsi
	leaq	64(%rsp), %rdi
	movsd	%xmm1, 32(%rsp)
	movsd	%xmm2, 24(%rsp)
	movsd	%xmm3, 16(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 152(%rsp)
	xorl	%eax, %eax
	leaq	40(%rsp), %rax
	movl	$1, 72(%rsp)
	movq	%rax, 112(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 144(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 64(%rsp)
	movq	%rax, 76(%rsp)
	movl	$1, 84(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	56(%rsp)
	.cfi_def_cfa_offset 184
	leaq	_Z10MandelbrotddddPi(%rip), %rdi
	pushq	56(%rsp)
	.cfi_def_cfa_offset 192
	movq	92(%rsp), %rcx
	movl	100(%rsp), %r8d
	movq	80(%rsp), %rsi
	movl	88(%rsp), %edx
	leaq	128(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 184
	popq	%rdx
	.cfi_def_cfa_offset 176
.L2:
	movq	152(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$168, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3962:
	.size	_Z34__device_stub__Z10MandelbrotddddPiddddPi, .-_Z34__device_stub__Z10MandelbrotddddPiddddPi
	.globl	_Z10MandelbrotddddPi
	.type	_Z10MandelbrotddddPi, @function
_Z10MandelbrotddddPi:
.LFB3963:
	.cfi_startproc
	endbr64
	jmp	_Z34__device_stub__Z10MandelbrotddddPiddddPi
	.cfi_endproc
.LFE3963:
	.size	_Z10MandelbrotddddPi, .-_Z10MandelbrotddddPi
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB3937:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movl	$4000000, %edi
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	subq	$56, %rsp
	.cfi_def_cfa_offset 80
	movq	%fs:40, %rax
	movq	%rax, 40(%rsp)
	xorl	%eax, %eax
	call	_Znam@PLT
	xorl	%esi, %esi
	movq	%rax, %rbx
	movq	%rax, %rcx
.L9:
	movq	%rcx, %rdx
	movl	$1000, %eax
.L10:
	xorl	%edi, %edi
	addq	$4000, %rdx
	movl	%edi, -4000(%rdx)
	decl	%eax
	jne	.L10
	incl	%esi
	addq	$4, %rcx
	cmpl	$1000, %esi
	jne	.L9
	leaq	8(%rsp), %rdi
	movl	$4000000, %esi
	movl	$10, %ebp
	call	cudaMalloc@PLT
	movl	$1, 24(%rsp)
	movabsq	$4294967297000, %rax
	movq	%rax, 16(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 28(%rsp)
	movl	$1, 36(%rsp)
.L13:
	movl	36(%rsp), %ecx
	movq	28(%rsp), %rdx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	16(%rsp), %rdi
	movl	24(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L12
	movq	8(%rsp), %rdi
	movsd	.LC0(%rip), %xmm3
	movsd	.LC1(%rip), %xmm2
	movsd	.LC2(%rip), %xmm1
	movsd	.LC3(%rip), %xmm0
	call	_Z34__device_stub__Z10MandelbrotddddPiddddPi
.L12:
	decl	%ebp
	jne	.L13
	movq	8(%rsp), %rsi
	movl	$2, %ecx
	movl	$4000000, %edx
	movq	%rbx, %rdi
	call	cudaMemcpy@PLT
	movq	40(%rsp), %rax
	subq	%fs:40, %rax
	je	.L14
	call	__stack_chk_fail@PLT
.L14:
	addq	$56, %rsp
	.cfi_def_cfa_offset 24
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3937:
	.size	main, .-main
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC4:
	.string	"_Z10MandelbrotddddPi"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB3965:
	.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	.LC4(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z10MandelbrotddddPi(%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
.LFE3965:
	.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	0
	.long	1073217536
	.align 8
.LC1:
	.long	0
	.long	-1074266112
	.align 8
.LC2:
	.long	0
	.long	1072693248
	.align 8
.LC3:
	.long	0
	.long	-1073741824
	.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:

/*

Calculating which points in the specified window of the complex plane
lie inside the Mandelbrot set

Can be easily extended using OpenGL to include graphics
and actually plot the results on screen

*/


#include <hip/hip_runtime.h>
#include <iostream>
#include <complex>
#include <cmath>
#include <cstdlib>

// -----------------------------
// declare some global variables

const int MAX_ITER = 200;    // maximum number of iterations

const int NX = 1000;  // number of points along X inside the window
const int NY = 1000;  // number of points along Y inside the window

// -------------------------------------------------------
// define a class for complex numbers and their operations

class dcmplx
{
public:
    double re;   // real component
    double im;   // imaginary component

// function to calculate the magnitude or absolute value of the complex number
// this function is called from and executes on the device (GPU) 

__device__
double magnitude()
{
    return pow((re*re + im*im),0.5);
}

};

// ----------------------------------------------------------------------------------
// function to check all points inside the specified window for membership in the set
// this function is called from the host (CPU) but executes on the device (GPU)

__global__
void Mandelbrot(double xmin, double xmax, double ymin, double ymax, int *dev_color)
{
    double dx = (xmax - xmin)/NX; // grid spacing along X
    double dy = (ymax - ymin)/NY; // grid spacing along Y

    int i = 10*blockIdx.x + threadIdx.x;
    int j = 10*blockIdx.y + threadIdx.y;

    double x = xmin + (double) i*dx;   // actual x coordinate (real component)
    double y = ymin + (double) j*dy;   // actual y coordinate (imaginary component)

    dcmplx c;
    c.re = x;
    c.im = y;

    dcmplx z;
    z.re = 0.0;
    z.im = 0.0;

    // ---------------
    // z <---- z*z + c

    int iter = 0;

    while(iter<MAX_ITER) {
        iter++;
        dcmplx temp = z;
        z.re = temp.re*temp.re - temp.im*temp.im  +  c.re;
        z.im = 2.0*temp.re*temp.im                +  c.im;
        
        if (z.magnitude() > 2.0) break;
    }

    // the 2D array "dev_color" stores how many iterations were required for divergence
    // for points outside the Mandelbrot set, this is typically a small number
    // points inside the set do not diverge and thus iter is a large number for such points

    dev_color[i+j*NX] = iter;

    // GRAPHICS: the value of iter can be used to pick a suitable "color" when drawing the set
}

// ---------------------------------------
// main program executes on the host (CPU)

int main(int argc, char* argv[])
{   
    // allocate memory for a 2D array on the host (CPU)
    int *color = new int[NX*NY];

    // fill the 2D array with zeroes
    for(int i=0;i<NX;i++) {
        for(int j=0;j<NY;j++) {
            color[i+j*NX] = 0;
        }
    }

    // allocate memory on the device (GPU)
    int *dev_color;
    hipMalloc((void **) &dev_color, NX*NY*sizeof(int));

    // initialize parameters on the host
    double xmin = -2, xmax = 1, ymin = -1.5, ymax = 1.5;

    // check each pixel for membership and fill the 2D array accordingly
    // this kernel is executed in parallel on the GPU using a grid of 100 x 100 blocks
    // where each block contains 10 x 10 threads

    dim3 dimGrid (1000, 1000, 1);   // 1000 x 1000 blocks in a grid
    dim3 dimBlock(1   , 1   , 1);   // 1 x 1 x 1 threads in each block

    for(int k=0;k<10;k++) {
        Mandelbrot<<<dimGrid,dimBlock>>>(xmin,xmax,ymin,ymax,dev_color);
    }

    // copy the 2D array from device (GPU) to host (CPU)
    hipMemcpy(color, dev_color, NX*NY*sizeof(int), hipMemcpyDeviceToHost);
    
    return 0;
}

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

                                        # End of file scope inline assembly
	.globl	_Z25__device_stub__MandelbrotddddPi # -- Begin function _Z25__device_stub__MandelbrotddddPi
	.type	_Z25__device_stub__MandelbrotddddPi,@function
_Z25__device_stub__MandelbrotddddPi:    # @_Z25__device_stub__MandelbrotddddPi
	.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	$144, %rsp
	.cfi_def_cfa_offset 192
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	56(%rsp), %rax
	movsd	%xmm0, (%rax)
	leaq	48(%rsp), %rcx
	movsd	%xmm1, (%rcx)
	leaq	40(%rsp), %rdx
	movsd	%xmm2, (%rdx)
	leaq	32(%rsp), %rsi
	movsd	%xmm3, (%rsi)
	leaq	24(%rsp), %r8
	movq	%rdi, (%r8)
	leaq	96(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	movq	%rdx, 16(%rbx)
	movq	%rsi, 24(%rbx)
	movq	%r8, 32(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%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	$_Z10MandelbrotddddPi, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$160, %rsp
	.cfi_adjust_cfa_offset -160
	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__MandelbrotddddPi, .Lfunc_end0-_Z25__device_stub__MandelbrotddddPi
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function main
.LCPI1_0:
	.quad	0xc000000000000000              # double -2
.LCPI1_1:
	.quad	0x3ff0000000000000              # double 1
.LCPI1_2:
	.quad	0xbff8000000000000              # double -1.5
.LCPI1_3:
	.quad	0x3ff8000000000000              # double 1.5
	.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	%rbx
	.cfi_def_cfa_offset 40
	pushq	%rax
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	$4000000, %edi                  # imm = 0x3D0900
	callq	_Znam
	movq	%rax, %rbx
	xorl	%eax, %eax
	movq	%rbx, %rcx
.LBB1_1:                                # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB1_2 Depth 2
	movq	$-4000000, %rdx                 # imm = 0xFFC2F700
.LBB1_2:                                #   Parent Loop BB1_1 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movl	$0, 4000000(%rcx,%rdx)
	addq	$4000, %rdx                     # imm = 0xFA0
	jne	.LBB1_2
# %bb.3:                                #   in Loop: Header=BB1_1 Depth=1
	incq	%rax
	addq	$4, %rcx
	cmpq	$1000, %rax                     # imm = 0x3E8
	jne	.LBB1_1
# %bb.4:
	movq	%rsp, %rdi
	movl	$4000000, %esi                  # imm = 0x3D0900
	callq	hipMalloc
	movl	$10, %ebp
	movabsq	$4294967297000, %r14            # imm = 0x3E8000003E8
	movabsq	$4294967297, %r15               # imm = 0x100000001
.LBB1_5:                                # =>This Inner Loop Header: Depth=1
	movq	%r14, %rdi
	movl	$1, %esi
	movq	%r15, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB1_7
# %bb.6:                                #   in Loop: Header=BB1_5 Depth=1
	movq	(%rsp), %rdi
	movsd	.LCPI1_0(%rip), %xmm0           # xmm0 = mem[0],zero
	movsd	.LCPI1_1(%rip), %xmm1           # xmm1 = mem[0],zero
	movsd	.LCPI1_2(%rip), %xmm2           # xmm2 = mem[0],zero
	movsd	.LCPI1_3(%rip), %xmm3           # xmm3 = mem[0],zero
	callq	_Z25__device_stub__MandelbrotddddPi
.LBB1_7:                                #   in Loop: Header=BB1_5 Depth=1
	decl	%ebp
	jne	.LBB1_5
# %bb.8:
	movq	(%rsp), %rsi
	movl	$4000000, %edx                  # imm = 0x3D0900
	movq	%rbx, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	xorl	%eax, %eax
	addq	$8, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	main, .Lfunc_end1-main
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	jne	.LBB2_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB2_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z10MandelbrotddddPi, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end2:
	.size	__hip_module_ctor, .Lfunc_end2-__hip_module_ctor
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_dtor
	.type	__hip_module_dtor,@function
__hip_module_dtor:                      # @__hip_module_dtor
	.cfi_startproc
# %bb.0:
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB3_2
# %bb.1:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__hipUnregisterFatBinary
	movq	$0, __hip_gpubin_handle(%rip)
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB3_2:
	retq
.Lfunc_end3:
	.size	__hip_module_dtor, .Lfunc_end3-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z10MandelbrotddddPi,@object    # @_Z10MandelbrotddddPi
	.section	.rodata,"a",@progbits
	.globl	_Z10MandelbrotddddPi
	.p2align	3, 0x0
_Z10MandelbrotddddPi:
	.quad	_Z25__device_stub__MandelbrotddddPi
	.size	_Z10MandelbrotddddPi, 8

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

	.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__MandelbrotddddPi
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z10MandelbrotddddPi
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z10MandelbrotddddPi    ; -- Begin function _Z10MandelbrotddddPi
	.globl	_Z10MandelbrotddddPi
	.p2align	8
	.type	_Z10MandelbrotddddPi,@function
_Z10MandelbrotddddPi:                   ; @_Z10MandelbrotddddPi
; %bb.0:
	s_load_b256 s[4:11], s[0:1], 0x0
	s_mov_b32 s12, 0x55555555
	s_mov_b32 s16, 0x4222de17
	s_mov_b32 s18, 0x3abe935a
	s_mov_b32 s20, 0x47e6c9c2
	s_mov_b32 s22, 0xcfa74449
	s_mov_b32 s24, 0x71bf3c30
	s_mov_b32 s26, 0x1c7792ce
	s_mov_b32 s28, 0xfefa39ef
	s_mov_b32 s30, 0x3b39803f
	s_mov_b32 s34, 0xd5df274d
	s_mov_b32 s36, 0x652b82fe
	s_mov_b32 s38, 0xfca7ab0c
	s_mov_b32 s40, 0x6a5dcb37
	s_mov_b32 s42, 0x623fde64
	s_mov_b32 s44, 0x7c89e6b0
	s_mov_b32 s46, 0x14761f6e
	s_mov_b32 s48, 0x1852b7b0
	s_mov_b32 s50, 0x11122322
	s_mov_b32 s52, 0x555502a1
	s_waitcnt lgkmcnt(0)
	v_add_f64 v[3:4], s[6:7], -s[4:5]
	v_add_f64 v[5:6], s[10:11], -s[8:9]
	s_load_b64 s[10:11], s[0:1], 0x20
	s_mov_b32 s6, 11
	s_mov_b32 s54, 0x55555511
	s_mov_b32 s33, 0
	s_mov_b32 s7, 0x3fe00000
	s_mov_b32 s13, 0x3fe55555
	s_mov_b32 s17, 0x3fbdee67
	s_mov_b32 s19, 0x3fbe25e4
	s_mov_b32 s21, 0x3fc110ef
	s_mov_b32 s23, 0x3fc3b13b
	s_mov_b32 s25, 0x3fc745d1
	s_mov_b32 s27, 0x3fcc71c7
	s_mov_b32 s29, 0x3fe62e42
	s_mov_b32 s31, 0x3c7abc9e
	s_mov_b32 s3, 0xbfe55555
	s_mov_b32 s35, 0x3c8543b0
	s_mov_b32 s37, 0x3ff71547
	s_mov_b32 s39, 0x3e928af3
	s_mov_b32 s41, 0x3e5ade15
	s_mov_b32 s43, 0x3ec71dee
	s_mov_b32 s45, 0x3efa0199
	s_mov_b32 s47, 0x3f2a01a0
	s_mov_b32 s49, 0x3f56c16c
	s_mov_b32 s51, 0x3f811111
	s_mov_b32 s53, 0x3fa55555
	s_mov_b32 s55, 0x3fc55555
	s_mov_b32 s57, 0xbfe62e42
	s_mov_b32 s59, 0xbc7abc9e
	s_mov_b32 s60, 0
                                        ; implicit-def: $sgpr61
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_div_scale_f64 v[1:2], null, 0x408f4000, 0x408f4000, v[3:4]
	v_div_scale_f64 v[7:8], null, 0x408f4000, 0x408f4000, v[5:6]
	v_div_scale_f64 v[17:18], vcc_lo, v[3:4], 0x408f4000, v[3:4]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_rcp_f64_e32 v[9:10], v[1:2]
	v_rcp_f64_e32 v[11:12], v[7:8]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[13:14], -v[1:2], v[9:10], 1.0
	v_fma_f64 v[15:16], -v[7:8], v[11:12], 1.0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[9:10], v[9:10], v[13:14], v[9:10]
	v_fma_f64 v[11:12], v[11:12], v[15:16], v[11:12]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[13:14], -v[1:2], v[9:10], 1.0
	v_fma_f64 v[15:16], -v[7:8], v[11:12], 1.0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_fma_f64 v[9:10], v[9:10], v[13:14], v[9:10]
	v_div_scale_f64 v[13:14], s2, v[5:6], 0x408f4000, v[5:6]
	v_fma_f64 v[11:12], v[11:12], v[15:16], v[11:12]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_f64 v[15:16], v[17:18], v[9:10]
	v_mul_f64 v[19:20], v[13:14], v[11:12]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[1:2], -v[1:2], v[15:16], v[17:18]
	v_fma_f64 v[7:8], -v[7:8], v[19:20], v[13:14]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_div_fmas_f64 v[9:10], v[1:2], v[9:10], v[15:16]
	s_mov_b32 vcc_lo, s2
	v_and_b32_e32 v2, 0x3ff, v0
	v_div_fmas_f64 v[7:8], v[7:8], v[11:12], v[19:20]
	v_bfe_u32 v11, v0, 10, 10
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mad_u64_u32 v[0:1], null, s14, 10, v[2:3]
	v_mad_u64_u32 v[1:2], null, s15, 10, v[11:12]
	s_mov_b32 s14, 0x968915a9
	s_mov_b32 s15, 0x3fba6564
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cvt_f64_i32_e32 v[11:12], v0
	v_cvt_f64_i32_e32 v[13:14], v1
	v_div_fixup_f64 v[9:10], v[9:10], 0x408f4000, v[3:4]
	v_mov_b32_e32 v2, 0
	v_mov_b32_e32 v3, 0
	v_div_fixup_f64 v[6:7], v[7:8], 0x408f4000, v[5:6]
	v_mov_b32_e32 v8, 0
	v_fma_f64 v[4:5], v[9:10], v[11:12], s[4:5]
	s_delay_alu instid0(VALU_DEP_4)
	v_dual_mov_b32 v11, 0x3ff00000 :: v_dual_mov_b32 v10, v3
	v_mov_b32_e32 v9, v2
	v_fma_f64 v[6:7], v[6:7], v[13:14], s[8:9]
	s_mov_b32 s4, 0x924920da
	s_mov_b32 s8, 0x9999999c
	s_mov_b32 s5, 0x3fd24924
	s_mov_b32 s9, 0x3fd99999
.LBB0_1:                                ; =>This Inner Loop Header: Depth=1
	s_or_b32 s61, s61, exec_lo
	s_cmpk_eq_i32 s60, 0xc8
	s_movk_i32 s0, 0xc8
	s_cbranch_scc1 .LBB0_3
; %bb.2:                                ;   in Loop: Header=BB0_1 Depth=1
	v_mul_f64 v[12:13], v[2:3], v[2:3]
	v_add_f64 v[14:15], v[9:10], v[9:10]
	s_mov_b32 s2, s12
	s_mov_b32 s56, s28
	s_mov_b32 s58, s30
	s_add_i32 s60, s60, 1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[9:10], v[9:10], v[9:10], -v[12:13]
	v_fma_f64 v[2:3], v[2:3], v[14:15], v[6:7]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[12:13], v[4:5], v[9:10]
	v_mul_f64 v[9:10], v[2:3], v[2:3]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[14:15], v[12:13], v[12:13], v[9:10]
	v_frexp_mant_f64_e64 v[9:10], |v[14:15]|
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cmp_gt_f64_e32 vcc_lo, s[12:13], v[9:10]
	v_cndmask_b32_e64 v16, 0, 1, vcc_lo
	v_ldexp_f64 v[9:10], v[9:10], v16
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_f64 v[16:17], v[9:10], 1.0
	v_add_f64 v[22:23], v[9:10], -1.0
	v_rcp_f64_e32 v[18:19], v[16:17]
	v_add_f64 v[24:25], v[16:17], -1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_add_f64 v[9:10], v[9:10], -v[24:25]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[20:21], -v[16:17], v[18:19], 1.0
	v_fma_f64 v[18:19], v[20:21], v[18:19], v[18:19]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[20:21], -v[16:17], v[18:19], 1.0
	v_fma_f64 v[18:19], v[20:21], v[18:19], v[18:19]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[20:21], v[22:23], v[18:19]
	v_mul_f64 v[26:27], v[16:17], v[20:21]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[16:17], v[20:21], v[16:17], -v[26:27]
	v_fma_f64 v[9:10], v[20:21], v[9:10], v[16:17]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[16:17], v[26:27], v[9:10]
	v_add_f64 v[24:25], v[22:23], -v[16:17]
	v_add_f64 v[26:27], v[16:17], -v[26:27]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[22:23], v[22:23], -v[24:25]
	v_add_f64 v[9:10], v[26:27], -v[9:10]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[16:17], v[22:23], -v[16:17]
	v_add_f64 v[9:10], v[9:10], v[16:17]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[9:10], v[24:25], v[9:10]
	v_mul_f64 v[9:10], v[18:19], v[9:10]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[16:17], v[20:21], v[9:10]
	v_add_f64 v[18:19], v[16:17], -v[20:21]
	v_mul_f64 v[20:21], v[16:17], v[16:17]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[9:10], v[9:10], -v[18:19]
	v_fma_f64 v[18:19], v[16:17], v[16:17], -v[20:21]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[22:23], v[9:10], v[9:10]
	v_fma_f64 v[18:19], v[16:17], v[22:23], v[18:19]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[22:23], v[20:21], v[18:19]
	v_fma_f64 v[24:25], v[22:23], s[16:17], s[14:15]
	v_add_f64 v[20:21], v[22:23], -v[20:21]
	v_mul_f64 v[30:31], v[16:17], v[22:23]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[24:25], v[22:23], v[24:25], s[18:19]
	v_add_f64 v[18:19], v[18:19], -v[20:21]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[24:25], v[22:23], v[24:25], s[20:21]
	v_fma_f64 v[24:25], v[22:23], v[24:25], s[22:23]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[24:25], v[22:23], v[24:25], s[24:25]
	v_fma_f64 v[24:25], v[22:23], v[24:25], s[26:27]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[24:25], v[22:23], v[24:25], s[4:5]
	v_fma_f64 v[24:25], v[22:23], v[24:25], s[8:9]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[26:27], v[22:23], v[24:25]
	v_fma_f64 v[20:21], v[22:23], v[24:25], -v[26:27]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[20:21], v[18:19], v[24:25], v[20:21]
	v_add_f64 v[24:25], v[26:27], v[20:21]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_f64 v[28:29], v[24:25], s[12:13]
	v_add_f64 v[26:27], v[24:25], -v[26:27]
	v_add_f64 v[32:33], v[28:29], s[2:3]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_4)
	v_add_f64 v[20:21], v[20:21], -v[26:27]
	v_fma_f64 v[26:27], v[22:23], v[16:17], -v[30:31]
	v_cmp_eq_f64_e64 s2, 0, v[14:15]
	v_add_f64 v[24:25], v[24:25], -v[32:33]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_f64 v[20:21], v[20:21], s[34:35]
	v_fma_f64 v[22:23], v[22:23], v[9:10], v[26:27]
	v_ldexp_f64 v[9:10], v[9:10], 1
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f64 v[20:21], v[20:21], v[24:25]
	v_fma_f64 v[18:19], v[18:19], v[16:17], v[22:23]
	v_ldexp_f64 v[16:17], v[16:17], 1
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f64 v[22:23], v[28:29], v[20:21]
	v_add_f64 v[24:25], v[30:31], v[18:19]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[26:27], v[28:29], -v[22:23]
	v_mul_f64 v[28:29], v[24:25], v[22:23]
	v_add_f64 v[30:31], v[24:25], -v[30:31]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f64 v[20:21], v[20:21], v[26:27]
	v_fma_f64 v[26:27], v[24:25], v[22:23], -v[28:29]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[18:19], v[18:19], -v[30:31]
	v_fma_f64 v[20:21], v[24:25], v[20:21], v[26:27]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_fma_f64 v[18:19], v[18:19], v[22:23], v[20:21]
	v_frexp_exp_i32_f64_e32 v22, v[14:15]
	v_add_f64 v[20:21], v[28:29], v[18:19]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_subrev_co_ci_u32_e32 v22, vcc_lo, 0, v22, vcc_lo
	v_cmp_eq_f64_e32 vcc_lo, 1.0, v[14:15]
	v_cvt_f64_i32_e32 v[22:23], v22
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_f64 v[24:25], v[16:17], v[20:21]
	v_add_f64 v[26:27], v[20:21], -v[28:29]
	v_mul_f64 v[28:29], v[22:23], s[28:29]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f64 v[16:17], v[24:25], -v[16:17]
	v_add_f64 v[18:19], v[18:19], -v[26:27]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[26:27], v[22:23], s[28:29], -v[28:29]
	v_add_f64 v[16:17], v[20:21], -v[16:17]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f64 v[9:10], v[9:10], v[18:19]
	v_fma_f64 v[18:19], v[22:23], s[30:31], v[26:27]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[9:10], v[9:10], v[16:17]
	v_add_f64 v[16:17], v[28:29], v[18:19]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[20:21], v[24:25], v[9:10]
	v_add_f64 v[28:29], v[16:17], -v[28:29]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_f64 v[22:23], v[16:17], v[20:21]
	v_add_f64 v[24:25], v[20:21], -v[24:25]
	v_add_f64 v[18:19], v[18:19], -v[28:29]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f64 v[26:27], v[22:23], -v[16:17]
	v_add_f64 v[9:10], v[9:10], -v[24:25]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_f64 v[30:31], v[22:23], -v[26:27]
	v_add_f64 v[20:21], v[20:21], -v[26:27]
	v_add_f64 v[24:25], v[18:19], v[9:10]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[16:17], v[16:17], -v[30:31]
	v_add_f64 v[16:17], v[20:21], v[16:17]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[20:21], v[24:25], -v[18:19]
	v_add_f64 v[16:17], v[24:25], v[16:17]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_f64 v[24:25], v[24:25], -v[20:21]
	v_add_f64 v[9:10], v[9:10], -v[20:21]
	v_add_f64 v[26:27], v[22:23], v[16:17]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[18:19], v[18:19], -v[24:25]
	v_add_f64 v[20:21], v[26:27], -v[22:23]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[9:10], v[9:10], v[18:19]
	v_add_f64 v[16:17], v[16:17], -v[20:21]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_f64 v[16:17], v[9:10], v[16:17]
	v_cndmask_b32_e32 v9, 0x3fe00000, v11, vcc_lo
	v_add_f64 v[18:19], v[26:27], v[16:17]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f64 v[20:21], v[18:19], -v[26:27]
	v_mul_f64 v[22:23], v[8:9], v[18:19]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[16:17], v[16:17], -v[20:21]
	v_fma_f64 v[18:19], v[8:9], v[18:19], -v[22:23]
	v_cmp_class_f64_e64 vcc_lo, v[22:23], 0x204
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[16:17], v[8:9], v[16:17], v[18:19]
	v_add_f64 v[18:19], v[22:23], v[16:17]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_dual_cndmask_b32 v21, v19, v23 :: v_dual_cndmask_b32 v20, v18, v22
	v_add_f64 v[18:19], v[18:19], -v[22:23]
	v_mul_f64 v[24:25], v[20:21], s[36:37]
	v_cmp_nlt_f64_e64 s0, 0x40900000, v[20:21]
	v_cmp_neq_f64_e64 vcc_lo, 0x7ff00000, |v[20:21]|
	v_cmp_ngt_f64_e64 s1, 0xc090cc00, v[20:21]
	v_add_f64 v[16:17], v[16:17], -v[18:19]
	v_trunc_f64_e32 v[18:19], v[8:9]
	v_rndne_f64_e32 v[24:25], v[24:25]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_dual_cndmask_b32 v17, 0, v17 :: v_dual_cndmask_b32 v16, 0, v16
	s_and_b32 vcc_lo, s1, s0
	v_fma_f64 v[26:27], v[24:25], s[56:57], v[20:21]
	v_cvt_i32_f64_e32 v10, v[24:25]
	v_cmp_class_f64_e64 s56, v[14:15], 0x204
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[26:27], v[24:25], s[58:59], v[26:27]
	v_fma_f64 v[28:29], v[26:27], s[40:41], s[38:39]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[28:29], v[26:27], v[28:29], s[42:43]
	v_fma_f64 v[28:29], v[26:27], v[28:29], s[44:45]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[28:29], v[26:27], v[28:29], s[46:47]
	v_fma_f64 v[28:29], v[26:27], v[28:29], s[48:49]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[28:29], v[26:27], v[28:29], s[50:51]
	v_fma_f64 v[28:29], v[26:27], v[28:29], s[52:53]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[28:29], v[26:27], v[28:29], s[54:55]
	v_fma_f64 v[28:29], v[26:27], v[28:29], s[6:7]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[28:29], v[26:27], v[28:29], 1.0
	v_fma_f64 v[24:25], v[26:27], v[28:29], 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_ldexp_f64 v[22:23], v[24:25], v10
	v_mul_f64 v[24:25], v[8:9], 0.5
	v_cndmask_b32_e64 v10, 0x7ff00000, v23, s0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_trunc_f64_e32 v[20:21], v[24:25]
	v_cndmask_b32_e32 v22, 0, v22, vcc_lo
	v_cmp_eq_f64_e32 vcc_lo, v[18:19], v[8:9]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e64 v23, 0, v10, s1
	v_fma_f64 v[9:10], v[22:23], v[16:17], v[22:23]
	v_cmp_class_f64_e64 s1, v[22:23], 0x204
	v_cndmask_b32_e64 v16, 0x7ff00000, 0, s2
	v_cmp_neq_f64_e64 s0, v[20:21], v[24:25]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_cndmask_b32_e64 v10, v10, v23, s1
	v_cndmask_b32_e64 v9, v9, v22, s1
	s_and_b32 vcc_lo, vcc_lo, s0
	s_and_not1_b32 s0, s61, exec_lo
	v_cndmask_b32_e32 v17, 0x3ff00000, v15, vcc_lo
	v_cndmask_b32_e32 v18, 0, v15, vcc_lo
	s_or_b32 vcc_lo, s2, s56
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cndmask_b32_e64 v9, v9, 0, vcc_lo
	v_bfi_b32 v10, 0x7fffffff, v10, v17
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_bfi_b32 v16, 0x7fffffff, v16, v18
	v_cndmask_b32_e32 v10, v10, v16, vcc_lo
	v_cmp_o_f64_e32 vcc_lo, v[14:15], v[14:15]
	v_cndmask_b32_e32 v9, 0, v9, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v10, 0x7ff80000, v10, vcc_lo
	v_cmp_lt_f64_e32 vcc_lo, 2.0, v[9:10]
	v_dual_mov_b32 v9, v12 :: v_dual_mov_b32 v10, v13
	s_and_b32 s1, vcc_lo, exec_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 s61, s0, s1
	s_mov_b32 s0, s60
	s_branch .LBB0_4
.LBB0_3:                                ;   in Loop: Header=BB0_1 Depth=1
                                        ; implicit-def: $sgpr60
                                        ; implicit-def: $vgpr9_vgpr10
                                        ; implicit-def: $vgpr2_vgpr3
.LBB0_4:                                ; %Flow
                                        ;   in Loop: Header=BB0_1 Depth=1
	v_mov_b32_e32 v12, s0
	s_and_b32 s1, exec_lo, s61
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 s33, s1, s33
	s_and_not1_b32 exec_lo, exec_lo, s33
	s_cbranch_execnz .LBB0_1
; %bb.5:
	s_or_b32 exec_lo, exec_lo, s33
	v_mad_u64_u32 v[2:3], null, 0x3e8, v1, v[0:1]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v3, 31, v2
	v_lshlrev_b64 v[0:1], 2, v[2:3]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v0, vcc_lo, s10, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s11, v1, vcc_lo
	global_store_b32 v[0:1], v12, off
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z10MandelbrotddddPi
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 40
		.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 34
		.amdhsa_next_free_sgpr 62
		.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	_Z10MandelbrotddddPi, .Lfunc_end0-_Z10MandelbrotddddPi
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 2580
; NumSgprs: 64
; NumVgprs: 34
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 7
; VGPRBlocks: 4
; NumSGPRsForWavesPerEU: 64
; NumVGPRsForWavesPerEU: 34
; 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:           8
        .value_kind:     by_value
      - .offset:         8
        .size:           8
        .value_kind:     by_value
      - .offset:         16
        .size:           8
        .value_kind:     by_value
      - .offset:         24
        .size:           8
        .value_kind:     by_value
      - .address_space:  global
        .offset:         32
        .size:           8
        .value_kind:     global_buffer
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 40
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z10MandelbrotddddPi
    .private_segment_fixed_size: 0
    .sgpr_count:     64
    .sgpr_spill_count: 0
    .symbol:         _Z10MandelbrotddddPi.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     34
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
