#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include <pthread.h>
#include <time.h>
#include <sys/time.h>

/*Dados HOST*/
struct data{
    double c_x_min;
    double c_x_max;
    double c_y_min;
    double c_y_max;

    double pixel_width;
    double pixel_height;

    int i_x_max;
    int i_y_max;

    int n_blocks;
    int n_threads;
};

struct data * global_data;

int colors[17][3] = {
    {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 gradient_size = 16;
unsigned char **image_buffer;
int image_size;
int image_buffer_size;
//////////////////////////////////////////


/*Timer info*/
struct timer_info {
    clock_t c_start;
    clock_t c_end;
    struct timespec t_start;
    struct timespec t_end;
    struct timeval v_start;
    struct timeval v_end;
};

struct timer_info timer;
//////////////////////////////////////////

void allocate_image_buffer(){
    int rgb_size = 3;
    image_buffer = (unsigned char **) malloc(sizeof(unsigned char *) * image_buffer_size);
    for(int i = 0; i < image_buffer_size; i++){
        image_buffer[i] = (unsigned char *) malloc(sizeof(unsigned char) * rgb_size);
    };
};

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

    global_data = (struct data * )malloc(sizeof(struct data));

    if(argc < 8){
        printf("usage: ./mandelbrot_cuda c_x_min c_x_max c_y_min c_y_max image_size blocks threads\n");
        printf("examples with image_size = 11500:\n");
        printf("    Full Picture:         ./mandelbrot_cuda -2.5 1.5 -2.0 2.0 11500 4 256\n");
        printf("    Seahorse Valley:      ./mandelbrot_cuda -0.8 -0.7 0.05 0.15 11500 4 256\n");
        printf("    Elephant Valley:      ./mandelbrot_cuda 0.175 0.375 -0.1 0.1 11500 4 256\n");
        printf("    Triple Spiral Valley: ./mandelbrot_cuda -0.188 -0.012 0.554 0.754 11500 4 256\n");
        exit(0);
    }
    else{
        sscanf(argv[1], "%lf", &global_data->c_x_min);
        sscanf(argv[2], "%lf", &global_data->c_x_max);
        sscanf(argv[3], "%lf", &global_data->c_y_min);
        sscanf(argv[4], "%lf", &global_data->c_y_max);
        sscanf(argv[5], "%d", &image_size);
        sscanf(argv[6], "%d", &global_data->n_blocks);
        sscanf(argv[6], "%d", &global_data->n_threads);

        global_data->i_x_max           = image_size;
        global_data->i_y_max           = image_size;
        image_buffer_size = image_size * image_size;

        global_data->pixel_width       = (global_data->c_x_max - global_data->c_x_min) / global_data->i_x_max;
        global_data->pixel_height      = (global_data->c_y_max - global_data->c_y_min) / global_data->i_y_max;
    };

};

void update_rgb_buffer(int iteration, int x, int y){
    int color;

    int i_y_max = global_data->i_y_max;
    int iteration_max = 200;

    if(iteration == iteration_max){
        image_buffer[(i_y_max * y) + x][0] = colors[gradient_size][0];
        image_buffer[(i_y_max * y) + x][1] = colors[gradient_size][1];
        image_buffer[(i_y_max * y) + x][2] = colors[gradient_size][2];
    }
    else{
        color = iteration % gradient_size;

        image_buffer[(i_y_max * y) + x][0] = colors[color][0];
        image_buffer[(i_y_max * y) + x][1] = colors[color][1];
        image_buffer[(i_y_max * y) + x][2] = colors[color][2];
    };
};

void write_to_file(){
    FILE * file;
    char * filename               = "Output/cuda.ppm";
    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,
    global_data->i_x_max, global_data->i_y_max, max_color_component_value);

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

    fclose(file);
};

__global__
void compute_mandelbrot_thread(struct data * gpu_data, int * out){
    
    int i_y_max = gpu_data->i_y_max;
    int n_threads = gpu_data->n_threads * gpu_data->n_blocks; //Numero de threads total
    int ind = blockIdx.x * blockDim.x + threadIdx.x;
    int chunkSize = i_y_max / n_threads;
    int leftover = i_y_max % n_threads;
    int begin = ind * chunkSize + ((ind != 0) ? leftover : 0);
    int end = begin + chunkSize + (ind == 0 ? leftover : 0);
    double z_x;
    double z_y;
    double z_x_squared;
    double z_y_squared;
    double escape_radius_squared = 4;
    int iteration;
    int i_x;
    int i_y;

    double c_x;
    double c_y;

    double c_y_min = gpu_data->c_y_min;
    double c_x_min = gpu_data->c_x_min;
    int i_x_max = gpu_data->i_x_max;
    int iteration_max = 200;
    double pixel_height = gpu_data->pixel_height;
    double pixel_width = gpu_data->pixel_width;

    for(i_y = begin; i_y < end; i_y++){
        c_y = c_y_min + i_y * pixel_height;

        if(fabs(c_y) < pixel_height / 2){
            c_y = 0.0;
        };

        for(i_x = 0; i_x < i_x_max; i_x++){
            c_x         = c_x_min + i_x * pixel_width;

            z_x         = 0.0;
            z_y         = 0.0;

            z_x_squared = 0.0;
            z_y_squared = 0.0;

            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;
            };
            out[ (i_x)*i_y_max + i_y       ] = iteration;
        };
    };

    
};


int main(int argc, char *argv[]){
    int * d_out;
    int * out;
    struct data * gpu_data;

    init(argc, argv);
    allocate_image_buffer();

    cudaSetDevice(0);
    cudaMalloc((void **)&d_out, global_data->i_x_max * global_data->i_y_max * sizeof(int));
    cudaMalloc((void **)&gpu_data, sizeof(struct data));
    cudaMemcpy(gpu_data, global_data, sizeof(struct data), cudaMemcpyHostToDevice);
    
    /*Lançamento da gpu*/
    int n_threads = global_data->n_threads;
    int n_blocks = global_data->n_blocks;

    clock_gettime(CLOCK_MONOTONIC, &timer.t_start);
    compute_mandelbrot_thread<<<n_blocks, n_threads>>>(gpu_data, d_out);
    cudaDeviceSynchronize();
    
    /*Copia dados para cpu*/
    out = (int *) malloc(global_data->i_x_max * global_data->i_y_max * sizeof(int));

    cudaMemcpy(out, d_out, global_data->i_x_max * global_data->i_y_max * sizeof(int), cudaMemcpyDeviceToHost);

    for(int i = 0; i<global_data->i_x_max * global_data->i_y_max; i++){
        int linha = i/global_data->i_y_max;
        int coluna = i%global_data->i_y_max;
        int iteration = out[i];
        update_rgb_buffer(iteration, linha, coluna);
    }

    clock_gettime(CLOCK_MONOTONIC, &timer.t_end);
    printf("%f\n",
               (double) (timer.t_end.tv_sec - timer.t_start.tv_sec) +
               (double) (timer.t_end.tv_nsec - timer.t_start.tv_nsec) / 1000000000.0);
    
    write_to_file();
    cudaFree(d_out);
    cudaFree(gpu_data);

    return 0;
};

	.file	"tmpxft_003ab600_00000000-6_mandelbrot_cuda.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2043:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2043:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z21allocate_image_bufferv
	.type	_Z21allocate_image_bufferv, @function
_Z21allocate_image_bufferv:
.LFB2036:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	movslq	image_buffer_size(%rip), %rdi
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	movq	%rdi, %r12
	salq	$3, %rdi
	xorl	%ebx, %ebx
	call	malloc@PLT
	movq	%rax, image_buffer(%rip)
	movq	%rax, %rbp
.L3:
	cmpl	%ebx, %r12d
	jle	.L7
	movl	$3, %edi
	call	malloc@PLT
	movq	%rax, 0(%rbp,%rbx,8)
	incq	%rbx
	jmp	.L3
.L7:
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2036:
	.size	_Z21allocate_image_bufferv, .-_Z21allocate_image_bufferv
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"usage: ./mandelbrot_cuda c_x_min c_x_max c_y_min c_y_max image_size blocks threads\n"
.LC1:
	.string	"examples with image_size = 11500:\n"
.LC2:
	.string	"    Full Picture:         ./mandelbrot_cuda -2.5 1.5 -2.0 2.0 11500 4 256\n"
.LC3:
	.string	"    Seahorse Valley:      ./mandelbrot_cuda -0.8 -0.7 0.05 0.15 11500 4 256\n"
.LC4:
	.string	"    Elephant Valley:      ./mandelbrot_cuda 0.175 0.375 -0.1 0.1 11500 4 256\n"
.LC5:
	.string	"    Triple Spiral Valley: ./mandelbrot_cuda -0.188 -0.012 0.554 0.754 11500 4 256\n"
.LC6:
	.string	"%lf"
.LC7:
	.string	"%d"
	.text
	.globl	_Z4initiPPc
	.type	_Z4initiPPc, @function
_Z4initiPPc:
.LFB2037:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movl	%edi, %ebp
	movl	$64, %edi
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movq	%rsi, %rbx
	pushq	%rdx
	.cfi_def_cfa_offset 32
	call	malloc@PLT
	movq	%rax, global_data(%rip)
	cmpl	$7, %ebp
	jg	.L9
	leaq	.LC0(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC1(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC2(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC3(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC4(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	$2, %edi
	leaq	.LC5(%rip), %rsi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	xorl	%edi, %edi
	call	exit@PLT
.L9:
	leaq	.LC6(%rip), %rbp
	movq	8(%rbx), %rdi
	movq	%rax, %rdx
	xorl	%eax, %eax
	movq	%rbp, %rsi
	call	__isoc23_sscanf@PLT
	movq	global_data(%rip), %rax
	movq	16(%rbx), %rdi
	movq	%rbp, %rsi
	leaq	8(%rax), %rdx
	xorl	%eax, %eax
	call	__isoc23_sscanf@PLT
	movq	global_data(%rip), %rax
	movq	24(%rbx), %rdi
	movq	%rbp, %rsi
	leaq	16(%rax), %rdx
	xorl	%eax, %eax
	call	__isoc23_sscanf@PLT
	movq	global_data(%rip), %rax
	movq	32(%rbx), %rdi
	movq	%rbp, %rsi
	leaq	.LC7(%rip), %rbp
	leaq	24(%rax), %rdx
	xorl	%eax, %eax
	call	__isoc23_sscanf@PLT
	movq	40(%rbx), %rdi
	movq	%rbp, %rsi
	xorl	%eax, %eax
	leaq	image_size(%rip), %rdx
	call	__isoc23_sscanf@PLT
	movq	global_data(%rip), %rax
	movq	48(%rbx), %rdi
	movq	%rbp, %rsi
	leaq	56(%rax), %rdx
	xorl	%eax, %eax
	call	__isoc23_sscanf@PLT
	movq	global_data(%rip), %rax
	movq	48(%rbx), %rdi
	movq	%rbp, %rsi
	leaq	60(%rax), %rdx
	xorl	%eax, %eax
	call	__isoc23_sscanf@PLT
	movl	image_size(%rip), %edx
	movq	global_data(%rip), %rax
	cvtsi2sdl	%edx, %xmm1
	movsd	8(%rax), %xmm0
	subsd	(%rax), %xmm0
	movl	%edx, %ecx
	imull	%edx, %ecx
	movl	%edx, 48(%rax)
	movl	%edx, 52(%rax)
	divsd	%xmm1, %xmm0
	movl	%ecx, image_buffer_size(%rip)
	movsd	%xmm0, 32(%rax)
	movsd	24(%rax), %xmm0
	subsd	16(%rax), %xmm0
	divsd	%xmm1, %xmm0
	movsd	%xmm0, 40(%rax)
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2037:
	.size	_Z4initiPPc, .-_Z4initiPPc
	.globl	_Z17update_rgb_bufferiii
	.type	_Z17update_rgb_bufferiii, @function
_Z17update_rgb_bufferiii:
.LFB2038:
	.cfi_startproc
	endbr64
	movl	%edi, %eax
	movq	global_data(%rip), %rdi
	movl	%esi, %ecx
	movslq	gradient_size(%rip), %rsi
	imull	52(%rdi), %edx
	leaq	colors(%rip), %rdi
	addl	%ecx, %edx
	cmpl	$200, %eax
	movslq	%edx, %rcx
	movq	image_buffer(%rip), %rdx
	movq	(%rdx,%rcx,8), %r8
	jne	.L12
	imulq	$12, %rsi, %rsi
	movl	(%rdi,%rsi), %eax
	movb	%al, (%r8)
	movq	image_buffer(%rip), %rax
	movq	(%rax,%rcx,8), %rdx
	movslq	gradient_size(%rip), %rax
	imulq	$12, %rax, %rax
	movl	4(%rdi,%rax), %eax
	movb	%al, 1(%rdx)
	movq	image_buffer(%rip), %rax
	movq	(%rax,%rcx,8), %rdx
	movslq	gradient_size(%rip), %rax
	imulq	$12, %rax, %rax
	movl	8(%rdi,%rax), %eax
	jmp	.L14
.L12:
	cltd
	idivl	%esi
	movslq	%edx, %rdx
	imulq	$12, %rdx, %rdx
	leaq	(%rdi,%rdx), %rax
	movl	(%rax), %edx
	movb	%dl, (%r8)
	movq	image_buffer(%rip), %rdx
	movl	4(%rax), %esi
	movq	(%rdx,%rcx,8), %rdx
	movb	%sil, 1(%rdx)
	movq	image_buffer(%rip), %rdx
	movl	8(%rax), %eax
	movq	(%rdx,%rcx,8), %rdx
.L14:
	movb	%al, 2(%rdx)
	ret
	.cfi_endproc
.LFE2038:
	.size	_Z17update_rgb_bufferiii, .-_Z17update_rgb_bufferiii
	.section	.rodata.str1.1
.LC8:
	.string	"wb"
.LC9:
	.string	"Output/cuda.ppm"
.LC10:
	.string	"# "
.LC11:
	.string	"P6\n %s\n %d\n %d\n %d\n"
	.text
	.globl	_Z13write_to_filev
	.type	_Z13write_to_filev, @function
_Z13write_to_filev:
.LFB2039:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	leaq	.LC8(%rip), %rsi
	leaq	.LC9(%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
	leaq	.LC11(%rip), %rdx
	movq	%rax, %rbx
	movq	global_data(%rip), %rax
	movq	%rbx, %rdi
	movl	52(%rax), %r9d
	movl	48(%rax), %r8d
	pushq	%rcx
	.cfi_def_cfa_offset 40
	xorl	%eax, %eax
	pushq	$255
	.cfi_def_cfa_offset 48
	leaq	.LC10(%rip), %rcx
	call	__fprintf_chk@PLT
	popq	%rsi
	.cfi_def_cfa_offset 40
	popq	%rdi
	.cfi_def_cfa_offset 32
.L16:
	cmpl	%ebp, image_buffer_size(%rip)
	jle	.L19
	movq	image_buffer(%rip), %rax
	movq	%rbx, %rcx
	movl	$3, %edx
	movl	$1, %esi
	movq	(%rax,%rbp,8), %rdi
	incq	%rbp
	call	fwrite@PLT
	jmp	.L16
.L19:
	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
.LFE2039:
	.size	_Z13write_to_filev, .-_Z13write_to_filev
	.globl	_Z51__device_stub__Z25compute_mandelbrot_threadP4dataPiP4dataPi
	.type	_Z51__device_stub__Z25compute_mandelbrot_threadP4dataPiP4dataPi, @function
_Z51__device_stub__Z25compute_mandelbrot_threadP4dataPiP4dataPi:
.LFB2065:
	.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	.L20
	pushq	32(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z25compute_mandelbrot_threadP4dataPi(%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
.L20:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L22
	call	__stack_chk_fail@PLT
.L22:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2065:
	.size	_Z51__device_stub__Z25compute_mandelbrot_threadP4dataPiP4dataPi, .-_Z51__device_stub__Z25compute_mandelbrot_threadP4dataPiP4dataPi
	.globl	_Z25compute_mandelbrot_threadP4dataPi
	.type	_Z25compute_mandelbrot_threadP4dataPi, @function
_Z25compute_mandelbrot_threadP4dataPi:
.LFB2066:
	.cfi_startproc
	endbr64
	jmp	_Z51__device_stub__Z25compute_mandelbrot_threadP4dataPiP4dataPi
	.cfi_endproc
.LFE2066:
	.size	_Z25compute_mandelbrot_threadP4dataPi, .-_Z25compute_mandelbrot_threadP4dataPi
	.section	.rodata.str1.1
.LC13:
	.string	"%f\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2040:
	.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
	subq	$56, %rsp
	.cfi_def_cfa_offset 80
	movq	%fs:40, %rax
	movq	%rax, 40(%rsp)
	xorl	%eax, %eax
	call	_Z4initiPPc
	call	_Z21allocate_image_bufferv
	xorl	%edi, %edi
	call	cudaSetDevice@PLT
	movq	global_data(%rip), %rax
	movq	%rsp, %rdi
	movl	48(%rax), %esi
	imull	52(%rax), %esi
	movslq	%esi, %rsi
	salq	$2, %rsi
	call	cudaMalloc@PLT
	leaq	8(%rsp), %rdi
	movl	$64, %esi
	call	cudaMalloc@PLT
	movq	8(%rsp), %rdi
	movl	$1, %ecx
	movq	global_data(%rip), %rsi
	movl	$64, %edx
	call	cudaMemcpy@PLT
	movq	global_data(%rip), %rax
	movl	$1, %edi
	leaq	16+timer(%rip), %rsi
	movl	60(%rax), %ebp
	movl	56(%rax), %ebx
	call	clock_gettime@PLT
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movabsq	$4294967297, %rax
	movq	%rax, 32(%rsp)
	movl	36(%rsp), %ecx
	movq	%rax, 20(%rsp)
	movl	24(%rsp), %esi
	movl	%ebp, 28(%rsp)
	movq	28(%rsp), %rdx
	movl	%ebx, 16(%rsp)
	movq	16(%rsp), %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L26
	movq	(%rsp), %rsi
	movq	8(%rsp), %rdi
	call	_Z51__device_stub__Z25compute_mandelbrot_threadP4dataPiP4dataPi
.L26:
	call	cudaDeviceSynchronize@PLT
	movq	global_data(%rip), %rax
	movl	48(%rax), %ebx
	imull	52(%rax), %ebx
	movslq	%ebx, %rbx
	salq	$2, %rbx
	movq	%rbx, %rdi
	call	malloc@PLT
	movq	(%rsp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdx
	movq	%rax, %rdi
	movq	%rax, %rbp
	call	cudaMemcpy@PLT
	xorl	%r9d, %r9d
.L27:
	movq	global_data(%rip), %rdx
	movl	%r9d, %eax
	movl	52(%rdx), %ecx
	movl	48(%rdx), %ebx
	imull	%ecx, %ebx
	cmpl	%r9d, %ebx
	jle	.L31
	cltd
	movl	0(%rbp,%r9,4), %edi
	incq	%r9
	idivl	%ecx
	movl	%eax, %esi
	call	_Z17update_rgb_bufferiii
	jmp	.L27
.L31:
	leaq	32+timer(%rip), %rsi
	movl	$1, %edi
	call	clock_gettime@PLT
	movq	40+timer(%rip), %rax
	movl	$2, %edi
	subq	24+timer(%rip), %rax
	cvtsi2sdq	%rax, %xmm0
	movq	32+timer(%rip), %rax
	subq	16+timer(%rip), %rax
	divsd	.LC12(%rip), %xmm0
	cvtsi2sdq	%rax, %xmm1
	leaq	.LC13(%rip), %rsi
	movb	$1, %al
	addsd	%xmm1, %xmm0
	call	__printf_chk@PLT
	call	_Z13write_to_filev
	movq	(%rsp), %rdi
	call	cudaFree@PLT
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	movq	40(%rsp), %rax
	subq	%fs:40, %rax
	je	.L29
	call	__stack_chk_fail@PLT
.L29:
	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
.LFE2040:
	.size	main, .-main
	.section	.rodata.str1.1
.LC14:
	.string	"_Z25compute_mandelbrot_threadP4dataPi"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2068:
	.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	.LC14(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z25compute_mandelbrot_threadP4dataPi(%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
.LFE2068:
	.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	timer
	.bss
	.align 32
	.type	timer, @object
	.size	timer, 80
timer:
	.zero	80
	.globl	image_buffer_size
	.align 4
	.type	image_buffer_size, @object
	.size	image_buffer_size, 4
image_buffer_size:
	.zero	4
	.globl	image_size
	.align 4
	.type	image_size, @object
	.size	image_size, 4
image_size:
	.zero	4
	.globl	image_buffer
	.align 8
	.type	image_buffer, @object
	.size	image_buffer, 8
image_buffer:
	.zero	8
	.globl	gradient_size
	.data
	.align 4
	.type	gradient_size, @object
	.size	gradient_size, 4
gradient_size:
	.long	16
	.globl	colors
	.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
	.globl	global_data
	.bss
	.align 8
	.type	global_data, @object
	.size	global_data, 8
global_data:
	.zero	8
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC12:
	.long	0
	.long	1104006501
	.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>
#include <pthread.h>
#include <time.h>
#include <sys/time.h>

/*Dados HOST*/
struct data{
    double c_x_min;
    double c_x_max;
    double c_y_min;
    double c_y_max;

    double pixel_width;
    double pixel_height;

    int i_x_max;
    int i_y_max;

    int n_blocks;
    int n_threads;
};

struct data * global_data;

int colors[17][3] = {
    {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 gradient_size = 16;
unsigned char **image_buffer;
int image_size;
int image_buffer_size;
//////////////////////////////////////////


/*Timer info*/
struct timer_info {
    clock_t c_start;
    clock_t c_end;
    struct timespec t_start;
    struct timespec t_end;
    struct timeval v_start;
    struct timeval v_end;
};

struct timer_info timer;
//////////////////////////////////////////

void allocate_image_buffer(){
    int rgb_size = 3;
    image_buffer = (unsigned char **) malloc(sizeof(unsigned char *) * image_buffer_size);
    for(int i = 0; i < image_buffer_size; i++){
        image_buffer[i] = (unsigned char *) malloc(sizeof(unsigned char) * rgb_size);
    };
};

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

    global_data = (struct data * )malloc(sizeof(struct data));

    if(argc < 8){
        printf("usage: ./mandelbrot_cuda c_x_min c_x_max c_y_min c_y_max image_size blocks threads\n");
        printf("examples with image_size = 11500:\n");
        printf("    Full Picture:         ./mandelbrot_cuda -2.5 1.5 -2.0 2.0 11500 4 256\n");
        printf("    Seahorse Valley:      ./mandelbrot_cuda -0.8 -0.7 0.05 0.15 11500 4 256\n");
        printf("    Elephant Valley:      ./mandelbrot_cuda 0.175 0.375 -0.1 0.1 11500 4 256\n");
        printf("    Triple Spiral Valley: ./mandelbrot_cuda -0.188 -0.012 0.554 0.754 11500 4 256\n");
        exit(0);
    }
    else{
        sscanf(argv[1], "%lf", &global_data->c_x_min);
        sscanf(argv[2], "%lf", &global_data->c_x_max);
        sscanf(argv[3], "%lf", &global_data->c_y_min);
        sscanf(argv[4], "%lf", &global_data->c_y_max);
        sscanf(argv[5], "%d", &image_size);
        sscanf(argv[6], "%d", &global_data->n_blocks);
        sscanf(argv[6], "%d", &global_data->n_threads);

        global_data->i_x_max           = image_size;
        global_data->i_y_max           = image_size;
        image_buffer_size = image_size * image_size;

        global_data->pixel_width       = (global_data->c_x_max - global_data->c_x_min) / global_data->i_x_max;
        global_data->pixel_height      = (global_data->c_y_max - global_data->c_y_min) / global_data->i_y_max;
    };

};

void update_rgb_buffer(int iteration, int x, int y){
    int color;

    int i_y_max = global_data->i_y_max;
    int iteration_max = 200;

    if(iteration == iteration_max){
        image_buffer[(i_y_max * y) + x][0] = colors[gradient_size][0];
        image_buffer[(i_y_max * y) + x][1] = colors[gradient_size][1];
        image_buffer[(i_y_max * y) + x][2] = colors[gradient_size][2];
    }
    else{
        color = iteration % gradient_size;

        image_buffer[(i_y_max * y) + x][0] = colors[color][0];
        image_buffer[(i_y_max * y) + x][1] = colors[color][1];
        image_buffer[(i_y_max * y) + x][2] = colors[color][2];
    };
};

void write_to_file(){
    FILE * file;
    char * filename               = "Output/cuda.ppm";
    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,
    global_data->i_x_max, global_data->i_y_max, max_color_component_value);

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

    fclose(file);
};

__global__
void compute_mandelbrot_thread(struct data * gpu_data, int * out){
    
    int i_y_max = gpu_data->i_y_max;
    int n_threads = gpu_data->n_threads * gpu_data->n_blocks; //Numero de threads total
    int ind = blockIdx.x * blockDim.x + threadIdx.x;
    int chunkSize = i_y_max / n_threads;
    int leftover = i_y_max % n_threads;
    int begin = ind * chunkSize + ((ind != 0) ? leftover : 0);
    int end = begin + chunkSize + (ind == 0 ? leftover : 0);
    double z_x;
    double z_y;
    double z_x_squared;
    double z_y_squared;
    double escape_radius_squared = 4;
    int iteration;
    int i_x;
    int i_y;

    double c_x;
    double c_y;

    double c_y_min = gpu_data->c_y_min;
    double c_x_min = gpu_data->c_x_min;
    int i_x_max = gpu_data->i_x_max;
    int iteration_max = 200;
    double pixel_height = gpu_data->pixel_height;
    double pixel_width = gpu_data->pixel_width;

    for(i_y = begin; i_y < end; i_y++){
        c_y = c_y_min + i_y * pixel_height;

        if(fabs(c_y) < pixel_height / 2){
            c_y = 0.0;
        };

        for(i_x = 0; i_x < i_x_max; i_x++){
            c_x         = c_x_min + i_x * pixel_width;

            z_x         = 0.0;
            z_y         = 0.0;

            z_x_squared = 0.0;
            z_y_squared = 0.0;

            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;
            };
            out[ (i_x)*i_y_max + i_y       ] = iteration;
        };
    };

    
};


int main(int argc, char *argv[]){
    int * d_out;
    int * out;
    struct data * gpu_data;

    init(argc, argv);
    allocate_image_buffer();

    hipSetDevice(0);
    hipMalloc((void **)&d_out, global_data->i_x_max * global_data->i_y_max * sizeof(int));
    hipMalloc((void **)&gpu_data, sizeof(struct data));
    hipMemcpy(gpu_data, global_data, sizeof(struct data), hipMemcpyHostToDevice);
    
    /*Lançamento da gpu*/
    int n_threads = global_data->n_threads;
    int n_blocks = global_data->n_blocks;

    clock_gettime(CLOCK_MONOTONIC, &timer.t_start);
    compute_mandelbrot_thread<<<n_blocks, n_threads>>>(gpu_data, d_out);
    hipDeviceSynchronize();
    
    /*Copia dados para cpu*/
    out = (int *) malloc(global_data->i_x_max * global_data->i_y_max * sizeof(int));

    hipMemcpy(out, d_out, global_data->i_x_max * global_data->i_y_max * sizeof(int), hipMemcpyDeviceToHost);

    for(int i = 0; i<global_data->i_x_max * global_data->i_y_max; i++){
        int linha = i/global_data->i_y_max;
        int coluna = i%global_data->i_y_max;
        int iteration = out[i];
        update_rgb_buffer(iteration, linha, coluna);
    }

    clock_gettime(CLOCK_MONOTONIC, &timer.t_end);
    printf("%f\n",
               (double) (timer.t_end.tv_sec - timer.t_start.tv_sec) +
               (double) (timer.t_end.tv_nsec - timer.t_start.tv_nsec) / 1000000000.0);
    
    write_to_file();
    hipFree(d_out);
    hipFree(gpu_data);

    return 0;
};

	.text
	.file	"mandelbrot_cuda.hip"
	.globl	_Z21allocate_image_bufferv      # -- Begin function _Z21allocate_image_bufferv
	.type	_Z21allocate_image_bufferv,@function
_Z21allocate_image_bufferv:             # @_Z21allocate_image_bufferv
	.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
	movslq	image_buffer_size(%rip), %rbx
	leaq	(,%rbx,8), %rdi
	callq	malloc
	movq	%rax, image_buffer(%rip)
	testq	%rbx, %rbx
	jle	.LBB0_3
# %bb.1:                                # %.lr.ph.preheader
	movl	%ebx, %ebx
	xorl	%r14d, %r14d
.LBB0_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movl	$3, %edi
	callq	malloc
	movq	image_buffer(%rip), %rcx
	movq	%rax, (%rcx,%r14,8)
	incq	%r14
	cmpq	%r14, %rbx
	jne	.LBB0_2
.LBB0_3:                                # %._crit_edge
	addq	$8, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z21allocate_image_bufferv, .Lfunc_end0-_Z21allocate_image_bufferv
	.cfi_endproc
                                        # -- End function
	.globl	_Z4initiPPc                     # -- Begin function _Z4initiPPc
	.type	_Z4initiPPc,@function
_Z4initiPPc:                            # @_Z4initiPPc
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.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 %rbp, -16
	movq	%rsi, %rbx
	movl	%edi, %ebp
	movl	$64, %edi
	callq	malloc
	movq	%rax, global_data(%rip)
	cmpl	$7, %ebp
	jle	.LBB1_2
# %bb.1:
	movq	8(%rbx), %rdi
	movl	$.L.str.6, %esi
	movq	%rax, %rdx
	xorl	%eax, %eax
	callq	__isoc23_sscanf
	movq	16(%rbx), %rdi
	movq	global_data(%rip), %rdx
	addq	$8, %rdx
	movl	$.L.str.6, %esi
	xorl	%eax, %eax
	callq	__isoc23_sscanf
	movq	24(%rbx), %rdi
	movq	global_data(%rip), %rdx
	addq	$16, %rdx
	movl	$.L.str.6, %esi
	xorl	%eax, %eax
	callq	__isoc23_sscanf
	movq	32(%rbx), %rdi
	movq	global_data(%rip), %rdx
	addq	$24, %rdx
	movl	$.L.str.6, %esi
	xorl	%eax, %eax
	callq	__isoc23_sscanf
	movq	40(%rbx), %rdi
	movl	$.L.str.7, %esi
	movl	$image_size, %edx
	xorl	%eax, %eax
	callq	__isoc23_sscanf
	movq	48(%rbx), %rdi
	movq	global_data(%rip), %rdx
	addq	$56, %rdx
	movl	$.L.str.7, %esi
	xorl	%eax, %eax
	callq	__isoc23_sscanf
	movq	48(%rbx), %rdi
	movq	global_data(%rip), %rdx
	addq	$60, %rdx
	movl	$.L.str.7, %esi
	xorl	%eax, %eax
	callq	__isoc23_sscanf
	movl	image_size(%rip), %eax
	movq	global_data(%rip), %rcx
	movl	%eax, 48(%rcx)
	movl	%eax, 52(%rcx)
	cvtsi2sd	%eax, %xmm0
	imull	%eax, %eax
	movl	%eax, image_buffer_size(%rip)
	movsd	8(%rcx), %xmm1                  # xmm1 = mem[0],zero
	movsd	24(%rcx), %xmm2                 # xmm2 = mem[0],zero
	subsd	(%rcx), %xmm1
	divsd	%xmm0, %xmm1
	movsd	%xmm1, 32(%rcx)
	subsd	16(%rcx), %xmm2
	divsd	%xmm0, %xmm2
	movsd	%xmm2, 40(%rcx)
	addq	$8, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.LBB1_2:
	.cfi_def_cfa_offset 32
	movl	$.Lstr, %edi
	callq	puts@PLT
	movl	$.Lstr.1, %edi
	callq	puts@PLT
	movl	$.Lstr.2, %edi
	callq	puts@PLT
	movl	$.Lstr.3, %edi
	callq	puts@PLT
	movl	$.Lstr.4, %edi
	callq	puts@PLT
	movl	$.Lstr.5, %edi
	callq	puts@PLT
	xorl	%edi, %edi
	callq	exit
.Lfunc_end1:
	.size	_Z4initiPPc, .Lfunc_end1-_Z4initiPPc
	.cfi_endproc
                                        # -- End function
	.globl	_Z17update_rgb_bufferiii        # -- Begin function _Z17update_rgb_bufferiii
	.type	_Z17update_rgb_bufferiii,@function
_Z17update_rgb_bufferiii:               # @_Z17update_rgb_bufferiii
	.cfi_startproc
# %bb.0:
	movl	%edx, %ecx
	movq	global_data(%rip), %rax
	movl	52(%rax), %r8d
	movslq	gradient_size(%rip), %r9
	cmpl	$200, %edi
	jne	.LBB2_2
# %bb.1:
	leaq	(%r9,%r9,2), %rax
	movb	colors(,%rax,4), %dl
	movq	image_buffer(%rip), %rdi
	imull	%ecx, %r8d
	addl	%esi, %r8d
	movslq	%r8d, %rax
	movq	(%rdi,%rax,8), %rcx
	movb	%dl, (%rcx)
	movslq	gradient_size(%rip), %rcx
	leaq	(%rcx,%rcx,2), %rcx
	movb	colors+4(,%rcx,4), %cl
	movq	image_buffer(%rip), %rdx
	movq	(%rdx,%rax,8), %rdx
	movb	%cl, 1(%rdx)
	movslq	gradient_size(%rip), %rdx
	jmp	.LBB2_3
.LBB2_2:
	movl	%edi, %eax
	cltd
	idivl	%r9d
	movslq	%edx, %rdx
	leaq	(,%rdx,4), %rdi
	movb	colors(%rdi,%rdi,2), %r9b
	movq	image_buffer(%rip), %r10
	imull	%ecx, %r8d
	addl	%esi, %r8d
	movslq	%r8d, %rax
	movq	(%r10,%rax,8), %rcx
	movb	%r9b, (%rcx)
	movb	colors+4(%rdi,%rdi,2), %cl
	movq	image_buffer(%rip), %rsi
	movq	(%rsi,%rax,8), %rsi
	movb	%cl, 1(%rsi)
.LBB2_3:
	leaq	(%rdx,%rdx,2), %rcx
	movb	colors+8(,%rcx,4), %cl
	movq	image_buffer(%rip), %rdx
	movq	(%rdx,%rax,8), %rax
	movb	%cl, 2(%rax)
	retq
.Lfunc_end2:
	.size	_Z17update_rgb_bufferiii, .Lfunc_end2-_Z17update_rgb_bufferiii
	.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.8, %edi
	movl	$.L.str.10, %esi
	callq	fopen
	movq	%rax, %rbx
	movq	global_data(%rip), %rax
	movl	48(%rax), %ecx
	movl	52(%rax), %r8d
	movl	$.L.str.11, %esi
	movl	$.L.str.9, %edx
	movq	%rbx, %rdi
	movl	$255, %r9d
	xorl	%eax, %eax
	callq	fprintf
	cmpl	$0, image_buffer_size(%rip)
	jle	.LBB3_3
# %bb.1:                                # %.lr.ph.preheader
	xorl	%r14d, %r14d
.LBB3_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movq	image_buffer(%rip), %rax
	movq	(%rax,%r14,8), %rdi
	movl	$1, %esi
	movl	$3, %edx
	movq	%rbx, %rcx
	callq	fwrite
	incq	%r14
	movslq	image_buffer_size(%rip), %rax
	cmpq	%rax, %r14
	jl	.LBB3_2
.LBB3_3:                                # %._crit_edge
	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_end3:
	.size	_Z13write_to_filev, .Lfunc_end3-_Z13write_to_filev
	.cfi_endproc
                                        # -- End function
	.globl	_Z40__device_stub__compute_mandelbrot_threadP4dataPi # -- Begin function _Z40__device_stub__compute_mandelbrot_threadP4dataPi
	.type	_Z40__device_stub__compute_mandelbrot_threadP4dataPi,@function
_Z40__device_stub__compute_mandelbrot_threadP4dataPi: # @_Z40__device_stub__compute_mandelbrot_threadP4dataPi
	.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	$_Z25compute_mandelbrot_threadP4dataPi, %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_end4:
	.size	_Z40__device_stub__compute_mandelbrot_threadP4dataPi, .Lfunc_end4-_Z40__device_stub__compute_mandelbrot_threadP4dataPi
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function main
.LCPI5_0:
	.quad	0x41cdcd6500000000              # double 1.0E+9
	.text
	.globl	main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%r14
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	subq	$24, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -24
	.cfi_offset %r14, -16
	callq	_Z4initiPPc
	callq	_Z21allocate_image_bufferv
	xorl	%edi, %edi
	callq	hipSetDevice
	movq	global_data(%rip), %rax
	movslq	48(%rax), %rcx
	movslq	52(%rax), %rsi
	imulq	%rcx, %rsi
	shlq	$2, %rsi
	leaq	8(%rsp), %rdi
	callq	hipMalloc
	leaq	16(%rsp), %rbx
	movl	$64, %esi
	movq	%rbx, %rdi
	callq	hipMalloc
	movq	(%rbx), %rdi
	movq	global_data(%rip), %rsi
	movl	$64, %edx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	global_data(%rip), %rax
	movl	56(%rax), %ebx
	movl	60(%rax), %r14d
	movl	$timer+16, %esi
	movl	$1, %edi
	callq	clock_gettime
	btsq	$32, %rbx
	btsq	$32, %r14
	movq	%rbx, %rdi
	movl	$1, %esi
	movq	%r14, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB5_2
# %bb.1:
	movq	16(%rsp), %rdi
	movq	8(%rsp), %rsi
	callq	_Z40__device_stub__compute_mandelbrot_threadP4dataPi
.LBB5_2:
	callq	hipDeviceSynchronize
	movq	global_data(%rip), %rax
	movslq	48(%rax), %rcx
	movslq	52(%rax), %r14
	imulq	%rcx, %r14
	shlq	$2, %r14
	movq	%r14, %rdi
	callq	malloc
	movq	%rax, %rbx
	movq	8(%rsp), %rsi
	movq	%rax, %rdi
	movq	%r14, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movq	global_data(%rip), %rax
	movl	52(%rax), %ecx
	movl	48(%rax), %eax
	imull	%ecx, %eax
	testl	%eax, %eax
	jle	.LBB5_5
# %bb.3:                                # %.lr.ph.preheader
	xorl	%r14d, %r14d
.LBB5_4:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movl	%r14d, %eax
	cltd
	idivl	%ecx
	movl	(%rbx,%r14,4), %edi
	movl	%eax, %esi
	callq	_Z17update_rgb_bufferiii
	incq	%r14
	movq	global_data(%rip), %rax
	movslq	48(%rax), %rdx
	movslq	52(%rax), %rcx
	imulq	%rcx, %rdx
	cmpq	%rdx, %r14
	jl	.LBB5_4
.LBB5_5:                                # %._crit_edge
	movl	$timer+32, %esi
	movl	$1, %edi
	callq	clock_gettime
	movq	timer+32(%rip), %rax
	subq	timer+16(%rip), %rax
	cvtsi2sd	%rax, %xmm1
	movq	timer+40(%rip), %rax
	subq	timer+24(%rip), %rax
	cvtsi2sd	%rax, %xmm0
	divsd	.LCPI5_0(%rip), %xmm0
	addsd	%xmm1, %xmm0
	movl	$.L.str.12, %edi
	movb	$1, %al
	callq	printf
	callq	_Z13write_to_filev
	movq	8(%rsp), %rdi
	callq	hipFree
	movq	16(%rsp), %rdi
	callq	hipFree
	xorl	%eax, %eax
	addq	$24, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end5:
	.size	main, .Lfunc_end5-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	.LBB6_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB6_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z25compute_mandelbrot_threadP4dataPi, %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_end6:
	.size	__hip_module_ctor, .Lfunc_end6-__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	.LBB7_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
.LBB7_2:
	retq
.Lfunc_end7:
	.size	__hip_module_dtor, .Lfunc_end7-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	global_data,@object             # @global_data
	.bss
	.globl	global_data
	.p2align	3, 0x0
global_data:
	.quad	0
	.size	global_data, 8

	.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	gradient_size,@object           # @gradient_size
	.globl	gradient_size
	.p2align	2, 0x0
gradient_size:
	.long	16                              # 0x10
	.size	gradient_size, 4

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

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

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

	.type	timer,@object                   # @timer
	.globl	timer
	.p2align	3, 0x0
timer:
	.zero	80
	.size	timer, 80

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

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

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	"Output/cuda.ppm"
	.size	.L.str.8, 16

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

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

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

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

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

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z25compute_mandelbrot_threadP4dataPi"
	.size	.L__unnamed_1, 38

	.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

	.type	.Lstr,@object                   # @str
	.section	.rodata.str1.1,"aMS",@progbits,1
.Lstr:
	.asciz	"usage: ./mandelbrot_cuda c_x_min c_x_max c_y_min c_y_max image_size blocks threads"
	.size	.Lstr, 83

	.type	.Lstr.1,@object                 # @str.1
.Lstr.1:
	.asciz	"examples with image_size = 11500:"
	.size	.Lstr.1, 34

	.type	.Lstr.2,@object                 # @str.2
.Lstr.2:
	.asciz	"    Full Picture:         ./mandelbrot_cuda -2.5 1.5 -2.0 2.0 11500 4 256"
	.size	.Lstr.2, 74

	.type	.Lstr.3,@object                 # @str.3
.Lstr.3:
	.asciz	"    Seahorse Valley:      ./mandelbrot_cuda -0.8 -0.7 0.05 0.15 11500 4 256"
	.size	.Lstr.3, 76

	.type	.Lstr.4,@object                 # @str.4
.Lstr.4:
	.asciz	"    Elephant Valley:      ./mandelbrot_cuda 0.175 0.375 -0.1 0.1 11500 4 256"
	.size	.Lstr.4, 77

	.type	.Lstr.5,@object                 # @str.5
.Lstr.5:
	.asciz	"    Triple Spiral Valley: ./mandelbrot_cuda -0.188 -0.012 0.554 0.754 11500 4 256"
	.size	.Lstr.5, 82

	.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 _Z40__device_stub__compute_mandelbrot_threadP4dataPi
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym image_size
	.addrsig_sym timer
	.addrsig_sym _Z25compute_mandelbrot_threadP4dataPi
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z25compute_mandelbrot_threadP4dataPi ; -- Begin function _Z25compute_mandelbrot_threadP4dataPi
	.globl	_Z25compute_mandelbrot_threadP4dataPi
	.p2align	8
	.type	_Z25compute_mandelbrot_threadP4dataPi,@function
_Z25compute_mandelbrot_threadP4dataPi:  ; @_Z25compute_mandelbrot_threadP4dataPi
; %bb.0:
	s_clause 0x1
	s_load_b128 s[4:7], s[0:1], 0x0
	s_load_b32 s8, s[0:1], 0x1c
	s_waitcnt lgkmcnt(0)
	s_load_b128 s[0:3], s[4:5], 0x30
	s_and_b32 s8, s8, 0xffff
	s_waitcnt lgkmcnt(0)
	s_mul_i32 s2, s2, s3
	s_ashr_i32 s12, s1, 31
	s_ashr_i32 s3, s2, 31
	s_add_i32 s13, s1, s12
	s_add_i32 s9, s2, s3
	s_xor_b32 s13, s13, s12
	s_xor_b32 s9, s9, s3
	s_xor_b32 s3, s12, s3
	v_cvt_f32_u32_e32 v1, s9
	s_sub_i32 s11, 0, s9
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_iflag_f32_e32 v1, v1
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v1, 0x4f7ffffe, v1
	v_cvt_u32_f32_e32 v1, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_readfirstlane_b32 s10, v1
	v_mad_u64_u32 v[2:3], null, s15, s8, v[0:1]
	s_mul_i32 s11, s11, s10
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_mul_hi_u32 s11, s10, s11
	v_cmp_eq_u32_e32 vcc_lo, 0, v2
	s_add_i32 s10, s10, s11
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_mul_hi_u32 s10, s13, s10
	s_mul_i32 s11, s10, s9
	s_add_i32 s12, s10, 1
	s_sub_i32 s11, s13, s11
	s_delay_alu instid0(SALU_CYCLE_1)
	s_sub_i32 s13, s11, s9
	s_cmp_ge_u32 s11, s9
	s_cselect_b32 s10, s12, s10
	s_cselect_b32 s11, s13, s11
	s_add_i32 s12, s10, 1
	s_cmp_ge_u32 s11, s9
	s_cselect_b32 s9, s12, s10
	s_mov_b32 s12, 0
	s_xor_b32 s8, s9, s3
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_sub_i32 s3, s8, s3
	s_mul_i32 s2, s3, s2
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_sub_i32 s2, s1, s2
	v_cndmask_b32_e64 v3, s2, 0, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[0:1], null, v2, s3, v[3:4]
	v_cndmask_b32_e64 v1, 0, s2, vcc_lo
	s_mov_b32 s2, exec_lo
	v_add3_u32 v19, v1, s3, v0
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_lt_i32_e64 v0, v19
	s_cbranch_execz .LBB0_8
; %bb.1:                                ; %.lr.ph73
	s_clause 0x2
	s_load_b128 s[8:11], s[4:5], 0x20
	s_load_b64 s[2:3], s[4:5], 0x0
	s_load_b64 s[4:5], s[4:5], 0x10
	s_cmp_gt_i32 s0, 0
	s_cselect_b32 s13, -1, 0
	s_waitcnt lgkmcnt(0)
	v_mul_f64 v[1:2], s[10:11], 0.5
.LBB0_2:                                ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_4 Depth 2
                                        ;       Child Loop BB0_5 Depth 3
	s_and_not1_b32 vcc_lo, exec_lo, s13
	s_cbranch_vccnz .LBB0_7
; %bb.3:                                ; %.lr.ph.preheader
                                        ;   in Loop: Header=BB0_2 Depth=1
	v_cvt_f64_i32_e32 v[3:4], v0
	s_mov_b32 s14, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[5:6], s[10:11], v[3:4], s[4:5]
	v_cmp_nlt_f64_e64 vcc_lo, |v[5:6]|, v[1:2]
	v_mov_b32_e32 v3, 0
	v_dual_mov_b32 v4, 0 :: v_dual_cndmask_b32 v5, 0, v5
	v_cndmask_b32_e32 v6, 0, v6, vcc_lo
	s_delay_alu instid0(VALU_DEP_2)
	v_dual_mov_b32 v8, v4 :: v_dual_mov_b32 v7, v3
.LBB0_4:                                ; %.lr.ph
                                        ;   Parent Loop BB0_2 Depth=1
                                        ; =>  This Loop Header: Depth=2
                                        ;       Child Loop BB0_5 Depth 3
	s_delay_alu instid0(VALU_DEP_1)
	v_fma_f64 v[9:10], s[8:9], v[7:8], s[2:3]
	v_dual_mov_b32 v12, v4 :: v_dual_mov_b32 v11, v3
	v_dual_mov_b32 v14, v4 :: v_dual_mov_b32 v13, v3
	v_dual_mov_b32 v16, v4 :: v_dual_mov_b32 v15, v3
	v_dual_mov_b32 v18, v4 :: v_dual_mov_b32 v17, v3
	s_mov_b32 s15, 0
	s_mov_b32 s16, 0
.LBB0_5:                                ;   Parent Loop BB0_2 Depth=1
                                        ;     Parent Loop BB0_4 Depth=2
                                        ; =>    This Inner Loop Header: Depth=3
	s_delay_alu instid0(VALU_DEP_1)
	v_add_f64 v[15:16], v[15:16], -v[17:18]
	v_add_f64 v[17:18], v[11:12], v[11:12]
	s_add_i32 s17, s16, 1
	s_cmpk_gt_u32 s16, 0xc6
	v_mov_b32_e32 v20, s17
	s_cselect_b32 s16, -1, 0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f64 v[11:12], v[9:10], v[15:16]
	v_fma_f64 v[13:14], v[13:14], v[17:18], v[5:6]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[15:16], v[11:12], v[11:12]
	v_fma_f64 v[17:18], v[13:14], v[13:14], v[15:16]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(SALU_CYCLE_1)
	v_cmp_ngt_f64_e32 vcc_lo, 4.0, v[17:18]
	v_mul_f64 v[17:18], v[13:14], v[13:14]
	s_or_b32 s16, s16, vcc_lo
	s_and_b32 s16, exec_lo, s16
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 s15, s16, s15
	s_mov_b32 s16, s17
	s_and_not1_b32 exec_lo, exec_lo, s15
	s_cbranch_execnz .LBB0_5
; %bb.6:                                ; %.critedge
                                        ;   in Loop: Header=BB0_4 Depth=2
	s_or_b32 exec_lo, exec_lo, s15
	v_add_f64 v[7:8], v[7:8], 1.0
	v_mad_u64_u32 v[9:10], null, s14, s1, v[0:1]
	s_add_i32 s14, s14, 1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_cmp_lg_u32 s14, s0
	v_ashrrev_i32_e32 v10, 31, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[9:10], 2, v[9:10]
	v_add_co_u32 v9, vcc_lo, s6, v9
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v10, vcc_lo, s7, v10, vcc_lo
	global_store_b32 v[9:10], v20, off
	s_cbranch_scc1 .LBB0_4
.LBB0_7:                                ; %._crit_edge
                                        ;   in Loop: Header=BB0_2 Depth=1
	v_add_nc_u32_e32 v0, 1, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_cmp_eq_u32_e32 vcc_lo, v0, v19
	s_or_b32 s12, vcc_lo, s12
	s_and_not1_b32 exec_lo, exec_lo, s12
	s_cbranch_execnz .LBB0_2
.LBB0_8:                                ; %._crit_edge74
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z25compute_mandelbrot_threadP4dataPi
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 272
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 21
		.amdhsa_next_free_sgpr 18
		.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	_Z25compute_mandelbrot_threadP4dataPi, .Lfunc_end0-_Z25compute_mandelbrot_threadP4dataPi
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 676
; NumSgprs: 20
; NumVgprs: 21
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 2
; NumSGPRsForWavesPerEU: 20
; NumVGPRsForWavesPerEU: 21
; Occupancy: 16
; WaveLimiterHint : 1
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.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:           _Z25compute_mandelbrot_threadP4dataPi
    .private_segment_fixed_size: 0
    .sgpr_count:     20
    .sgpr_spill_count: 0
    .symbol:         _Z25compute_mandelbrot_threadP4dataPi.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     21
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
