#include <stdio.h>
#include <stdlib.h>
#include <math.h>

#define GRADIENT_SIZE 16

#define C_X_MIN -0.188
#define C_X_MAX -0.012
#define C_Y_MIN 0.554
#define C_Y_MAX 0.754

#define IMAGE_SIZE 4096
#define ARRAY_SIZE (3 * IMAGE_SIZE * IMAGE_SIZE * sizeof(unsigned char)) 

#define PIXEL_WIDTH ((C_X_MAX - C_X_MIN) / IMAGE_SIZE)
#define PIXEL_HEIGHT ((C_Y_MAX - C_Y_MIN) / IMAGE_SIZE)
#define ITERATION_MAX 200

int colors[51] = {
                    66, 30, 15,
                    25, 7, 26,
                    9, 1, 47,
                    4, 4, 73,
                    0, 7, 100,
                    12, 44, 138,
                    24, 82, 177,
                    57, 125, 209,
                    134, 181, 229,
                    211, 236, 248,
                    241, 233, 191,
                    248, 201, 95,
                    255, 170, 0,
                    204, 128, 0,
                    153, 87, 0,
                    106, 52, 3,
                    16, 16, 16,
                    };
int *d_colors;

int x_grid;
int y_grid;
int x_block;
int y_block;

unsigned char *image_buffer;
unsigned char *d_image_buffer;

cudaEvent_t allocation_start, allocation_end;
cudaEvent_t computing_start, computing_end;
cudaEvent_t memcpy_start, memcpy_end;
float allocation_time, computing_time, memcpy_time;

void init(int argc, char *argv[])
{
    if (argc != 5) {
        printf("usage: ./mandelbrot_cu x_grid y_grid x_blocks y_blocks");
        exit(0);
    } 
    sscanf(argv[1], "%d", &x_grid);
    sscanf(argv[2], "%d", &y_grid);
    sscanf(argv[3], "%d", &x_block);
    sscanf(argv[4], "%d", &y_block);

    cudaEventCreate(&allocation_start); cudaEventCreate(&allocation_end); 
    cudaEventCreate(&computing_start); cudaEventCreate(&computing_end); 
    cudaEventCreate(&memcpy_start); cudaEventCreate(&memcpy_end); 

    cudaEventRecord(allocation_start, 0);

    cudaMallocHost((void **) &image_buffer, ARRAY_SIZE);
    cudaMalloc((void **) &d_image_buffer, ARRAY_SIZE);

    cudaMalloc((void **) &d_colors, 51 * sizeof(int));
    cudaMemcpy(d_colors, colors, 51 * sizeof(int), cudaMemcpyHostToDevice);
    
    cudaEventRecord(allocation_end, 0);
}

void write_to_file()
{
    FILE * file;
    const char *filename = "output.ppm";
    const char *comment  = "# ";

    int max_color_component_value = 255;

    file = fopen(filename,"wb");

    fprintf(file, "P6\n %s\n %d\n %d\n %d\n", comment,
            IMAGE_SIZE, IMAGE_SIZE, max_color_component_value);

    for(int i = 0; i < IMAGE_SIZE * IMAGE_SIZE; i++){
        fwrite(&image_buffer[3 * i], 1 , 3, file);
    };

    fclose(file);
}

__global__ void gpu_compute_mandelbrot(unsigned char *buffer, int *colors_d)
{
    double z_x = 0.0;
    double z_y = 0.0;
    double z_x_squared = 0.0;
    double z_y_squared = 0.0;
    double escape_radius_squared = 4;

    double c_x;
    double c_y;
    
    int i_y = blockIdx.y * blockDim.y + threadIdx.y; 
    int i_x = blockIdx.x * blockDim.x + threadIdx.x;

    int color;
    int iteration;

    c_y = C_Y_MIN + i_y * PIXEL_HEIGHT;
    if (fabs(c_y) < PIXEL_HEIGHT / 2)
        c_y = 0.0;
    c_x = C_X_MIN + i_x * PIXEL_WIDTH;
    for (iteration = 0;
                iteration < ITERATION_MAX && \
                ((z_x_squared + z_y_squared) < escape_radius_squared);
                iteration++) {
                z_y         = 2 * z_x * z_y + c_y;
                z_x         = z_x_squared - z_y_squared + c_x;
                z_x_squared = z_x * z_x;
                z_y_squared = z_y * z_y;
    }
    color = (iteration == ITERATION_MAX) ? GRADIENT_SIZE : (iteration % GRADIENT_SIZE);
    for (int i = 0; i < 3; i++) {
        buffer[3 * ((IMAGE_SIZE * i_y) + i_x) + i] = colors_d[(3 * color) + i];
    }
}

void compute_mandelbrot()
{
    gpu_compute_mandelbrot<<<dim3(x_grid, y_grid), dim3(x_block, y_block)>>>(d_image_buffer, d_colors);
    cudaDeviceSynchronize();
    cudaEventRecord(memcpy_start, 0);
    cudaMemcpy(image_buffer, d_image_buffer, ARRAY_SIZE, cudaMemcpyDeviceToHost);
    cudaEventRecord(memcpy_end, 0);
}

int main(int argc, char *argv[])
{
    init(argc, argv);

    cudaEventRecord(computing_start, 0);
    compute_mandelbrot();
    cudaEventRecord(computing_end, 0);

    write_to_file();

    cudaFreeHost(image_buffer); cudaFree(d_image_buffer);
    cudaFree(d_colors);

    cudaEventElapsedTime(&allocation_time, allocation_start, allocation_end);
    cudaEventElapsedTime(&computing_time, computing_start, computing_end);
    cudaEventElapsedTime(&memcpy_time, memcpy_start, memcpy_end);

    printf("%.6f\n", computing_time / 1000.0);
    return 0;
}
	.file	"tmpxft_0024cb06_00000000-6_mandelbrot_cu.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
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"usage: ./mandelbrot_cu x_grid y_grid x_blocks y_blocks"
.LC1:
	.string	"%d"
	.text
	.globl	_Z4initiPPc
	.type	_Z4initiPPc, @function
_Z4initiPPc:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	pushq	%rdx
	.cfi_def_cfa_offset 32
	cmpl	$5, %edi
	je	.L3
	movl	$2, %edi
	leaq	.LC0(%rip), %rsi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	xorl	%edi, %edi
	call	exit@PLT
.L3:
	leaq	.LC1(%rip), %rbp
	movq	8(%rsi), %rdi
	movq	%rsi, %rbx
	xorl	%eax, %eax
	movq	%rbp, %rsi
	leaq	x_grid(%rip), %rdx
	call	__isoc23_sscanf@PLT
	movq	16(%rbx), %rdi
	movq	%rbp, %rsi
	xorl	%eax, %eax
	leaq	y_grid(%rip), %rdx
	call	__isoc23_sscanf@PLT
	movq	24(%rbx), %rdi
	movq	%rbp, %rsi
	xorl	%eax, %eax
	leaq	x_block(%rip), %rdx
	call	__isoc23_sscanf@PLT
	movq	32(%rbx), %rdi
	movq	%rbp, %rsi
	xorl	%eax, %eax
	leaq	y_block(%rip), %rdx
	call	__isoc23_sscanf@PLT
	leaq	allocation_start(%rip), %rdi
	call	cudaEventCreate@PLT
	leaq	allocation_end(%rip), %rdi
	call	cudaEventCreate@PLT
	leaq	computing_start(%rip), %rdi
	call	cudaEventCreate@PLT
	leaq	computing_end(%rip), %rdi
	call	cudaEventCreate@PLT
	leaq	memcpy_start(%rip), %rdi
	call	cudaEventCreate@PLT
	leaq	memcpy_end(%rip), %rdi
	call	cudaEventCreate@PLT
	movq	allocation_start(%rip), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movl	$50331648, %esi
	leaq	image_buffer(%rip), %rdi
	call	cudaMallocHost@PLT
	movl	$50331648, %esi
	leaq	d_image_buffer(%rip), %rdi
	call	cudaMalloc@PLT
	movl	$204, %esi
	leaq	d_colors(%rip), %rdi
	call	cudaMalloc@PLT
	movl	$1, %ecx
	movl	$204, %edx
	movq	d_colors(%rip), %rdi
	leaq	colors(%rip), %rsi
	call	cudaMemcpy@PLT
	movq	allocation_end(%rip), %rdi
	xorl	%esi, %esi
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	jmp	cudaEventRecord@PLT
	.cfi_endproc
.LFE2027:
	.size	_Z4initiPPc, .-_Z4initiPPc
	.section	.rodata.str1.1
.LC2:
	.string	"wb"
.LC3:
	.string	"output.ppm"
.LC4:
	.string	"# "
.LC5:
	.string	"P6\n %s\n %d\n %d\n %d\n"
	.text
	.globl	_Z13write_to_filev
	.type	_Z13write_to_filev, @function
_Z13write_to_filev:
.LFB2028:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	leaq	.LC2(%rip), %rsi
	leaq	.LC3(%rip), %rdi
	xorl	%ebp, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	pushq	%rdx
	.cfi_def_cfa_offset 32
	call	fopen@PLT
	movl	$2, %esi
	movl	$4096, %r9d
	movl	$4096, %r8d
	pushq	%rcx
	.cfi_def_cfa_offset 40
	movq	%rax, %rdi
	movq	%rax, %rbx
	leaq	.LC4(%rip), %rcx
	pushq	$255
	.cfi_def_cfa_offset 48
	leaq	.LC5(%rip), %rdx
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
	popq	%rsi
	.cfi_def_cfa_offset 40
	popq	%rdi
	.cfi_def_cfa_offset 32
.L7:
	movq	image_buffer(%rip), %rdi
	movq	%rbx, %rcx
	movl	$3, %edx
	movl	$1, %esi
	addq	%rbp, %rdi
	addq	$3, %rbp
	call	fwrite@PLT
	cmpq	$50331648, %rbp
	jne	.L7
	popq	%rax
	.cfi_def_cfa_offset 24
	movq	%rbx, %rdi
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	jmp	fclose@PLT
	.cfi_endproc
.LFE2028:
	.size	_Z13write_to_filev, .-_Z13write_to_filev
	.globl	_Z44__device_stub__Z22gpu_compute_mandelbrotPhPiPhPi
	.type	_Z44__device_stub__Z22gpu_compute_mandelbrotPhPiPhPi, @function
_Z44__device_stub__Z22gpu_compute_mandelbrotPhPiPhPi:
.LFB2055:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 8(%rsp)
	leaq	32(%rsp), %rcx
	leaq	24(%rsp), %rdx
	movq	%rsi, (%rsp)
	leaq	40(%rsp), %rdi
	leaq	52(%rsp), %rsi
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 48(%rsp)
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 40(%rsp)
	movq	%rax, 52(%rsp)
	movl	$1, 60(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L10
	pushq	32(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z22gpu_compute_mandelbrotPhPi(%rip), %rdi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 144
	movq	68(%rsp), %rcx
	movl	76(%rsp), %r8d
	movq	56(%rsp), %rsi
	movl	64(%rsp), %edx
	leaq	104(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L10:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L12
	call	__stack_chk_fail@PLT
.L12:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2055:
	.size	_Z44__device_stub__Z22gpu_compute_mandelbrotPhPiPhPi, .-_Z44__device_stub__Z22gpu_compute_mandelbrotPhPiPhPi
	.globl	_Z22gpu_compute_mandelbrotPhPi
	.type	_Z22gpu_compute_mandelbrotPhPi, @function
_Z22gpu_compute_mandelbrotPhPi:
.LFB2056:
	.cfi_startproc
	endbr64
	jmp	_Z44__device_stub__Z22gpu_compute_mandelbrotPhPiPhPi
	.cfi_endproc
.LFE2056:
	.size	_Z22gpu_compute_mandelbrotPhPi, .-_Z22gpu_compute_mandelbrotPhPi
	.globl	_Z18compute_mandelbrotv
	.type	_Z18compute_mandelbrotv, @function
_Z18compute_mandelbrotv:
.LFB2029:
	.cfi_startproc
	endbr64
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movl	x_block(%rip), %eax
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	movl	$1, %esi
	movl	%eax, 20(%rsp)
	movl	y_block(%rip), %eax
	movl	%eax, 24(%rsp)
	movl	x_grid(%rip), %eax
	movq	20(%rsp), %rdx
	movl	%eax, 8(%rsp)
	movl	y_grid(%rip), %eax
	movl	%eax, 12(%rsp)
	movq	8(%rsp), %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L16
	movq	d_colors(%rip), %rsi
	movq	d_image_buffer(%rip), %rdi
	call	_Z44__device_stub__Z22gpu_compute_mandelbrotPhPiPhPi
.L16:
	call	cudaDeviceSynchronize@PLT
	movq	memcpy_start(%rip), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movl	$2, %ecx
	movl	$50331648, %edx
	movq	d_image_buffer(%rip), %rsi
	movq	image_buffer(%rip), %rdi
	call	cudaMemcpy@PLT
	movq	memcpy_end(%rip), %rdi
	xorl	%esi, %esi
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	jmp	cudaEventRecord@PLT
	.cfi_endproc
.LFE2029:
	.size	_Z18compute_mandelbrotv, .-_Z18compute_mandelbrotv
	.section	.rodata.str1.1
.LC7:
	.string	"%.6f\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2030:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_Z4initiPPc
	movq	computing_start(%rip), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	call	_Z18compute_mandelbrotv
	movq	computing_end(%rip), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	call	_Z13write_to_filev
	movq	image_buffer(%rip), %rdi
	call	cudaFreeHost@PLT
	movq	d_image_buffer(%rip), %rdi
	call	cudaFree@PLT
	movq	d_colors(%rip), %rdi
	call	cudaFree@PLT
	movq	allocation_end(%rip), %rdx
	movq	allocation_start(%rip), %rsi
	leaq	allocation_time(%rip), %rdi
	call	cudaEventElapsedTime@PLT
	movq	computing_end(%rip), %rdx
	movq	computing_start(%rip), %rsi
	leaq	computing_time(%rip), %rdi
	call	cudaEventElapsedTime@PLT
	movq	memcpy_end(%rip), %rdx
	movq	memcpy_start(%rip), %rsi
	leaq	memcpy_time(%rip), %rdi
	call	cudaEventElapsedTime@PLT
	leaq	.LC7(%rip), %rsi
	movb	$1, %al
	movl	$2, %edi
	cvtss2sd	computing_time(%rip), %xmm0
	divsd	.LC6(%rip), %xmm0
	call	__printf_chk@PLT
	xorl	%eax, %eax
	popq	%rdx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2030:
	.size	main, .-main
	.section	.rodata.str1.1
.LC8:
	.string	"_Z22gpu_compute_mandelbrotPhPi"
	.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	.LC8(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z22gpu_compute_mandelbrotPhPi(%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
	.globl	memcpy_time
	.bss
	.align 4
	.type	memcpy_time, @object
	.size	memcpy_time, 4
memcpy_time:
	.zero	4
	.globl	computing_time
	.align 4
	.type	computing_time, @object
	.size	computing_time, 4
computing_time:
	.zero	4
	.globl	allocation_time
	.align 4
	.type	allocation_time, @object
	.size	allocation_time, 4
allocation_time:
	.zero	4
	.globl	memcpy_end
	.align 8
	.type	memcpy_end, @object
	.size	memcpy_end, 8
memcpy_end:
	.zero	8
	.globl	memcpy_start
	.align 8
	.type	memcpy_start, @object
	.size	memcpy_start, 8
memcpy_start:
	.zero	8
	.globl	computing_end
	.align 8
	.type	computing_end, @object
	.size	computing_end, 8
computing_end:
	.zero	8
	.globl	computing_start
	.align 8
	.type	computing_start, @object
	.size	computing_start, 8
computing_start:
	.zero	8
	.globl	allocation_end
	.align 8
	.type	allocation_end, @object
	.size	allocation_end, 8
allocation_end:
	.zero	8
	.globl	allocation_start
	.align 8
	.type	allocation_start, @object
	.size	allocation_start, 8
allocation_start:
	.zero	8
	.globl	d_image_buffer
	.align 8
	.type	d_image_buffer, @object
	.size	d_image_buffer, 8
d_image_buffer:
	.zero	8
	.globl	image_buffer
	.align 8
	.type	image_buffer, @object
	.size	image_buffer, 8
image_buffer:
	.zero	8
	.globl	y_block
	.align 4
	.type	y_block, @object
	.size	y_block, 4
y_block:
	.zero	4
	.globl	x_block
	.align 4
	.type	x_block, @object
	.size	x_block, 4
x_block:
	.zero	4
	.globl	y_grid
	.align 4
	.type	y_grid, @object
	.size	y_grid, 4
y_grid:
	.zero	4
	.globl	x_grid
	.align 4
	.type	x_grid, @object
	.size	x_grid, 4
x_grid:
	.zero	4
	.globl	d_colors
	.align 8
	.type	d_colors, @object
	.size	d_colors, 8
d_colors:
	.zero	8
	.globl	colors
	.data
	.align 32
	.type	colors, @object
	.size	colors, 204
colors:
	.long	66
	.long	30
	.long	15
	.long	25
	.long	7
	.long	26
	.long	9
	.long	1
	.long	47
	.long	4
	.long	4
	.long	73
	.long	0
	.long	7
	.long	100
	.long	12
	.long	44
	.long	138
	.long	24
	.long	82
	.long	177
	.long	57
	.long	125
	.long	209
	.long	134
	.long	181
	.long	229
	.long	211
	.long	236
	.long	248
	.long	241
	.long	233
	.long	191
	.long	248
	.long	201
	.long	95
	.long	255
	.long	170
	.long	0
	.long	204
	.long	128
	.long	0
	.long	153
	.long	87
	.long	0
	.long	106
	.long	52
	.long	3
	.long	16
	.long	16
	.long	16
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC6:
	.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 <stdio.h>
#include <stdlib.h>
#include <math.h>

#define GRADIENT_SIZE 16

#define C_X_MIN -0.188
#define C_X_MAX -0.012
#define C_Y_MIN 0.554
#define C_Y_MAX 0.754

#define IMAGE_SIZE 4096
#define ARRAY_SIZE (3 * IMAGE_SIZE * IMAGE_SIZE * sizeof(unsigned char)) 

#define PIXEL_WIDTH ((C_X_MAX - C_X_MIN) / IMAGE_SIZE)
#define PIXEL_HEIGHT ((C_Y_MAX - C_Y_MIN) / IMAGE_SIZE)
#define ITERATION_MAX 200

int colors[51] = {
                    66, 30, 15,
                    25, 7, 26,
                    9, 1, 47,
                    4, 4, 73,
                    0, 7, 100,
                    12, 44, 138,
                    24, 82, 177,
                    57, 125, 209,
                    134, 181, 229,
                    211, 236, 248,
                    241, 233, 191,
                    248, 201, 95,
                    255, 170, 0,
                    204, 128, 0,
                    153, 87, 0,
                    106, 52, 3,
                    16, 16, 16,
                    };
int *d_colors;

int x_grid;
int y_grid;
int x_block;
int y_block;

unsigned char *image_buffer;
unsigned char *d_image_buffer;

hipEvent_t allocation_start, allocation_end;
hipEvent_t computing_start, computing_end;
hipEvent_t memcpy_start, memcpy_end;
float allocation_time, computing_time, memcpy_time;

void init(int argc, char *argv[])
{
    if (argc != 5) {
        printf("usage: ./mandelbrot_cu x_grid y_grid x_blocks y_blocks");
        exit(0);
    } 
    sscanf(argv[1], "%d", &x_grid);
    sscanf(argv[2], "%d", &y_grid);
    sscanf(argv[3], "%d", &x_block);
    sscanf(argv[4], "%d", &y_block);

    hipEventCreate(&allocation_start); hipEventCreate(&allocation_end); 
    hipEventCreate(&computing_start); hipEventCreate(&computing_end); 
    hipEventCreate(&memcpy_start); hipEventCreate(&memcpy_end); 

    hipEventRecord(allocation_start, 0);

    hipHostMalloc((void **) &image_buffer, ARRAY_SIZE, hipHostMallocDefault);
    hipMalloc((void **) &d_image_buffer, ARRAY_SIZE);

    hipMalloc((void **) &d_colors, 51 * sizeof(int));
    hipMemcpy(d_colors, colors, 51 * sizeof(int), hipMemcpyHostToDevice);
    
    hipEventRecord(allocation_end, 0);
}

void write_to_file()
{
    FILE * file;
    const char *filename = "output.ppm";
    const char *comment  = "# ";

    int max_color_component_value = 255;

    file = fopen(filename,"wb");

    fprintf(file, "P6\n %s\n %d\n %d\n %d\n", comment,
            IMAGE_SIZE, IMAGE_SIZE, max_color_component_value);

    for(int i = 0; i < IMAGE_SIZE * IMAGE_SIZE; i++){
        fwrite(&image_buffer[3 * i], 1 , 3, file);
    };

    fclose(file);
}

__global__ void gpu_compute_mandelbrot(unsigned char *buffer, int *colors_d)
{
    double z_x = 0.0;
    double z_y = 0.0;
    double z_x_squared = 0.0;
    double z_y_squared = 0.0;
    double escape_radius_squared = 4;

    double c_x;
    double c_y;
    
    int i_y = blockIdx.y * blockDim.y + threadIdx.y; 
    int i_x = blockIdx.x * blockDim.x + threadIdx.x;

    int color;
    int iteration;

    c_y = C_Y_MIN + i_y * PIXEL_HEIGHT;
    if (fabs(c_y) < PIXEL_HEIGHT / 2)
        c_y = 0.0;
    c_x = C_X_MIN + i_x * PIXEL_WIDTH;
    for (iteration = 0;
                iteration < ITERATION_MAX && \
                ((z_x_squared + z_y_squared) < escape_radius_squared);
                iteration++) {
                z_y         = 2 * z_x * z_y + c_y;
                z_x         = z_x_squared - z_y_squared + c_x;
                z_x_squared = z_x * z_x;
                z_y_squared = z_y * z_y;
    }
    color = (iteration == ITERATION_MAX) ? GRADIENT_SIZE : (iteration % GRADIENT_SIZE);
    for (int i = 0; i < 3; i++) {
        buffer[3 * ((IMAGE_SIZE * i_y) + i_x) + i] = colors_d[(3 * color) + i];
    }
}

void compute_mandelbrot()
{
    gpu_compute_mandelbrot<<<dim3(x_grid, y_grid), dim3(x_block, y_block)>>>(d_image_buffer, d_colors);
    hipDeviceSynchronize();
    hipEventRecord(memcpy_start, 0);
    hipMemcpy(image_buffer, d_image_buffer, ARRAY_SIZE, hipMemcpyDeviceToHost);
    hipEventRecord(memcpy_end, 0);
}

int main(int argc, char *argv[])
{
    init(argc, argv);

    hipEventRecord(computing_start, 0);
    compute_mandelbrot();
    hipEventRecord(computing_end, 0);

    write_to_file();

    hipHostFree(image_buffer); hipFree(d_image_buffer);
    hipFree(d_colors);

    hipEventElapsedTime(&allocation_time, allocation_start, allocation_end);
    hipEventElapsedTime(&computing_time, computing_start, computing_end);
    hipEventElapsedTime(&memcpy_time, memcpy_start, memcpy_end);

    printf("%.6f\n", computing_time / 1000.0);
    return 0;
}
	.text
	.file	"mandelbrot_cu.hip"
	.globl	_Z4initiPPc                     # -- Begin function _Z4initiPPc
	.type	_Z4initiPPc,@function
_Z4initiPPc:                            # @_Z4initiPPc
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset %rbx, -16
	cmpl	$5, %edi
	jne	.LBB0_1
# %bb.2:
	movq	%rsi, %rbx
	movq	8(%rsi), %rdi
	movl	$.L.str.1, %esi
	movl	$x_grid, %edx
	xorl	%eax, %eax
	callq	__isoc23_sscanf
	movq	16(%rbx), %rdi
	movl	$.L.str.1, %esi
	movl	$y_grid, %edx
	xorl	%eax, %eax
	callq	__isoc23_sscanf
	movq	24(%rbx), %rdi
	movl	$.L.str.1, %esi
	movl	$x_block, %edx
	xorl	%eax, %eax
	callq	__isoc23_sscanf
	movq	32(%rbx), %rdi
	movl	$.L.str.1, %esi
	movl	$y_block, %edx
	xorl	%eax, %eax
	callq	__isoc23_sscanf
	movl	$allocation_start, %edi
	callq	hipEventCreate
	movl	$allocation_end, %edi
	callq	hipEventCreate
	movl	$computing_start, %edi
	callq	hipEventCreate
	movl	$computing_end, %edi
	callq	hipEventCreate
	movl	$memcpy_start, %edi
	callq	hipEventCreate
	movl	$memcpy_end, %edi
	callq	hipEventCreate
	movq	allocation_start(%rip), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movl	$image_buffer, %edi
	movl	$50331648, %esi                 # imm = 0x3000000
	xorl	%edx, %edx
	callq	hipHostMalloc
	movl	$d_image_buffer, %edi
	movl	$50331648, %esi                 # imm = 0x3000000
	callq	hipMalloc
	movl	$d_colors, %edi
	movl	$204, %esi
	callq	hipMalloc
	movq	d_colors(%rip), %rdi
	movl	$colors, %esi
	movl	$204, %edx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	allocation_end(%rip), %rdi
	xorl	%esi, %esi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	hipEventRecord                  # TAILCALL
.LBB0_1:
	.cfi_def_cfa_offset 16
	movl	$.L.str, %edi
	xorl	%eax, %eax
	callq	printf
	xorl	%edi, %edi
	callq	exit
.Lfunc_end0:
	.size	_Z4initiPPc, .Lfunc_end0-_Z4initiPPc
	.cfi_endproc
                                        # -- End function
	.globl	_Z13write_to_filev              # -- Begin function _Z13write_to_filev
	.type	_Z13write_to_filev,@function
_Z13write_to_filev:                     # @_Z13write_to_filev
	.cfi_startproc
# %bb.0:
	pushq	%r14
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	pushq	%rax
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -24
	.cfi_offset %r14, -16
	movl	$.L.str.2, %edi
	movl	$.L.str.4, %esi
	callq	fopen
	movq	%rax, %rbx
	xorl	%r14d, %r14d
	movl	$.L.str.5, %esi
	movl	$.L.str.3, %edx
	movq	%rax, %rdi
	movl	$4096, %ecx                     # imm = 0x1000
	movl	$4096, %r8d                     # imm = 0x1000
	movl	$255, %r9d
	xorl	%eax, %eax
	callq	fprintf
.LBB1_1:                                # =>This Inner Loop Header: Depth=1
	movq	image_buffer(%rip), %rdi
	addq	%r14, %rdi
	movl	$1, %esi
	movl	$3, %edx
	movq	%rbx, %rcx
	callq	fwrite
	addq	$3, %r14
	cmpq	$50331648, %r14                 # imm = 0x3000000
	jne	.LBB1_1
# %bb.2:
	movq	%rbx, %rdi
	addq	$8, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	jmp	fclose                          # TAILCALL
.Lfunc_end1:
	.size	_Z13write_to_filev, .Lfunc_end1-_Z13write_to_filev
	.cfi_endproc
                                        # -- End function
	.globl	_Z37__device_stub__gpu_compute_mandelbrotPhPi # -- Begin function _Z37__device_stub__gpu_compute_mandelbrotPhPi
	.type	_Z37__device_stub__gpu_compute_mandelbrotPhPi,@function
_Z37__device_stub__gpu_compute_mandelbrotPhPi: # @_Z37__device_stub__gpu_compute_mandelbrotPhPi
	.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	$80, %rsp
	.cfi_def_cfa_offset 128
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	24(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	16(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	64(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%rsp), %r15
	leaq	8(%rsp), %r12
	movq	%rsp, %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z22gpu_compute_mandelbrotPhPi, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$96, %rsp
	.cfi_adjust_cfa_offset -96
	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_end2:
	.size	_Z37__device_stub__gpu_compute_mandelbrotPhPi, .Lfunc_end2-_Z37__device_stub__gpu_compute_mandelbrotPhPi
	.cfi_endproc
                                        # -- End function
	.globl	_Z18compute_mandelbrotv         # -- Begin function _Z18compute_mandelbrotv
	.type	_Z18compute_mandelbrotv,@function
_Z18compute_mandelbrotv:                # @_Z18compute_mandelbrotv
	.cfi_startproc
# %bb.0:
	pushq	%rax
	.cfi_def_cfa_offset 16
	movl	x_grid(%rip), %eax
	movl	x_block(%rip), %ecx
	movl	y_grid(%rip), %edi
	shlq	$32, %rdi
	orq	%rax, %rdi
	movl	y_block(%rip), %edx
	shlq	$32, %rdx
	orq	%rcx, %rdx
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB3_2
# %bb.1:
	movq	d_image_buffer(%rip), %rdi
	movq	d_colors(%rip), %rsi
	callq	_Z37__device_stub__gpu_compute_mandelbrotPhPi
.LBB3_2:
	callq	hipDeviceSynchronize
	movq	memcpy_start(%rip), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	image_buffer(%rip), %rdi
	movq	d_image_buffer(%rip), %rsi
	movl	$50331648, %edx                 # imm = 0x3000000
	movl	$2, %ecx
	callq	hipMemcpy
	movq	memcpy_end(%rip), %rdi
	xorl	%esi, %esi
	popq	%rax
	.cfi_def_cfa_offset 8
	jmp	hipEventRecord                  # TAILCALL
.Lfunc_end3:
	.size	_Z18compute_mandelbrotv, .Lfunc_end3-_Z18compute_mandelbrotv
	.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	%rax
	.cfi_def_cfa_offset 16
	callq	_Z4initiPPc
	movq	computing_start(%rip), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	callq	_Z18compute_mandelbrotv
	movq	computing_end(%rip), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	callq	_Z13write_to_filev
	movq	image_buffer(%rip), %rdi
	callq	hipHostFree
	movq	d_image_buffer(%rip), %rdi
	callq	hipFree
	movq	d_colors(%rip), %rdi
	callq	hipFree
	movq	allocation_start(%rip), %rsi
	movq	allocation_end(%rip), %rdx
	movl	$allocation_time, %edi
	callq	hipEventElapsedTime
	movq	computing_start(%rip), %rsi
	movq	computing_end(%rip), %rdx
	movl	$computing_time, %edi
	callq	hipEventElapsedTime
	movq	memcpy_start(%rip), %rsi
	movq	memcpy_end(%rip), %rdx
	movl	$memcpy_time, %edi
	callq	hipEventElapsedTime
	cvtss2sd	computing_time(%rip), %xmm0
	divsd	.LCPI4_0(%rip), %xmm0
	movl	$.L.str.6, %edi
	movb	$1, %al
	callq	printf
	xorl	%eax, %eax
	popq	%rcx
	.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	$_Z22gpu_compute_mandelbrotPhPi, %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	colors,@object                  # @colors
	.data
	.globl	colors
	.p2align	4, 0x0
colors:
	.long	66                              # 0x42
	.long	30                              # 0x1e
	.long	15                              # 0xf
	.long	25                              # 0x19
	.long	7                               # 0x7
	.long	26                              # 0x1a
	.long	9                               # 0x9
	.long	1                               # 0x1
	.long	47                              # 0x2f
	.long	4                               # 0x4
	.long	4                               # 0x4
	.long	73                              # 0x49
	.long	0                               # 0x0
	.long	7                               # 0x7
	.long	100                             # 0x64
	.long	12                              # 0xc
	.long	44                              # 0x2c
	.long	138                             # 0x8a
	.long	24                              # 0x18
	.long	82                              # 0x52
	.long	177                             # 0xb1
	.long	57                              # 0x39
	.long	125                             # 0x7d
	.long	209                             # 0xd1
	.long	134                             # 0x86
	.long	181                             # 0xb5
	.long	229                             # 0xe5
	.long	211                             # 0xd3
	.long	236                             # 0xec
	.long	248                             # 0xf8
	.long	241                             # 0xf1
	.long	233                             # 0xe9
	.long	191                             # 0xbf
	.long	248                             # 0xf8
	.long	201                             # 0xc9
	.long	95                              # 0x5f
	.long	255                             # 0xff
	.long	170                             # 0xaa
	.long	0                               # 0x0
	.long	204                             # 0xcc
	.long	128                             # 0x80
	.long	0                               # 0x0
	.long	153                             # 0x99
	.long	87                              # 0x57
	.long	0                               # 0x0
	.long	106                             # 0x6a
	.long	52                              # 0x34
	.long	3                               # 0x3
	.long	16                              # 0x10
	.long	16                              # 0x10
	.long	16                              # 0x10
	.size	colors, 204

	.type	d_colors,@object                # @d_colors
	.bss
	.globl	d_colors
	.p2align	3, 0x0
d_colors:
	.quad	0
	.size	d_colors, 8

	.type	x_grid,@object                  # @x_grid
	.globl	x_grid
	.p2align	2, 0x0
x_grid:
	.long	0                               # 0x0
	.size	x_grid, 4

	.type	y_grid,@object                  # @y_grid
	.globl	y_grid
	.p2align	2, 0x0
y_grid:
	.long	0                               # 0x0
	.size	y_grid, 4

	.type	x_block,@object                 # @x_block
	.globl	x_block
	.p2align	2, 0x0
x_block:
	.long	0                               # 0x0
	.size	x_block, 4

	.type	y_block,@object                 # @y_block
	.globl	y_block
	.p2align	2, 0x0
y_block:
	.long	0                               # 0x0
	.size	y_block, 4

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

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

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

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

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

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

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

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

	.type	allocation_time,@object         # @allocation_time
	.globl	allocation_time
	.p2align	2, 0x0
allocation_time:
	.long	0x00000000                      # float 0
	.size	allocation_time, 4

	.type	computing_time,@object          # @computing_time
	.globl	computing_time
	.p2align	2, 0x0
computing_time:
	.long	0x00000000                      # float 0
	.size	computing_time, 4

	.type	memcpy_time,@object             # @memcpy_time
	.globl	memcpy_time
	.p2align	2, 0x0
memcpy_time:
	.long	0x00000000                      # float 0
	.size	memcpy_time, 4

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"usage: ./mandelbrot_cu x_grid y_grid x_blocks y_blocks"
	.size	.L.str, 55

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"%d"
	.size	.L.str.1, 3

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"output.ppm"
	.size	.L.str.2, 11

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"# "
	.size	.L.str.3, 3

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"wb"
	.size	.L.str.4, 3

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"P6\n %s\n %d\n %d\n %d\n"
	.size	.L.str.5, 20

	.type	_Z22gpu_compute_mandelbrotPhPi,@object # @_Z22gpu_compute_mandelbrotPhPi
	.section	.rodata,"a",@progbits
	.globl	_Z22gpu_compute_mandelbrotPhPi
	.p2align	3, 0x0
_Z22gpu_compute_mandelbrotPhPi:
	.quad	_Z37__device_stub__gpu_compute_mandelbrotPhPi
	.size	_Z22gpu_compute_mandelbrotPhPi, 8

	.type	.L.str.6,@object                # @.str.6
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.6:
	.asciz	"%.6f\n"
	.size	.L.str.6, 6

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z22gpu_compute_mandelbrotPhPi"
	.size	.L__unnamed_1, 31

	.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 _Z37__device_stub__gpu_compute_mandelbrotPhPi
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym colors
	.addrsig_sym d_colors
	.addrsig_sym x_grid
	.addrsig_sym y_grid
	.addrsig_sym x_block
	.addrsig_sym y_block
	.addrsig_sym image_buffer
	.addrsig_sym d_image_buffer
	.addrsig_sym allocation_start
	.addrsig_sym allocation_end
	.addrsig_sym computing_start
	.addrsig_sym computing_end
	.addrsig_sym memcpy_start
	.addrsig_sym memcpy_end
	.addrsig_sym allocation_time
	.addrsig_sym computing_time
	.addrsig_sym memcpy_time
	.addrsig_sym _Z22gpu_compute_mandelbrotPhPi
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z22gpu_compute_mandelbrotPhPi ; -- Begin function _Z22gpu_compute_mandelbrotPhPi
	.globl	_Z22gpu_compute_mandelbrotPhPi
	.p2align	8
	.type	_Z22gpu_compute_mandelbrotPhPi,@function
_Z22gpu_compute_mandelbrotPhPi:         ; @_Z22gpu_compute_mandelbrotPhPi
; %bb.0:
	s_load_b32 s4, s[0:1], 0x1c
	v_bfe_u32 v3, v0, 10, 10
	v_and_b32_e32 v0, 0x3ff, v0
	s_mov_b32 s5, 0x3f099999
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s2, s4, 16
	s_and_b32 s4, s4, 0xffff
	v_mad_u64_u32 v[1:2], null, s15, s2, v[3:4]
	s_mov_b32 s2, 0x353f7cee
	s_mov_b32 s3, 0x3fe1ba5e
	s_delay_alu instid0(VALU_DEP_1)
	v_cvt_f64_i32_e32 v[4:5], v1
	v_mad_u64_u32 v[2:3], null, s14, s4, v[0:1]
	s_mov_b32 s4, 0x99999998
	s_delay_alu instid0(VALU_DEP_2) | instid1(SALU_CYCLE_1)
	v_fma_f64 v[9:10], v[4:5], s[4:5], s[2:3]
	v_mov_b32_e32 v5, 0
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_3) | instid1(VALU_DEP_3)
	v_cvt_f64_i32_e32 v[3:4], v2
	s_mov_b32 s5, 0x3ef99999
	s_mov_b32 s2, 0x4dd2f1aa
	s_mov_b32 s3, 0xbfc81062
	v_cmp_nlt_f64_e64 vcc_lo, |v[9:10]|, s[4:5]
	s_mov_b32 s4, 0x20c49ba
	s_mov_b32 s5, 0x3f06872b
	v_mov_b32_e32 v6, 0
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_fma_f64 v[3:4], v[3:4], s[4:5], s[2:3]
	s_mov_b32 s2, 0
	s_mov_b32 s3, 0
	v_dual_cndmask_b32 v9, 0, v9 :: v_dual_mov_b32 v14, v6
	v_dual_mov_b32 v8, v6 :: v_dual_mov_b32 v7, v5
	v_dual_mov_b32 v12, v6 :: v_dual_mov_b32 v11, v5
	v_dual_cndmask_b32 v10, 0, v10 :: v_dual_mov_b32 v13, v5
.LBB0_1:                                ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(VALU_DEP_1)
	v_add_f64 v[11:12], v[11:12], -v[13:14]
	v_add_f64 v[13:14], v[5:6], v[5:6]
	s_add_i32 s4, s3, 1
	s_cmpk_gt_u32 s3, 0xc6
	v_mov_b32_e32 v0, s4
	s_cselect_b32 s3, -1, 0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f64 v[5:6], v[3:4], v[11:12]
	v_fma_f64 v[7:8], v[7:8], v[13:14], v[9:10]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[11:12], v[5:6], v[5:6]
	v_fma_f64 v[13:14], v[7:8], v[7:8], v[11:12]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(SALU_CYCLE_1)
	v_cmp_ngt_f64_e32 vcc_lo, 4.0, v[13:14]
	v_mul_f64 v[13:14], v[7:8], v[7:8]
	s_or_b32 s3, s3, vcc_lo
	s_and_b32 s3, exec_lo, s3
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 s2, s3, s2
	s_mov_b32 s3, s4
	s_and_not1_b32 exec_lo, exec_lo, s2
	s_cbranch_execnz .LBB0_1
; %bb.2:                                ; %.critedge
	s_or_b32 exec_lo, exec_lo, s2
	s_load_b128 s[0:3], s[0:1], 0x0
	v_cmp_ne_u32_e32 vcc_lo, 0xc8, v0
	v_and_b32_e32 v3, 15, v0
	v_lshl_add_u32 v1, v1, 12, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_u32_u24_e32 v2, 3, v3
	v_lshl_add_u32 v1, v1, 1, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v0, 48, v2, vcc_lo
	v_ashrrev_i32_e32 v3, 31, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_lshlrev_b32_e32 v0, 2, v0
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v2, vcc_lo, s0, v1
	v_add_co_u32 v0, s0, s2, v0
	v_add_co_ci_u32_e32 v3, vcc_lo, s1, v3, vcc_lo
	v_add_co_ci_u32_e64 v1, null, s3, 0, s0
	s_mov_b64 s[0:1], 0
.LBB0_3:                                ; =>This Inner Loop Header: Depth=1
	global_load_b32 v6, v[0:1], off
	v_add_co_u32 v4, vcc_lo, v2, s0
	v_add_co_ci_u32_e32 v5, vcc_lo, s1, v3, vcc_lo
	v_add_co_u32 v0, vcc_lo, v0, 4
	v_add_co_ci_u32_e32 v1, vcc_lo, 0, v1, vcc_lo
	s_add_u32 s0, s0, 1
	s_addc_u32 s1, s1, 0
	s_cmp_eq_u32 s0, 3
	s_waitcnt vmcnt(0)
	global_store_b8 v[4:5], v6, off
	s_cbranch_scc0 .LBB0_3
; %bb.4:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z22gpu_compute_mandelbrotPhPi
		.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 15
		.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	_Z22gpu_compute_mandelbrotPhPi, .Lfunc_end0-_Z22gpu_compute_mandelbrotPhPi
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 520
; NumSgprs: 18
; NumVgprs: 15
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 15
; 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
      - .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:           _Z22gpu_compute_mandelbrotPhPi
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z22gpu_compute_mandelbrotPhPi.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     15
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
