/**
 * @file processQuatJulEle_linux.cu
 * 
 * Calculate julia sets for quaternions.
 *  
 */

/** Work out which piece of the global array this thread should operate on */ 
__device__ size_t calculateGlobalIndex() {
    // Which block are we?
    size_t const globalBlockIndex = blockIdx.x + blockIdx.y * gridDim.x;
    // Which thread are we within the block?
    size_t const localThreadIdx = threadIdx.x + blockDim.x * threadIdx.y;
    // How big is each block?
    size_t const threadsPerBlock = blockDim.x*blockDim.y;
    // Which thread are we overall?
    return localThreadIdx + globalBlockIndex*threadsPerBlock;

}

/** 
    Iterate for one point and return its count
 */ 
__device__ unsigned int doIterations( double const xPart0, 
                                      double const yPart0,
                                      double const zPart0,
                                      double const wPart0,
                                      double const cx,
                                      double const cy,
                                      double const cz,
                                      double const cw,
                                      unsigned int const maxIters ) {
    // Initialise: z = z0
    double xPart = xPart0;
    double yPart = yPart0;
    double zPart = zPart0;
    double wPart = wPart0;
    unsigned int count = 0;
    // Loop until escaped. Check Quaternion magnitude smaler than 4
    while ( ( count <= maxIters )
            && ((xPart*xPart + yPart*yPart + zPart*zPart + wPart*wPart) <= 16.0) ) {
        ++count;
        // Update: z = z*z + z0;
        double const oldXPart = xPart;
        double const oldYPart = yPart;
        double const oldZPart = zPart;
        double const oldWPart = wPart;
        // Quat mult and add constant
        xPart = oldXPart*oldXPart-oldYPart*oldYPart-oldZPart*oldZPart-oldWPart*oldWPart + cx;
        yPart = oldXPart*oldYPart+oldYPart*oldXPart-oldZPart*oldWPart+oldWPart*oldZPart + cy;
        zPart = oldXPart*oldZPart+oldYPart*oldWPart+oldZPart*oldXPart-oldWPart*oldYPart + cz;
        wPart = oldXPart*oldWPart-oldYPart*oldZPart+oldZPart*oldYPart+oldWPart*oldXPart + cw;
        //xPart = xPart*xPart - yPart*yPart + xPart0;
        //yPart = 2.0*oldRealPart*yPart + yPart0;
    }
    return count;
}


/** Main entry point.
 * Read data from arraybuffer and call doIterations for every point.
 */
__global__ void processMandelbrotElement( 
                      double * out, 
                      const double * x, 
                      const double * y,
                      const double * z,
                      const double * w,
                      const double cx,
                      const double cy,
                      const double cz,
                      const double cw,                      
                      const unsigned int maxIters, 
                      const unsigned int numel ) {
    // Work out which thread we are
    size_t const globalThreadIdx = calculateGlobalIndex();

    // If we're off the end, return now
    if (globalThreadIdx >= numel) {
        return;
    }
    
    // Get our X, Y, Z, W coords
    double const xPart0 = x[globalThreadIdx];
    double const yPart0 = y[globalThreadIdx];
    double const zPart0 = z[globalThreadIdx];
    double const wPart0 = w[globalThreadIdx];

    // Run the itearations on this location
    unsigned int const count = doIterations( xPart0, yPart0, zPart0, wPart0, cx, cy, cz, cw, maxIters );
    out[globalThreadIdx] =  double( count  );
}

	.file	"tmpxft_0030d9ff_00000000-6_processQuatJulEle_linux.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2012:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2012:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z20calculateGlobalIndexv
	.type	_Z20calculateGlobalIndexv, @function
_Z20calculateGlobalIndexv:
.LFB2008:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2008:
	.size	_Z20calculateGlobalIndexv, .-_Z20calculateGlobalIndexv
	.globl	_Z12doIterationsddddddddj
	.type	_Z12doIterationsddddddddj, @function
_Z12doIterationsddddddddj:
.LFB2009:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2009:
	.size	_Z12doIterationsddddddddj, .-_Z12doIterationsddddddddj
	.globl	_Z62__device_stub__Z24processMandelbrotElementPdPKdS1_S1_S1_ddddjjPdPKdS1_S1_S1_ddddjj
	.type	_Z62__device_stub__Z24processMandelbrotElementPdPKdS1_S1_S1_ddddjjPdPKdS1_S1_S1_ddddjj, @function
_Z62__device_stub__Z24processMandelbrotElementPdPKdS1_S1_S1_ddddjjPdPKdS1_S1_S1_ddddjj:
.LFB2034:
	.cfi_startproc
	endbr64
	subq	$248, %rsp
	.cfi_def_cfa_offset 256
	movq	%rdi, 72(%rsp)
	leaq	96(%rsp), %rdi
	movq	%rsi, 64(%rsp)
	leaq	108(%rsp), %rsi
	movq	%rdx, 56(%rsp)
	leaq	80(%rsp), %rdx
	movq	%rcx, 48(%rsp)
	leaq	88(%rsp), %rcx
	movq	%r8, 40(%rsp)
	movl	%r9d, 4(%rsp)
	movsd	%xmm0, 32(%rsp)
	movsd	%xmm1, 24(%rsp)
	movsd	%xmm2, 16(%rsp)
	movsd	%xmm3, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 232(%rsp)
	xorl	%eax, %eax
	leaq	72(%rsp), %rax
	movl	$1, 104(%rsp)
	movq	%rax, 144(%rsp)
	leaq	64(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	56(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 192(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 200(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 208(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 216(%rsp)
	leaq	256(%rsp), %rax
	movq	%rax, 224(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 96(%rsp)
	movq	%rax, 108(%rsp)
	movl	$1, 116(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L6
	pushq	88(%rsp)
	.cfi_def_cfa_offset 264
	leaq	_Z24processMandelbrotElementPdPKdS1_S1_S1_ddddjj(%rip), %rdi
	pushq	88(%rsp)
	.cfi_def_cfa_offset 272
	movq	124(%rsp), %rcx
	movl	132(%rsp), %r8d
	movq	112(%rsp), %rsi
	movl	120(%rsp), %edx
	leaq	160(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 264
	popq	%rdx
	.cfi_def_cfa_offset 256
.L6:
	movq	232(%rsp), %rax
	subq	%fs:40, %rax
	je	.L8
	call	__stack_chk_fail@PLT
.L8:
	addq	$248, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2034:
	.size	_Z62__device_stub__Z24processMandelbrotElementPdPKdS1_S1_S1_ddddjjPdPKdS1_S1_S1_ddddjj, .-_Z62__device_stub__Z24processMandelbrotElementPdPKdS1_S1_S1_ddddjjPdPKdS1_S1_S1_ddddjj
	.globl	_Z24processMandelbrotElementPdPKdS1_S1_S1_ddddjj
	.type	_Z24processMandelbrotElementPdPKdS1_S1_S1_ddddjj, @function
_Z24processMandelbrotElementPdPKdS1_S1_S1_ddddjj:
.LFB2035:
	.cfi_startproc
	endbr64
	jmp	_Z62__device_stub__Z24processMandelbrotElementPdPKdS1_S1_S1_ddddjjPdPKdS1_S1_S1_ddddjj
	.cfi_endproc
.LFE2035:
	.size	_Z24processMandelbrotElementPdPKdS1_S1_S1_ddddjj, .-_Z24processMandelbrotElementPdPKdS1_S1_S1_ddddjj
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"_Z24processMandelbrotElementPdPKdS1_S1_S1_ddddjj"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2037:
	.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	.LC0(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z24processMandelbrotElementPdPKdS1_S1_S1_ddddjj(%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
.LFE2037:
	.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
	.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>
/**
 * @file processQuatJulEle_linux.cu
 * 
 * Calculate julia sets for quaternions.
 *  
 */

/** Work out which piece of the global array this thread should operate on */ 
__device__ size_t calculateGlobalIndex() {
    // Which block are we?
    size_t const globalBlockIndex = blockIdx.x + blockIdx.y * gridDim.x;
    // Which thread are we within the block?
    size_t const localThreadIdx = threadIdx.x + blockDim.x * threadIdx.y;
    // How big is each block?
    size_t const threadsPerBlock = blockDim.x*blockDim.y;
    // Which thread are we overall?
    return localThreadIdx + globalBlockIndex*threadsPerBlock;

}

/** 
    Iterate for one point and return its count
 */ 
__device__ unsigned int doIterations( double const xPart0, 
                                      double const yPart0,
                                      double const zPart0,
                                      double const wPart0,
                                      double const cx,
                                      double const cy,
                                      double const cz,
                                      double const cw,
                                      unsigned int const maxIters ) {
    // Initialise: z = z0
    double xPart = xPart0;
    double yPart = yPart0;
    double zPart = zPart0;
    double wPart = wPart0;
    unsigned int count = 0;
    // Loop until escaped. Check Quaternion magnitude smaler than 4
    while ( ( count <= maxIters )
            && ((xPart*xPart + yPart*yPart + zPart*zPart + wPart*wPart) <= 16.0) ) {
        ++count;
        // Update: z = z*z + z0;
        double const oldXPart = xPart;
        double const oldYPart = yPart;
        double const oldZPart = zPart;
        double const oldWPart = wPart;
        // Quat mult and add constant
        xPart = oldXPart*oldXPart-oldYPart*oldYPart-oldZPart*oldZPart-oldWPart*oldWPart + cx;
        yPart = oldXPart*oldYPart+oldYPart*oldXPart-oldZPart*oldWPart+oldWPart*oldZPart + cy;
        zPart = oldXPart*oldZPart+oldYPart*oldWPart+oldZPart*oldXPart-oldWPart*oldYPart + cz;
        wPart = oldXPart*oldWPart-oldYPart*oldZPart+oldZPart*oldYPart+oldWPart*oldXPart + cw;
        //xPart = xPart*xPart - yPart*yPart + xPart0;
        //yPart = 2.0*oldRealPart*yPart + yPart0;
    }
    return count;
}


/** Main entry point.
 * Read data from arraybuffer and call doIterations for every point.
 */
__global__ void processMandelbrotElement( 
                      double * out, 
                      const double * x, 
                      const double * y,
                      const double * z,
                      const double * w,
                      const double cx,
                      const double cy,
                      const double cz,
                      const double cw,                      
                      const unsigned int maxIters, 
                      const unsigned int numel ) {
    // Work out which thread we are
    size_t const globalThreadIdx = calculateGlobalIndex();

    // If we're off the end, return now
    if (globalThreadIdx >= numel) {
        return;
    }
    
    // Get our X, Y, Z, W coords
    double const xPart0 = x[globalThreadIdx];
    double const yPart0 = y[globalThreadIdx];
    double const zPart0 = z[globalThreadIdx];
    double const wPart0 = w[globalThreadIdx];

    // Run the itearations on this location
    unsigned int const count = doIterations( xPart0, yPart0, zPart0, wPart0, cx, cy, cz, cw, maxIters );
    out[globalThreadIdx] =  double( count  );
}

	.text
	.file	"processQuatJulEle_linux.hip"
	.globl	_Z39__device_stub__processMandelbrotElementPdPKdS1_S1_S1_ddddjj # -- Begin function _Z39__device_stub__processMandelbrotElementPdPKdS1_S1_S1_ddddjj
	.type	_Z39__device_stub__processMandelbrotElementPdPKdS1_S1_S1_ddddjj,@function
_Z39__device_stub__processMandelbrotElementPdPKdS1_S1_S1_ddddjj: # @_Z39__device_stub__processMandelbrotElementPdPKdS1_S1_S1_ddddjj
	.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	$224, %rsp
	.cfi_def_cfa_offset 272
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	88(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	80(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	72(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	64(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	56(%rsp), %rcx
	movq	%r8, (%rcx)
	leaq	48(%rsp), %r8
	movsd	%xmm0, (%r8)
	leaq	40(%rsp), %r10
	movsd	%xmm1, (%r10)
	leaq	32(%rsp), %r11
	movsd	%xmm2, (%r11)
	leaq	24(%rsp), %r14
	movsd	%xmm3, (%r14)
	leaq	4(%rsp), %r15
	movl	%r9d, (%r15)
	leaq	128(%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)
	movq	%r10, 48(%rbx)
	movq	%r11, 56(%rbx)
	movq	%r14, 64(%rbx)
	movq	%r15, 72(%rbx)
	leaq	272(%rsp), %rax
	movq	%rax, 80(%rbx)
	leaq	112(%rsp), %r14
	leaq	96(%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	$_Z24processMandelbrotElementPdPKdS1_S1_S1_ddddjj, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$240, %rsp
	.cfi_adjust_cfa_offset -240
	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	_Z39__device_stub__processMandelbrotElementPdPKdS1_S1_S1_ddddjj, .Lfunc_end0-_Z39__device_stub__processMandelbrotElementPdPKdS1_S1_S1_ddddjj
	.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	.LBB1_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB1_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z24processMandelbrotElementPdPKdS1_S1_S1_ddddjj, %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_end1:
	.size	__hip_module_ctor, .Lfunc_end1-__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	.LBB2_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
.LBB2_2:
	retq
.Lfunc_end2:
	.size	__hip_module_dtor, .Lfunc_end2-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z24processMandelbrotElementPdPKdS1_S1_S1_ddddjj,@object # @_Z24processMandelbrotElementPdPKdS1_S1_S1_ddddjj
	.section	.rodata,"a",@progbits
	.globl	_Z24processMandelbrotElementPdPKdS1_S1_S1_ddddjj
	.p2align	3, 0x0
_Z24processMandelbrotElementPdPKdS1_S1_S1_ddddjj:
	.quad	_Z39__device_stub__processMandelbrotElementPdPKdS1_S1_S1_ddddjj
	.size	_Z24processMandelbrotElementPdPKdS1_S1_S1_ddddjj, 8

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

	.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 _Z39__device_stub__processMandelbrotElementPdPKdS1_S1_S1_ddddjj
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z24processMandelbrotElementPdPKdS1_S1_S1_ddddjj
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z24processMandelbrotElementPdPKdS1_S1_S1_ddddjj ; -- Begin function _Z24processMandelbrotElementPdPKdS1_S1_S1_ddddjj
	.globl	_Z24processMandelbrotElementPdPKdS1_S1_S1_ddddjj
	.p2align	8
	.type	_Z24processMandelbrotElementPdPKdS1_S1_S1_ddddjj,@function
_Z24processMandelbrotElementPdPKdS1_S1_S1_ddddjj: ; @_Z24processMandelbrotElementPdPKdS1_S1_S1_ddddjj
; %bb.0:
	s_clause 0x1
	s_load_b32 s4, s[0:1], 0x50
	s_load_b32 s5, s[0:1], 0x5c
	s_add_u32 s2, s0, 0x50
	s_addc_u32 s3, s1, 0
	v_mov_b32_e32 v3, 0
	s_mov_b32 s21, 0
	v_and_b32_e32 v2, 0x3ff, v0
	v_bfe_u32 v0, v0, 10, 10
	s_waitcnt lgkmcnt(0)
	s_mul_i32 s6, s4, s15
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(SALU_CYCLE_1)
	s_add_i32 s6, s6, s14
	s_cmp_lt_u32 s14, s4
	s_cselect_b32 s4, 12, 18
	v_mov_b32_e32 v1, s4
	s_lshr_b32 s4, s5, 16
	global_load_u16 v1, v1, s[2:3]
	s_load_b64 s[2:3], s[0:1], 0x48
	s_waitcnt lgkmcnt(0)
	s_mov_b32 s20, s3
	s_mov_b32 s3, exec_lo
	s_waitcnt vmcnt(0)
	v_mul_lo_u32 v4, s4, v1
	v_mad_u32_u24 v2, v0, v1, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[0:1], null, v4, s6, v[2:3]
	v_cmpx_gt_u64_e64 s[20:21], v[0:1]
	s_cbranch_execz .LBB0_6
; %bb.1:
	s_load_b512 s[4:19], s[0:1], 0x0
	v_lshlrev_b64 v[0:1], 3, v[0:1]
	s_load_b64 s[0:1], s[0:1], 0x40
	s_mov_b32 s3, 0
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v2, vcc_lo, s6, v0
	v_add_co_ci_u32_e32 v3, vcc_lo, s7, v1, vcc_lo
	v_add_co_u32 v6, vcc_lo, s8, v0
	v_add_co_ci_u32_e32 v7, vcc_lo, s9, v1, vcc_lo
	v_add_co_u32 v10, vcc_lo, s10, v0
	v_add_co_ci_u32_e32 v11, vcc_lo, s11, v1, vcc_lo
	v_add_co_u32 v12, vcc_lo, s12, v0
	v_add_co_ci_u32_e32 v13, vcc_lo, s13, v1, vcc_lo
	global_load_b64 v[4:5], v[2:3], off
	global_load_b64 v[8:9], v[6:7], off
	global_load_b64 v[6:7], v[10:11], off
	global_load_b64 v[2:3], v[12:13], off
                                        ; implicit-def: $sgpr6
.LBB0_2:                                ; =>This Inner Loop Header: Depth=1
	s_waitcnt vmcnt(2)
	v_mul_f64 v[10:11], v[8:9], v[8:9]
	s_or_b32 s6, s6, exec_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_fma_f64 v[12:13], v[4:5], v[4:5], v[10:11]
	s_waitcnt vmcnt(1)
	v_fma_f64 v[12:13], v[6:7], v[6:7], v[12:13]
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[12:13], v[2:3], v[2:3], v[12:13]
	v_cmp_ge_f64_e32 vcc_lo, 0x40300000, v[12:13]
	v_mov_b32_e32 v12, s3
	s_and_saveexec_b32 s7, vcc_lo
	s_cbranch_execz .LBB0_4
; %bb.3:                                ;   in Loop: Header=BB0_2 Depth=1
	v_mul_f64 v[12:13], v[8:9], v[2:3]
	v_mul_f64 v[16:17], v[4:5], v[4:5]
	v_mul_f64 v[14:15], v[4:5], v[8:9]
	v_mul_f64 v[18:19], v[8:9], v[6:7]
	v_mul_f64 v[20:21], v[6:7], v[6:7]
	s_add_i32 s3, s3, 1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_3) | instid1(SALU_CYCLE_1)
	s_cmp_gt_u32 s3, s2
	s_cselect_b32 s8, -1, 0
	s_and_not1_b32 s6, s6, exec_lo
	s_and_b32 s8, s8, exec_lo
	s_or_b32 s6, s6, s8
	v_fma_f64 v[12:13], v[4:5], v[6:7], v[12:13]
	v_add_f64 v[10:11], v[16:17], -v[10:11]
	v_add_f64 v[14:15], v[14:15], v[14:15]
	v_fma_f64 v[16:17], v[4:5], v[2:3], -v[18:19]
	v_mul_f64 v[18:19], v[2:3], v[2:3]
	v_fma_f64 v[12:13], v[4:5], v[6:7], v[12:13]
	v_add_f64 v[10:11], v[10:11], -v[20:21]
	v_fma_f64 v[14:15], -v[6:7], v[2:3], v[14:15]
	v_fma_f64 v[16:17], v[8:9], v[6:7], v[16:17]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_fma_f64 v[8:9], -v[8:9], v[2:3], v[12:13]
	v_add_f64 v[10:11], v[10:11], -v[18:19]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_fma_f64 v[6:7], v[6:7], v[2:3], v[14:15]
	v_fma_f64 v[2:3], v[4:5], v[2:3], v[16:17]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_f64 v[4:5], v[8:9], s[18:19]
	v_add_f64 v[10:11], v[10:11], s[14:15]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_f64 v[8:9], v[6:7], s[16:17]
	v_add_f64 v[2:3], v[2:3], s[0:1]
	s_delay_alu instid0(VALU_DEP_4)
	v_dual_mov_b32 v12, s3 :: v_dual_mov_b32 v7, v5
	v_mov_b32_e32 v6, v4
	v_dual_mov_b32 v4, v10 :: v_dual_mov_b32 v5, v11
.LBB0_4:                                ; %Flow
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s7
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s7, exec_lo, s6
	s_or_b32 s21, s7, s21
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s21
	s_cbranch_execnz .LBB0_2
; %bb.5:                                ; %_Z12doIterationsddddddddj.exit
	s_or_b32 exec_lo, exec_lo, s21
	v_cvt_f64_u32_e32 v[2:3], v12
	v_add_co_u32 v0, vcc_lo, s4, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s5, v1, vcc_lo
	global_store_b64 v[0:1], v[2:3], off
.LBB0_6:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z24processMandelbrotElementPdPKdS1_S1_S1_ddddjj
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 336
		.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 22
		.amdhsa_next_free_sgpr 22
		.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	_Z24processMandelbrotElementPdPKdS1_S1_S1_ddddjj, .Lfunc_end0-_Z24processMandelbrotElementPdPKdS1_S1_S1_ddddjj
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 664
; NumSgprs: 24
; NumVgprs: 22
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 2
; NumSGPRsForWavesPerEU: 24
; NumVGPRsForWavesPerEU: 22
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 14
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 1
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .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
      - .offset:         40
        .size:           8
        .value_kind:     by_value
      - .offset:         48
        .size:           8
        .value_kind:     by_value
      - .offset:         56
        .size:           8
        .value_kind:     by_value
      - .offset:         64
        .size:           8
        .value_kind:     by_value
      - .offset:         72
        .size:           4
        .value_kind:     by_value
      - .offset:         76
        .size:           4
        .value_kind:     by_value
      - .offset:         80
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         84
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         88
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         92
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         94
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         96
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         98
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         100
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         102
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         120
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         128
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         136
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         144
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 336
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z24processMandelbrotElementPdPKdS1_S1_S1_ddddjj
    .private_segment_fixed_size: 0
    .sgpr_count:     24
    .sgpr_spill_count: 0
    .symbol:         _Z24processMandelbrotElementPdPKdS1_S1_S1_ddddjj.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     22
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
