#include <stdio.h>
#include <stdlib.h>
#include <stdint.h>
#include <string.h>
#include <unistd.h>
#include <sys/time.h>

#define WIDTH 1024
#define HEIGHT 768
#define CORE_COUNT 128

static int colors[] =
{
  0xff0000,  // f
  0xee3300,  // e
  0xcc5500,  // d
  0xaa5500,  // c
  0xaa3300,  // b
  0x666600,  // a
  0x999900,  // 9
  0x669900,  // 8
  0x339900,  // 7
  0x0099aa,  // 6
  0x0066aa,  // 5
  0x0033aa,  // 4
  0x0000aa,  // 3
  0x000099,  // 2
  0x000066,  // 1
  0x000000,  // 0
};

struct MandelInfo
{
  float r_step;
  float i_step;
  float real_start;
  float imaginary_start;
  int width;
  int height;
};

struct CoreInfo
{
  MandelInfo mandel_info;
  int *picture;
  uint8_t signal_start;
  uint8_t signal_done;
  int core_id;
};

int mandel_calc(
  int *picture,
  int width,
  int height,
  float real_start,
  float real_end,
  float imaginary_start,
  float imaginary_end)
{
  const int max_count = 127;
  int x, y;
  float r, i, r_step, i_step;
  float tr, ti, zr, zi;
  int ptr, count;

  r_step = (real_end - real_start) / (float)width;
  i_step = (imaginary_end - imaginary_start) / (float)height;
  ptr = 0;

  //printf("step = %f %f\n", r_step, i_step);

  i = imaginary_start;

  for (y = 0; y < height; y++)
  {
    r = real_start;

    for (x = 0; x < width; x++)
    {
      zr = 0;
      zi = 0;

      for (count = 0; count < max_count; count++)
      {
        tr = ((zr * zr) - (zi * zi));
        ti = (2 * zr * zi);
        zr = tr + r;
        zi = ti + i;
        if ((zr * zr) + (zi * zi) > 4) break;
      }

      picture[ptr] = colors[count >> 3];

      ptr++;
      r = r + r_step;
    }

    i = i + i_step;
  }

  return 0;
}

__global__
void mandel_calc_cuda_single(
  int *picture,
  int width,
  int height,
  float real_start,
  float real_end,
  float imaginary_start,
  float imaginary_end)
{
  const int max_count = 127;
  int x,y;
  float r, i, r_step, i_step;
  float tr, ti, zr, zi;
  int ptr, count;

  int colors[] =
  {
    0xff0000,  // f
    0xee3300,  // e
    0xcc5500,  // d
    0xaa5500,  // c
    0xaa3300,  // b
    0x666600,  // a
    0x999900,  // 9
    0x669900,  // 8
    0x339900,  // 7
    0x0099aa,  // 6
    0x0066aa,  // 5
    0x0033aa,  // 4
    0x0000aa,  // 3
    0x000099,  // 2
    0x000066,  // 1
    0x000000,  // 0
  };

  r_step = (real_end - real_start) / (float)width;
  i_step = (imaginary_end - imaginary_start) / (float)height;
  ptr = 0;

//printf("step = %f %f\n", r_step, i_step);

  i = imaginary_start;

  for (y = 0; y < height; y++)
  {
    r = real_start;

    for (x = 0; x < width; x++)
    {
      zr = 0;
      zi = 0;

      for (count = 0; count < max_count; count++)
      {
        tr = ((zr * zr) - (zi * zi));
        ti = (2 * zr * zi);
        zr = tr + r;
        zi = ti + i;
        if ((zr * zr) + (zi * zi) > 4) break;
      }

      picture[ptr] = colors[count >> 3];

      ptr++;
      r = r + r_step;
    }

    i = i + i_step;
  }
}

__global__
void mandel_calc_cuda_multi(
  int *picture,
  int width,
  int height,
  float real_start,
  float real_end,
  float imaginary_start,
  float imaginary_end)
{
  const int max_count = 127;
  int x,y;
  float r, i, r_step, i_step;
  float tr, ti, zr, zi;
  int ptr, count;

  int index = threadIdx.x;
  //int stride = blockDim.x;

//printf("index=%d stride=%d\n", index, stride);

  int colors[] =
  {
    0xff0000,  // f
    0xee3300,  // e
    0xcc5500,  // d
    0xaa5500,  // c
    0xaa3300,  // b
    0x666600,  // a
    0x999900,  // 9
    0x669900,  // 8
    0x339900,  // 7
    0x0099aa,  // 6
    0x0066aa,  // 5
    0x0033aa,  // 4
    0x0000aa,  // 3
    0x000099,  // 2
    0x000066,  // 1
    0x000000,  // 0
  };

  r_step = (real_end - real_start) / (float)width;
  i_step = (imaginary_end - imaginary_start) / (float)height;

//printf("step = %f %f\n", r_step, i_step);

  height = height / CORE_COUNT;

  y = index * height;

  ptr = y * width;

  i = imaginary_start + (i_step * y);
#if 0
  i = imaginary_start;

  for (y = 0; y < index * height; y++)
  {
    i = i + i_step;
  }
#endif

//printf("ptr=%d index=%d height=%d %f\n", ptr, index, height, i);

  for (y = 0; y < height; y++)
  {
    r = real_start;

    for (x = 0; x < width; x++)
    {
      zr = 0;
      zi = 0;

      for (count = 0; count < max_count; count++)
      {
        tr = ((zr * zr) - (zi * zi));
        ti = (2 * zr * zi);
        zr = tr + r;
        zi = ti + i;
        if ((zr * zr) + (zi * zi) > 4) break;
      }

      picture[ptr] = colors[count >> 3];

      ptr++;
      r = r + r_step;
    }

    i = i + i_step;
  }
}

int write_int32(FILE *out, int n)
{
  putc((n & 0xff), out);
  putc(((n >> 8) & 0xff), out);
  putc(((n >> 16) & 0xff), out);
  putc(((n >> 24) & 0xff), out);

  return 0;
}

int write_int16(FILE *out, int n)
{
  putc((n & 0xff), out);
  putc(((n >> 8) & 0xff), out);

  return 0;
}

void write_bmp(int *picture, int width, int height)
{
  FILE *out;
  int bmp_width;
  int bmp_size;
  int padding;
  int offset;
  int color;
  int x,y;

  out = fopen("out.bmp", "wb");
  if (out == NULL)
  {
    printf("Can't open file for writing.");
    return;
  }

  bmp_width = width * 3;
  bmp_width = (bmp_width + 3) & (~0x3);
  bmp_size = (bmp_width * height) + 14 + 40;
  padding = bmp_width - (width * 3);

  //printf("width=%d (%d)\n", width, width*3);
  //printf("bmp_width=%d\n", bmp_width);
  //printf("bmp_size=%d\n", bmp_size);

  /* size: 14 bytes */

  putc('B', out);
  putc('M', out);
  write_int32(out, bmp_size);
  write_int16(out, 0);
  write_int16(out, 0);
  write_int32(out, 54);

  /* head1: 14  head2: 40 */

  write_int32(out, 40);         /* biSize */
  write_int32(out, width);
  write_int32(out, height);
  write_int16(out, 1);
  write_int16(out, 24);
  write_int32(out, 0);          /* compression */
  write_int32(out, bmp_width*height);
  write_int32(out, 0);          /* biXPelsperMetre */
  write_int32(out, 0);          /* biYPelsperMetre */
  write_int32(out, 0);
  write_int32(out, 0);

  for (y = 0; y < height; y++)
  {
    offset = y * width;

    for (x = 0; x < width; x++)
    {
      color = picture[offset++];

      putc(color & 0xff, out);
      putc((color >> 8) & 0xff, out);
      putc((color >> 16) & 0xff, out);
    }
    for (x = 0; x < padding; x++) { putc(0, out); }
  }

  fclose(out);
}

int main(int argc, char *argv[])
{
  struct timeval tv_start, tv_end;
  int *picture;

  float real_start = 0.37 - 0.00;
  float real_end = 0.37 + 0.04;
  float imaginary_start = -0.2166 - 0.02;
  float imaginary_end = -0.2166 + 0.02;

  int do_cuda = 0;

  if (argc != 2)
  {
    printf("Usage: %s <normal/cuda/cuda128>\n", argv[0]);
    exit(0);
  }

  if (strcmp(argv[1], "normal") == 0)
  {
    do_cuda = 0;
  }
    else
  if (strcmp(argv[1], "cuda") == 0)
  {
    do_cuda = 1;
  }
    else
  if (strcmp(argv[1], "cuda128") == 0)
  {
    do_cuda = 2;
  }

  const int length = WIDTH * HEIGHT * sizeof(int);

  switch (do_cuda)
  {
    case 0:
      picture = (int *)malloc(length);
      break;
    case 1:
    case 2:
      cudaMallocManaged(&picture, length);
      break;
  }

  gettimeofday(&tv_start, NULL);

  if (do_cuda == 1)
  {
    mandel_calc_cuda_single<<<1,1>>>(picture, WIDTH, HEIGHT, real_start, real_end, imaginary_start, imaginary_end);

    cudaDeviceSynchronize();
  }
    else
  if (do_cuda == 2)
  {
    mandel_calc_cuda_multi<<<1,128>>>(picture, WIDTH, HEIGHT, real_start, real_end, imaginary_start, imaginary_end);

    cudaDeviceSynchronize();
  }
    else
  {
    mandel_calc(picture, WIDTH, HEIGHT, real_start, real_end, imaginary_start, imaginary_end);
  }

  gettimeofday(&tv_end, NULL);

  printf("%ld %ld\n", tv_end.tv_sec, tv_end.tv_usec);
  printf("%ld %ld\n", tv_start.tv_sec, tv_start.tv_usec);
  long time_diff = tv_end.tv_usec - tv_start.tv_usec;
  while(time_diff < 0) { tv_end.tv_sec--; time_diff += 1000000; }
  time_diff += (tv_end.tv_sec - tv_start.tv_sec) * 1000000;
  printf("time=%f\n", (float)time_diff / 1000000);

  switch (do_cuda)
  {
    case 0:
    {
      write_bmp(picture, WIDTH, HEIGHT);
      free(picture);
      break;
    }
    case 1:
    case 2:
    {
      int *image = (int *)malloc(length);
      cudaMemcpy(image, picture, length, cudaMemcpyDeviceToHost);
      write_bmp(image, WIDTH, HEIGHT);
      cudaFree(picture);
      break;
    }
  }

  return 0;
}


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

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2047:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2047:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z11mandel_calcPiiiffff
	.type	_Z11mandel_calcPiiiffff, @function
_Z11mandel_calcPiiiffff:
.LFB2040:
	.cfi_startproc
	endbr64
	movl	%esi, %ecx
	subss	%xmm0, %xmm1
	movaps	%xmm0, %xmm5
	xorl	%r10d, %r10d
	cvtsi2ssl	%ecx, %xmm0
	subss	%xmm2, %xmm3
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	testl	%ecx, %ecx
	movl	%edx, %esi
	leaq	_ZL6colors(%rip), %r11
	cmovns	%ecx, %r10d
	divss	%xmm0, %xmm1
	cvtsi2ssl	%edx, %xmm0
	divss	%xmm0, %xmm3
.L3:
	cmpl	%esi, %r9d
	jge	.L6
	movslq	%r8d, %rdx
	movaps	%xmm5, %xmm6
.L7:
	movl	%edx, %eax
	subl	%r8d, %eax
	cmpl	%ecx, %eax
	jge	.L11
	xorps	%xmm4, %xmm4
	xorl	%eax, %eax
	movaps	%xmm4, %xmm0
.L5:
	movaps	%xmm0, %xmm7
	movaps	%xmm4, %xmm8
	mulss	%xmm0, %xmm7
	addss	%xmm0, %xmm0
	mulss	%xmm4, %xmm8
	mulss	%xmm0, %xmm4
	subss	%xmm8, %xmm7
	addss	%xmm2, %xmm4
	addss	%xmm6, %xmm7
	movaps	%xmm4, %xmm8
	mulss	%xmm4, %xmm8
	movaps	%xmm7, %xmm0
	mulss	%xmm7, %xmm7
	addss	%xmm8, %xmm7
	comiss	.LC1(%rip), %xmm7
	ja	.L4
	incl	%eax
	cmpl	$127, %eax
	jne	.L5
.L4:
	sarl	$3, %eax
	addss	%xmm1, %xmm6
	cltq
	movl	(%r11,%rax,4), %eax
	movl	%eax, (%rdi,%rdx,4)
	incq	%rdx
	jmp	.L7
.L11:
	addss	%xmm3, %xmm2
	addl	%r10d, %r8d
	incl	%r9d
	jmp	.L3
.L6:
	xorl	%eax, %eax
	ret
	.cfi_endproc
.LFE2040:
	.size	_Z11mandel_calcPiiiffff, .-_Z11mandel_calcPiiiffff
	.globl	_Z11write_int32P8_IO_FILEi
	.type	_Z11write_int32P8_IO_FILEi, @function
_Z11write_int32P8_IO_FILEi:
.LFB2041:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movq	%rdi, %rbp
	movzbl	%sil, %edi
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movl	%esi, %ebx
	movq	%rbp, %rsi
	pushq	%rax
	.cfi_def_cfa_offset 32
	call	putc@PLT
	movzbl	%bh, %edi
	movq	%rbp, %rsi
	call	putc@PLT
	movl	%ebx, %edi
	movq	%rbp, %rsi
	shrl	$24, %ebx
	sarl	$16, %edi
	movzbl	%dil, %edi
	call	putc@PLT
	movl	%ebx, %edi
	movq	%rbp, %rsi
	call	putc@PLT
	popq	%rdx
	.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
.LFE2041:
	.size	_Z11write_int32P8_IO_FILEi, .-_Z11write_int32P8_IO_FILEi
	.globl	_Z11write_int16P8_IO_FILEi
	.type	_Z11write_int16P8_IO_FILEi, @function
_Z11write_int16P8_IO_FILEi:
.LFB2042:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movq	%rdi, %rbp
	movzbl	%sil, %edi
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movl	%esi, %ebx
	movq	%rbp, %rsi
	pushq	%rax
	.cfi_def_cfa_offset 32
	call	putc@PLT
	movzbl	%bh, %edi
	movq	%rbp, %rsi
	call	putc@PLT
	popq	%rdx
	.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
.LFE2042:
	.size	_Z11write_int16P8_IO_FILEi, .-_Z11write_int16P8_IO_FILEi
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC2:
	.string	"wb"
.LC3:
	.string	"out.bmp"
.LC4:
	.string	"Can't open file for writing."
	.text
	.globl	_Z9write_bmpPiii
	.type	_Z9write_bmpPiii, @function
_Z9write_bmpPiii:
.LFB2043:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movl	%edx, %r14d
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movl	%esi, %r12d
	leaq	.LC2(%rip), %rsi
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%rdi, %rbp
	leaq	.LC3(%rip), %rdi
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$40, %rsp
	.cfi_def_cfa_offset 96
	call	fopen@PLT
	testq	%rax, %rax
	jne	.L17
	addq	$40, %rsp
	.cfi_remember_state
	.cfi_def_cfa_offset 56
	leaq	.LC4(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	jmp	__printf_chk@PLT
.L17:
	.cfi_restore_state
	movq	%rax, %rbx
	leal	1(%r12), %eax
	movl	$66, %edi
	xorl	%r15d, %r15d
	imull	$-3, %r12d, %edx
	leal	(%rax,%rax,2), %eax
	movq	%rbx, %rsi
	andl	$-4, %eax
	movl	%eax, %r13d
	imull	%r14d, %r13d
	addl	%edx, %eax
	movl	%eax, 12(%rsp)
	call	putc@PLT
	movq	%rbx, %rsi
	movl	$77, %edi
	call	putc@PLT
	leal	54(%r13), %esi
	movq	%rbx, %rdi
	call	_Z11write_int32P8_IO_FILEi
	xorl	%esi, %esi
	movq	%rbx, %rdi
	call	_Z11write_int16P8_IO_FILEi
	xorl	%esi, %esi
	movq	%rbx, %rdi
	call	_Z11write_int16P8_IO_FILEi
	movl	$54, %esi
	movq	%rbx, %rdi
	call	_Z11write_int32P8_IO_FILEi
	movl	$40, %esi
	movq	%rbx, %rdi
	call	_Z11write_int32P8_IO_FILEi
	movl	%r12d, %esi
	movq	%rbx, %rdi
	call	_Z11write_int32P8_IO_FILEi
	movl	%r14d, %esi
	movq	%rbx, %rdi
	call	_Z11write_int32P8_IO_FILEi
	movl	$1, %esi
	movq	%rbx, %rdi
	call	_Z11write_int16P8_IO_FILEi
	movl	$24, %esi
	movq	%rbx, %rdi
	call	_Z11write_int16P8_IO_FILEi
	xorl	%esi, %esi
	movq	%rbx, %rdi
	call	_Z11write_int32P8_IO_FILEi
	movl	%r13d, %esi
	movq	%rbx, %rdi
	call	_Z11write_int32P8_IO_FILEi
	xorl	%esi, %esi
	movq	%rbx, %rdi
	call	_Z11write_int32P8_IO_FILEi
	xorl	%esi, %esi
	movq	%rbx, %rdi
	call	_Z11write_int32P8_IO_FILEi
	xorl	%esi, %esi
	movq	%rbx, %rdi
	call	_Z11write_int32P8_IO_FILEi
	xorl	%esi, %esi
	movq	%rbx, %rdi
	call	_Z11write_int32P8_IO_FILEi
	movslq	%r12d, %rax
	salq	$2, %rax
	movq	%rax, 16(%rsp)
.L18:
	cmpl	%r14d, %r15d
	jge	.L25
	xorl	%edx, %edx
.L22:
	cmpl	%edx, %r12d
	jle	.L26
	movl	0(%rbp,%rdx,4), %r13d
	movq	%rbx, %rsi
	movq	%rdx, 24(%rsp)
	movzbl	%r13b, %edi
	call	putc@PLT
	movl	%r13d, %eax
	movq	%rbx, %rsi
	sarl	$16, %r13d
	movzbl	%ah, %edi
	call	putc@PLT
	movzbl	%r13b, %edi
	movq	%rbx, %rsi
	call	putc@PLT
	movq	24(%rsp), %rdx
	incq	%rdx
	jmp	.L22
.L26:
	xorl	%r13d, %r13d
.L20:
	movl	12(%rsp), %eax
	cmpl	%eax, %r13d
	jge	.L27
	movq	%rbx, %rsi
	xorl	%edi, %edi
	incl	%r13d
	call	putc@PLT
	jmp	.L20
.L27:
	movq	16(%rsp), %rax
	incl	%r15d
	addq	%rax, %rbp
	jmp	.L18
.L25:
	addq	$40, %rsp
	.cfi_def_cfa_offset 56
	movq	%rbx, %rdi
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	jmp	fclose@PLT
	.cfi_endproc
.LFE2043:
	.size	_Z9write_bmpPiii, .-_Z9write_bmpPiii
	.globl	_Z49__device_stub__Z23mandel_calc_cuda_singlePiiiffffPiiiffff
	.type	_Z49__device_stub__Z23mandel_calc_cuda_singlePiiiffffPiiiffff, @function
_Z49__device_stub__Z23mandel_calc_cuda_singlePiiiffffPiiiffff:
.LFB2069:
	.cfi_startproc
	endbr64
	subq	$168, %rsp
	.cfi_def_cfa_offset 176
	movq	%rdi, 24(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movl	%esi, 20(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 16(%rsp)
	leaq	32(%rsp), %rdx
	movss	%xmm0, 12(%rsp)
	movss	%xmm1, 8(%rsp)
	movss	%xmm2, 4(%rsp)
	movss	%xmm3, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 152(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 136(%rsp)
	movq	%rsp, %rax
	movq	%rax, 144(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L28
	pushq	40(%rsp)
	.cfi_def_cfa_offset 184
	leaq	_Z23mandel_calc_cuda_singlePiiiffff(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 192
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 184
	popq	%rdx
	.cfi_def_cfa_offset 176
.L28:
	movq	152(%rsp), %rax
	subq	%fs:40, %rax
	je	.L30
	call	__stack_chk_fail@PLT
.L30:
	addq	$168, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2069:
	.size	_Z49__device_stub__Z23mandel_calc_cuda_singlePiiiffffPiiiffff, .-_Z49__device_stub__Z23mandel_calc_cuda_singlePiiiffffPiiiffff
	.globl	_Z23mandel_calc_cuda_singlePiiiffff
	.type	_Z23mandel_calc_cuda_singlePiiiffff, @function
_Z23mandel_calc_cuda_singlePiiiffff:
.LFB2070:
	.cfi_startproc
	endbr64
	jmp	_Z49__device_stub__Z23mandel_calc_cuda_singlePiiiffffPiiiffff
	.cfi_endproc
.LFE2070:
	.size	_Z23mandel_calc_cuda_singlePiiiffff, .-_Z23mandel_calc_cuda_singlePiiiffff
	.globl	_Z48__device_stub__Z22mandel_calc_cuda_multiPiiiffffPiiiffff
	.type	_Z48__device_stub__Z22mandel_calc_cuda_multiPiiiffffPiiiffff, @function
_Z48__device_stub__Z22mandel_calc_cuda_multiPiiiffffPiiiffff:
.LFB2071:
	.cfi_startproc
	endbr64
	subq	$168, %rsp
	.cfi_def_cfa_offset 176
	movq	%rdi, 24(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movl	%esi, 20(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 16(%rsp)
	leaq	32(%rsp), %rdx
	movss	%xmm0, 12(%rsp)
	movss	%xmm1, 8(%rsp)
	movss	%xmm2, 4(%rsp)
	movss	%xmm3, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 152(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 136(%rsp)
	movq	%rsp, %rax
	movq	%rax, 144(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L33
	pushq	40(%rsp)
	.cfi_def_cfa_offset 184
	leaq	_Z22mandel_calc_cuda_multiPiiiffff(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 192
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 184
	popq	%rdx
	.cfi_def_cfa_offset 176
.L33:
	movq	152(%rsp), %rax
	subq	%fs:40, %rax
	je	.L35
	call	__stack_chk_fail@PLT
.L35:
	addq	$168, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2071:
	.size	_Z48__device_stub__Z22mandel_calc_cuda_multiPiiiffffPiiiffff, .-_Z48__device_stub__Z22mandel_calc_cuda_multiPiiiffffPiiiffff
	.globl	_Z22mandel_calc_cuda_multiPiiiffff
	.type	_Z22mandel_calc_cuda_multiPiiiffff, @function
_Z22mandel_calc_cuda_multiPiiiffff:
.LFB2072:
	.cfi_startproc
	endbr64
	jmp	_Z48__device_stub__Z22mandel_calc_cuda_multiPiiiffffPiiiffff
	.cfi_endproc
.LFE2072:
	.size	_Z22mandel_calc_cuda_multiPiiiffff, .-_Z22mandel_calc_cuda_multiPiiiffff
	.section	.rodata.str1.1
.LC5:
	.string	"Usage: %s <normal/cuda/cuda128>\n"
.LC6:
	.string	"normal"
.LC7:
	.string	"cuda"
.LC8:
	.string	"cuda128"
.LC13:
	.string	"%ld %ld\n"
.LC15:
	.string	"time=%f\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2044:
	.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	$72, %rsp
	.cfi_def_cfa_offset 96
	movq	%fs:40, %rax
	movq	%rax, 56(%rsp)
	xorl	%eax, %eax
	cmpl	$2, %edi
	je	.L39
	movq	(%rsi), %rdx
	movl	$2, %edi
	leaq	.LC5(%rip), %rsi
	call	__printf_chk@PLT
	xorl	%edi, %edi
	call	exit@PLT
.L39:
	movq	8(%rsi), %rbp
	movl	%edi, %ebx
	leaq	.LC6(%rip), %rsi
	movq	%rbp, %rdi
	call	strcmp@PLT
	testl	%eax, %eax
	je	.L40
	leaq	.LC7(%rip), %rsi
	movq	%rbp, %rdi
	call	strcmp@PLT
	testl	%eax, %eax
	je	.L54
	leaq	.LC8(%rip), %rsi
	movq	%rbp, %rdi
	call	strcmp@PLT
	testl	%eax, %eax
	je	.L41
.L40:
	movl	$3145728, %edi
	xorl	%ebx, %ebx
	call	malloc@PLT
	movq	%rax, (%rsp)
	jmp	.L42
.L54:
	movl	$1, %ebx
.L41:
	movq	%rsp, %rdi
	movl	$1, %edx
	movl	$3145728, %esi
	call	cudaMallocManaged@PLT
.L42:
	xorl	%esi, %esi
	leaq	24(%rsp), %rdi
	call	gettimeofday@PLT
	cmpl	$1, %ebx
	jne	.L43
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	movl	$1, %esi
	movabsq	$4294967297, %rdi
	movl	$1, 48(%rsp)
	movq	%rdi, %rdx
	movq	%rdi, 40(%rsp)
	movq	%rdi, 12(%rsp)
	movl	$1, 20(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L47
	movss	.LC9(%rip), %xmm3
	movss	.LC10(%rip), %xmm2
	movl	$768, %edx
	movl	$1024, %esi
	movss	.LC11(%rip), %xmm1
	movss	.LC12(%rip), %xmm0
	movq	(%rsp), %rdi
	call	_Z49__device_stub__Z23mandel_calc_cuda_singlePiiiffffPiiiffff
	jmp	.L47
.L43:
	cmpl	$2, %ebx
	jne	.L46
	movl	$33554433, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	salq	$7, %rdx
	movl	$1, %esi
	movabsq	$4294967297, %rdi
	movl	$1, 48(%rsp)
	movq	%rdx, 40(%rsp)
	movq	%rdi, 12(%rsp)
	movl	$1, 20(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L47
	movss	.LC9(%rip), %xmm3
	movss	.LC10(%rip), %xmm2
	movl	$768, %edx
	movl	$1024, %esi
	movss	.LC11(%rip), %xmm1
	movss	.LC12(%rip), %xmm0
	movq	(%rsp), %rdi
	call	_Z48__device_stub__Z22mandel_calc_cuda_multiPiiiffffPiiiffff
.L47:
	call	cudaDeviceSynchronize@PLT
	jmp	.L45
.L46:
	movss	.LC9(%rip), %xmm3
	movss	.LC10(%rip), %xmm2
	movl	$768, %edx
	movl	$1024, %esi
	movss	.LC11(%rip), %xmm1
	movss	.LC12(%rip), %xmm0
	movq	(%rsp), %rdi
	call	_Z11mandel_calcPiiiffff
.L45:
	leaq	40(%rsp), %rdi
	xorl	%esi, %esi
	leaq	.LC13(%rip), %rbp
	call	gettimeofday@PLT
	movq	48(%rsp), %rcx
	movq	%rbp, %rsi
	xorl	%eax, %eax
	movq	40(%rsp), %rdx
	movl	$2, %edi
	call	__printf_chk@PLT
	movq	32(%rsp), %rcx
	movq	%rbp, %rsi
	xorl	%eax, %eax
	movq	24(%rsp), %rdx
	movl	$2, %edi
	call	__printf_chk@PLT
	movq	40(%rsp), %rax
	movq	48(%rsp), %rdx
	xorl	%ecx, %ecx
	subq	32(%rsp), %rdx
.L48:
	leaq	-1(%rax), %rsi
	testq	%rdx, %rdx
	jns	.L65
	addq	$1000000, %rdx
	movq	%rsi, %rax
	movb	$1, %cl
	jmp	.L48
.L65:
	testb	%cl, %cl
	je	.L50
	movq	%rax, 40(%rsp)
.L50:
	movq	40(%rsp), %rax
	subq	24(%rsp), %rax
	leaq	.LC15(%rip), %rsi
	decl	%ebx
	imulq	$1000000, %rax, %rax
	movl	$2, %edi
	addq	%rdx, %rax
	cvtsi2ssq	%rax, %xmm0
	movb	$1, %al
	divss	.LC14(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	call	__printf_chk@PLT
	cmpl	$1, %ebx
	jbe	.L51
	movq	(%rsp), %rdi
	movl	$768, %edx
	movl	$1024, %esi
	call	_Z9write_bmpPiii
	movq	(%rsp), %rdi
	call	free@PLT
	jmp	.L52
.L51:
	movl	$3145728, %edi
	call	malloc@PLT
	movq	(%rsp), %rsi
	movl	$2, %ecx
	movl	$3145728, %edx
	movq	%rax, %rbx
	movq	%rax, %rdi
	call	cudaMemcpy@PLT
	movq	%rbx, %rdi
	movl	$768, %edx
	movl	$1024, %esi
	call	_Z9write_bmpPiii
	movq	(%rsp), %rdi
	call	cudaFree@PLT
.L52:
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	je	.L53
	call	__stack_chk_fail@PLT
.L53:
	addq	$72, %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
.LFE2044:
	.size	main, .-main
	.section	.rodata.str1.1
.LC16:
	.string	"_Z22mandel_calc_cuda_multiPiiiffff"
.LC17:
	.string	"_Z23mandel_calc_cuda_singlePiiiffff"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2074:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -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	.LC16(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z22mandel_calc_cuda_multiPiiiffff(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC17(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z23mandel_calc_cuda_singlePiiiffff(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	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	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2074:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.section	.rodata
	.align 32
	.type	_ZL6colors, @object
	.size	_ZL6colors, 64
_ZL6colors:
	.long	16711680
	.long	15610624
	.long	13391104
	.long	11162880
	.long	11154176
	.long	6710784
	.long	10066176
	.long	6723840
	.long	3381504
	.long	39338
	.long	26282
	.long	13226
	.long	170
	.long	153
	.long	102
	.long	0
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC1:
	.long	1082130432
	.align 4
.LC9:
	.long	-1102491261
	.align 4
.LC10:
	.long	-1099806907
	.align 4
.LC11:
	.long	1053944709
	.align 4
.LC12:
	.long	1052602532
	.align 4
.LC14:
	.long	1232348160
	.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 <stdint.h>
#include <string.h>
#include <unistd.h>
#include <sys/time.h>

#define WIDTH 1024
#define HEIGHT 768
#define CORE_COUNT 128

static int colors[] =
{
  0xff0000,  // f
  0xee3300,  // e
  0xcc5500,  // d
  0xaa5500,  // c
  0xaa3300,  // b
  0x666600,  // a
  0x999900,  // 9
  0x669900,  // 8
  0x339900,  // 7
  0x0099aa,  // 6
  0x0066aa,  // 5
  0x0033aa,  // 4
  0x0000aa,  // 3
  0x000099,  // 2
  0x000066,  // 1
  0x000000,  // 0
};

struct MandelInfo
{
  float r_step;
  float i_step;
  float real_start;
  float imaginary_start;
  int width;
  int height;
};

struct CoreInfo
{
  MandelInfo mandel_info;
  int *picture;
  uint8_t signal_start;
  uint8_t signal_done;
  int core_id;
};

int mandel_calc(
  int *picture,
  int width,
  int height,
  float real_start,
  float real_end,
  float imaginary_start,
  float imaginary_end)
{
  const int max_count = 127;
  int x, y;
  float r, i, r_step, i_step;
  float tr, ti, zr, zi;
  int ptr, count;

  r_step = (real_end - real_start) / (float)width;
  i_step = (imaginary_end - imaginary_start) / (float)height;
  ptr = 0;

  //printf("step = %f %f\n", r_step, i_step);

  i = imaginary_start;

  for (y = 0; y < height; y++)
  {
    r = real_start;

    for (x = 0; x < width; x++)
    {
      zr = 0;
      zi = 0;

      for (count = 0; count < max_count; count++)
      {
        tr = ((zr * zr) - (zi * zi));
        ti = (2 * zr * zi);
        zr = tr + r;
        zi = ti + i;
        if ((zr * zr) + (zi * zi) > 4) break;
      }

      picture[ptr] = colors[count >> 3];

      ptr++;
      r = r + r_step;
    }

    i = i + i_step;
  }

  return 0;
}

__global__
void mandel_calc_cuda_single(
  int *picture,
  int width,
  int height,
  float real_start,
  float real_end,
  float imaginary_start,
  float imaginary_end)
{
  const int max_count = 127;
  int x,y;
  float r, i, r_step, i_step;
  float tr, ti, zr, zi;
  int ptr, count;

  int colors[] =
  {
    0xff0000,  // f
    0xee3300,  // e
    0xcc5500,  // d
    0xaa5500,  // c
    0xaa3300,  // b
    0x666600,  // a
    0x999900,  // 9
    0x669900,  // 8
    0x339900,  // 7
    0x0099aa,  // 6
    0x0066aa,  // 5
    0x0033aa,  // 4
    0x0000aa,  // 3
    0x000099,  // 2
    0x000066,  // 1
    0x000000,  // 0
  };

  r_step = (real_end - real_start) / (float)width;
  i_step = (imaginary_end - imaginary_start) / (float)height;
  ptr = 0;

//printf("step = %f %f\n", r_step, i_step);

  i = imaginary_start;

  for (y = 0; y < height; y++)
  {
    r = real_start;

    for (x = 0; x < width; x++)
    {
      zr = 0;
      zi = 0;

      for (count = 0; count < max_count; count++)
      {
        tr = ((zr * zr) - (zi * zi));
        ti = (2 * zr * zi);
        zr = tr + r;
        zi = ti + i;
        if ((zr * zr) + (zi * zi) > 4) break;
      }

      picture[ptr] = colors[count >> 3];

      ptr++;
      r = r + r_step;
    }

    i = i + i_step;
  }
}

__global__
void mandel_calc_cuda_multi(
  int *picture,
  int width,
  int height,
  float real_start,
  float real_end,
  float imaginary_start,
  float imaginary_end)
{
  const int max_count = 127;
  int x,y;
  float r, i, r_step, i_step;
  float tr, ti, zr, zi;
  int ptr, count;

  int index = threadIdx.x;
  //int stride = blockDim.x;

//printf("index=%d stride=%d\n", index, stride);

  int colors[] =
  {
    0xff0000,  // f
    0xee3300,  // e
    0xcc5500,  // d
    0xaa5500,  // c
    0xaa3300,  // b
    0x666600,  // a
    0x999900,  // 9
    0x669900,  // 8
    0x339900,  // 7
    0x0099aa,  // 6
    0x0066aa,  // 5
    0x0033aa,  // 4
    0x0000aa,  // 3
    0x000099,  // 2
    0x000066,  // 1
    0x000000,  // 0
  };

  r_step = (real_end - real_start) / (float)width;
  i_step = (imaginary_end - imaginary_start) / (float)height;

//printf("step = %f %f\n", r_step, i_step);

  height = height / CORE_COUNT;

  y = index * height;

  ptr = y * width;

  i = imaginary_start + (i_step * y);
#if 0
  i = imaginary_start;

  for (y = 0; y < index * height; y++)
  {
    i = i + i_step;
  }
#endif

//printf("ptr=%d index=%d height=%d %f\n", ptr, index, height, i);

  for (y = 0; y < height; y++)
  {
    r = real_start;

    for (x = 0; x < width; x++)
    {
      zr = 0;
      zi = 0;

      for (count = 0; count < max_count; count++)
      {
        tr = ((zr * zr) - (zi * zi));
        ti = (2 * zr * zi);
        zr = tr + r;
        zi = ti + i;
        if ((zr * zr) + (zi * zi) > 4) break;
      }

      picture[ptr] = colors[count >> 3];

      ptr++;
      r = r + r_step;
    }

    i = i + i_step;
  }
}

int write_int32(FILE *out, int n)
{
  putc((n & 0xff), out);
  putc(((n >> 8) & 0xff), out);
  putc(((n >> 16) & 0xff), out);
  putc(((n >> 24) & 0xff), out);

  return 0;
}

int write_int16(FILE *out, int n)
{
  putc((n & 0xff), out);
  putc(((n >> 8) & 0xff), out);

  return 0;
}

void write_bmp(int *picture, int width, int height)
{
  FILE *out;
  int bmp_width;
  int bmp_size;
  int padding;
  int offset;
  int color;
  int x,y;

  out = fopen("out.bmp", "wb");
  if (out == NULL)
  {
    printf("Can't open file for writing.");
    return;
  }

  bmp_width = width * 3;
  bmp_width = (bmp_width + 3) & (~0x3);
  bmp_size = (bmp_width * height) + 14 + 40;
  padding = bmp_width - (width * 3);

  //printf("width=%d (%d)\n", width, width*3);
  //printf("bmp_width=%d\n", bmp_width);
  //printf("bmp_size=%d\n", bmp_size);

  /* size: 14 bytes */

  putc('B', out);
  putc('M', out);
  write_int32(out, bmp_size);
  write_int16(out, 0);
  write_int16(out, 0);
  write_int32(out, 54);

  /* head1: 14  head2: 40 */

  write_int32(out, 40);         /* biSize */
  write_int32(out, width);
  write_int32(out, height);
  write_int16(out, 1);
  write_int16(out, 24);
  write_int32(out, 0);          /* compression */
  write_int32(out, bmp_width*height);
  write_int32(out, 0);          /* biXPelsperMetre */
  write_int32(out, 0);          /* biYPelsperMetre */
  write_int32(out, 0);
  write_int32(out, 0);

  for (y = 0; y < height; y++)
  {
    offset = y * width;

    for (x = 0; x < width; x++)
    {
      color = picture[offset++];

      putc(color & 0xff, out);
      putc((color >> 8) & 0xff, out);
      putc((color >> 16) & 0xff, out);
    }
    for (x = 0; x < padding; x++) { putc(0, out); }
  }

  fclose(out);
}

int main(int argc, char *argv[])
{
  struct timeval tv_start, tv_end;
  int *picture;

  float real_start = 0.37 - 0.00;
  float real_end = 0.37 + 0.04;
  float imaginary_start = -0.2166 - 0.02;
  float imaginary_end = -0.2166 + 0.02;

  int do_cuda = 0;

  if (argc != 2)
  {
    printf("Usage: %s <normal/cuda/cuda128>\n", argv[0]);
    exit(0);
  }

  if (strcmp(argv[1], "normal") == 0)
  {
    do_cuda = 0;
  }
    else
  if (strcmp(argv[1], "cuda") == 0)
  {
    do_cuda = 1;
  }
    else
  if (strcmp(argv[1], "cuda128") == 0)
  {
    do_cuda = 2;
  }

  const int length = WIDTH * HEIGHT * sizeof(int);

  switch (do_cuda)
  {
    case 0:
      picture = (int *)malloc(length);
      break;
    case 1:
    case 2:
      hipMallocManaged(&picture, length);
      break;
  }

  gettimeofday(&tv_start, NULL);

  if (do_cuda == 1)
  {
    mandel_calc_cuda_single<<<1,1>>>(picture, WIDTH, HEIGHT, real_start, real_end, imaginary_start, imaginary_end);

    hipDeviceSynchronize();
  }
    else
  if (do_cuda == 2)
  {
    mandel_calc_cuda_multi<<<1,128>>>(picture, WIDTH, HEIGHT, real_start, real_end, imaginary_start, imaginary_end);

    hipDeviceSynchronize();
  }
    else
  {
    mandel_calc(picture, WIDTH, HEIGHT, real_start, real_end, imaginary_start, imaginary_end);
  }

  gettimeofday(&tv_end, NULL);

  printf("%ld %ld\n", tv_end.tv_sec, tv_end.tv_usec);
  printf("%ld %ld\n", tv_start.tv_sec, tv_start.tv_usec);
  long time_diff = tv_end.tv_usec - tv_start.tv_usec;
  while(time_diff < 0) { tv_end.tv_sec--; time_diff += 1000000; }
  time_diff += (tv_end.tv_sec - tv_start.tv_sec) * 1000000;
  printf("time=%f\n", (float)time_diff / 1000000);

  switch (do_cuda)
  {
    case 0:
    {
      write_bmp(picture, WIDTH, HEIGHT);
      free(picture);
      break;
    }
    case 1:
    case 2:
    {
      int *image = (int *)malloc(length);
      hipMemcpy(image, picture, length, hipMemcpyDeviceToHost);
      write_bmp(image, WIDTH, HEIGHT);
      hipFree(picture);
      break;
    }
  }

  return 0;
}


	.text
	.file	"mandelbrot_cuda.hip"
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function _Z11mandel_calcPiiiffff
.LCPI0_0:
	.long	0x40800000                      # float 4
	.text
	.globl	_Z11mandel_calcPiiiffff
	.type	_Z11mandel_calcPiiiffff,@function
_Z11mandel_calcPiiiffff:                # @_Z11mandel_calcPiiiffff
	.cfi_startproc
# %bb.0:
	testl	%edx, %edx
	jle	.LBB0_10
# %bb.1:                                # %.preheader44.lr.ph
	subss	%xmm0, %xmm1
	cvtsi2ss	%esi, %xmm4
	divss	%xmm4, %xmm1
	subss	%xmm2, %xmm3
	xorps	%xmm4, %xmm4
	cvtsi2ss	%edx, %xmm4
	divss	%xmm4, %xmm3
	xorl	%eax, %eax
	movss	.LCPI0_0(%rip), %xmm4           # xmm4 = mem[0],zero,zero,zero
	xorl	%ecx, %ecx
.LBB0_2:                                # %.preheader44
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB0_4 Depth 2
                                        #       Child Loop BB0_5 Depth 3
	testl	%esi, %esi
	jle	.LBB0_9
# %bb.3:                                # %.preheader.preheader
                                        #   in Loop: Header=BB0_2 Depth=1
	cltq
	xorl	%r8d, %r8d
	movaps	%xmm0, %xmm5
.LBB0_4:                                # %.preheader
                                        #   Parent Loop BB0_2 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB0_5 Depth 3
	xorps	%xmm6, %xmm6
	xorl	%r9d, %r9d
	xorps	%xmm7, %xmm7
.LBB0_5:                                #   Parent Loop BB0_2 Depth=1
                                        #     Parent Loop BB0_4 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	movaps	%xmm7, %xmm8
	mulss	%xmm7, %xmm8
	movaps	%xmm6, %xmm9
	mulss	%xmm6, %xmm9
	subss	%xmm9, %xmm8
	addss	%xmm7, %xmm7
	mulss	%xmm7, %xmm6
	movaps	%xmm8, %xmm7
	addss	%xmm5, %xmm7
	addss	%xmm2, %xmm6
	movaps	%xmm7, %xmm8
	mulss	%xmm7, %xmm8
	movaps	%xmm6, %xmm9
	mulss	%xmm6, %xmm9
	addss	%xmm8, %xmm9
	ucomiss	%xmm4, %xmm9
	ja	.LBB0_8
# %bb.6:                                #   in Loop: Header=BB0_5 Depth=3
	incl	%r9d
	cmpl	$127, %r9d
	jne	.LBB0_5
# %bb.7:                                #   in Loop: Header=BB0_4 Depth=2
	movl	$127, %r9d
.LBB0_8:                                #   in Loop: Header=BB0_4 Depth=2
	shrl	$3, %r9d
	movl	_ZL6colors(,%r9,4), %r9d
	movl	%r9d, (%rdi,%rax,4)
	incq	%rax
	addss	%xmm1, %xmm5
	incl	%r8d
	cmpl	%esi, %r8d
	jne	.LBB0_4
.LBB0_9:                                # %._crit_edge
                                        #   in Loop: Header=BB0_2 Depth=1
	addss	%xmm3, %xmm2
	incl	%ecx
	cmpl	%edx, %ecx
	jne	.LBB0_2
.LBB0_10:                               # %._crit_edge54
	xorl	%eax, %eax
	retq
.Lfunc_end0:
	.size	_Z11mandel_calcPiiiffff, .Lfunc_end0-_Z11mandel_calcPiiiffff
	.cfi_endproc
                                        # -- End function
	.globl	_Z38__device_stub__mandel_calc_cuda_singlePiiiffff # -- Begin function _Z38__device_stub__mandel_calc_cuda_singlePiiiffff
	.type	_Z38__device_stub__mandel_calc_cuda_singlePiiiffff,@function
_Z38__device_stub__mandel_calc_cuda_singlePiiiffff: # @_Z38__device_stub__mandel_calc_cuda_singlePiiiffff
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$144, %rsp
	.cfi_def_cfa_offset 192
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	20(%rsp), %rcx
	movl	%esi, (%rcx)
	leaq	16(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	12(%rsp), %rdx
	movss	%xmm0, (%rdx)
	leaq	8(%rsp), %rdi
	movss	%xmm1, (%rdi)
	leaq	4(%rsp), %r8
	movss	%xmm2, (%r8)
	movq	%rsp, %r9
	movss	%xmm3, (%r9)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rdi, 32(%rbx)
	movq	%r8, 40(%rbx)
	movq	%r9, 48(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	32(%rsp), %r12
	leaq	24(%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	$_Z23mandel_calc_cuda_singlePiiiffff, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$160, %rsp
	.cfi_adjust_cfa_offset -160
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	_Z38__device_stub__mandel_calc_cuda_singlePiiiffff, .Lfunc_end1-_Z38__device_stub__mandel_calc_cuda_singlePiiiffff
	.cfi_endproc
                                        # -- End function
	.globl	_Z37__device_stub__mandel_calc_cuda_multiPiiiffff # -- Begin function _Z37__device_stub__mandel_calc_cuda_multiPiiiffff
	.type	_Z37__device_stub__mandel_calc_cuda_multiPiiiffff,@function
_Z37__device_stub__mandel_calc_cuda_multiPiiiffff: # @_Z37__device_stub__mandel_calc_cuda_multiPiiiffff
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$144, %rsp
	.cfi_def_cfa_offset 192
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	20(%rsp), %rcx
	movl	%esi, (%rcx)
	leaq	16(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	12(%rsp), %rdx
	movss	%xmm0, (%rdx)
	leaq	8(%rsp), %rdi
	movss	%xmm1, (%rdi)
	leaq	4(%rsp), %r8
	movss	%xmm2, (%r8)
	movq	%rsp, %r9
	movss	%xmm3, (%r9)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rdi, 32(%rbx)
	movq	%r8, 40(%rbx)
	movq	%r9, 48(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	32(%rsp), %r12
	leaq	24(%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	$_Z22mandel_calc_cuda_multiPiiiffff, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$160, %rsp
	.cfi_adjust_cfa_offset -160
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end2:
	.size	_Z37__device_stub__mandel_calc_cuda_multiPiiiffff, .Lfunc_end2-_Z37__device_stub__mandel_calc_cuda_multiPiiiffff
	.cfi_endproc
                                        # -- End function
	.globl	_Z11write_int32P8_IO_FILEi      # -- Begin function _Z11write_int32P8_IO_FILEi
	.type	_Z11write_int32P8_IO_FILEi,@function
_Z11write_int32P8_IO_FILEi:             # @_Z11write_int32P8_IO_FILEi
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %rbp, -16
	movl	%esi, %ebp
	movq	%rdi, %rbx
	movl	$255, %r14d
	movl	%esi, %edi
	andl	%r14d, %edi
	movq	%rbx, %rsi
	callq	putc
	movl	%ebp, %edi
	shrl	$8, %edi
	andl	%r14d, %edi
	movq	%rbx, %rsi
	callq	putc
	movl	%ebp, %edi
	shrl	$16, %edi
	andl	%r14d, %edi
	movq	%rbx, %rsi
	callq	putc
	shrl	$24, %ebp
	movl	%ebp, %edi
	movq	%rbx, %rsi
	callq	putc
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end3:
	.size	_Z11write_int32P8_IO_FILEi, .Lfunc_end3-_Z11write_int32P8_IO_FILEi
	.cfi_endproc
                                        # -- End function
	.globl	_Z11write_int16P8_IO_FILEi      # -- Begin function _Z11write_int16P8_IO_FILEi
	.type	_Z11write_int16P8_IO_FILEi,@function
_Z11write_int16P8_IO_FILEi:             # @_Z11write_int16P8_IO_FILEi
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %rbp, -16
	movl	%esi, %ebx
	movq	%rdi, %r14
	movl	$255, %ebp
	movl	%esi, %edi
	andl	%ebp, %edi
	movq	%r14, %rsi
	callq	putc
	shrl	$8, %ebx
	andl	%ebp, %ebx
	movl	%ebx, %edi
	movq	%r14, %rsi
	callq	putc
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end4:
	.size	_Z11write_int16P8_IO_FILEi, .Lfunc_end4-_Z11write_int16P8_IO_FILEi
	.cfi_endproc
                                        # -- End function
	.globl	_Z9write_bmpPiii                # -- Begin function _Z9write_bmpPiii
	.type	_Z9write_bmpPiii,@function
_Z9write_bmpPiii:                       # @_Z9write_bmpPiii
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$40, %rsp
	.cfi_def_cfa_offset 96
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	%edx, %r15d
	movl	%esi, %r13d
	movq	%rdi, 24(%rsp)                  # 8-byte Spill
	movl	$.L.str, %edi
	movl	$.L.str.1, %esi
	callq	fopen
	testq	%rax, %rax
	je	.LBB5_10
# %bb.1:
	movq	%rax, %r12
	leal	3(%r13,%r13,2), %eax
	andl	$-4, %eax
	movl	%eax, 4(%rsp)                   # 4-byte Spill
	movl	%eax, %ebx
	imull	%r15d, %ebx
	leal	54(%rbx), %ebp
	movl	$66, %edi
	movq	%r12, %rsi
	callq	putc
	movl	$77, %edi
	movq	%r12, %rsi
	callq	putc
	movq	%r12, %rdi
	movl	%ebp, %esi
	callq	_Z11write_int32P8_IO_FILEi
	xorl	%eax, %eax
	movq	%rax, 8(%rsp)                   # 8-byte Spill
	xorl	%edi, %edi
	movq	%r12, %rsi
	callq	putc
	xorl	%edi, %edi
	movq	%r12, %rsi
	callq	putc
	xorl	%edi, %edi
	movq	%r12, %rsi
	callq	putc
	xorl	%edi, %edi
	movq	%r12, %rsi
	callq	putc
	movq	%r12, %rdi
	movl	$54, %esi
	callq	_Z11write_int32P8_IO_FILEi
	movq	%r12, %rdi
	movl	$40, %esi
	callq	_Z11write_int32P8_IO_FILEi
	movq	%r12, %rdi
	movl	%r13d, %esi
	callq	_Z11write_int32P8_IO_FILEi
	movq	%r12, %rdi
	movl	%r15d, %esi
	callq	_Z11write_int32P8_IO_FILEi
	movl	$1, %edi
	movq	%r12, %rsi
	callq	putc
	xorl	%edi, %edi
	movq	%r12, %rsi
	callq	putc
	movl	$24, %edi
	movq	%r12, %rsi
	callq	putc
	xorl	%edi, %edi
	movq	%r12, %rsi
	callq	putc
	movq	%r12, %rdi
	xorl	%esi, %esi
	callq	_Z11write_int32P8_IO_FILEi
	movq	%r12, %rdi
	movl	%ebx, %esi
	callq	_Z11write_int32P8_IO_FILEi
	movq	%r12, %rdi
	xorl	%esi, %esi
	callq	_Z11write_int32P8_IO_FILEi
	movq	%r12, %rdi
	xorl	%esi, %esi
	callq	_Z11write_int32P8_IO_FILEi
	movq	%r12, %rdi
	xorl	%esi, %esi
	callq	_Z11write_int32P8_IO_FILEi
	movq	%r12, %rdi
	xorl	%esi, %esi
	callq	_Z11write_int32P8_IO_FILEi
	testl	%r15d, %r15d
	jle	.LBB5_9
# %bb.2:                                # %.lr.ph58
	leal	(%r13,%r13,2), %eax
	subl	%eax, 4(%rsp)                   # 4-byte Folded Spill
	movl	%r13d, %ebx
	movl	$255, %r14d
	xorl	%ecx, %ecx
	movl	%r15d, 16(%rsp)                 # 4-byte Spill
	movq	%r13, 32(%rsp)                  # 8-byte Spill
.LBB5_3:                                # =>This Loop Header: Depth=1
                                        #     Child Loop BB5_5 Depth 2
                                        #     Child Loop BB5_7 Depth 2
	movl	%ecx, 20(%rsp)                  # 4-byte Spill
	testl	%r13d, %r13d
	jle	.LBB5_6
# %bb.4:                                # %.lr.ph.preheader
                                        #   in Loop: Header=BB5_3 Depth=1
	movslq	8(%rsp), %rax                   # 4-byte Folded Reload
	movq	24(%rsp), %rcx                  # 8-byte Reload
	leaq	(%rcx,%rax,4), %r15
	xorl	%r13d, %r13d
.LBB5_5:                                # %.lr.ph
                                        #   Parent Loop BB5_3 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movl	(%r15,%r13,4), %ebp
	movl	%ebp, %edi
	andl	%r14d, %edi
	movq	%r12, %rsi
	callq	putc
	movl	%ebp, %edi
	shrl	$8, %edi
	andl	%r14d, %edi
	movq	%r12, %rsi
	callq	putc
	shrl	$16, %ebp
	andl	%r14d, %ebp
	movl	%ebp, %edi
	movq	%r12, %rsi
	callq	putc
	incq	%r13
	cmpl	%r13d, %ebx
	jne	.LBB5_5
.LBB5_6:                                # %.preheader
                                        #   in Loop: Header=BB5_3 Depth=1
	movl	4(%rsp), %eax                   # 4-byte Reload
	movl	%eax, %ebp
	testl	%eax, %eax
	jle	.LBB5_8
.LBB5_7:                                # %.lr.ph55
                                        #   Parent Loop BB5_3 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	xorl	%edi, %edi
	movq	%r12, %rsi
	callq	putc
	decl	%ebp
	jne	.LBB5_7
.LBB5_8:                                # %._crit_edge
                                        #   in Loop: Header=BB5_3 Depth=1
	movl	20(%rsp), %ecx                  # 4-byte Reload
	incl	%ecx
	movq	32(%rsp), %r13                  # 8-byte Reload
	movq	8(%rsp), %rax                   # 8-byte Reload
	addl	%r13d, %eax
	movq	%rax, 8(%rsp)                   # 8-byte Spill
	movl	16(%rsp), %r15d                 # 4-byte Reload
	cmpl	%r15d, %ecx
	jne	.LBB5_3
.LBB5_9:                                # %._crit_edge59
	movq	%r12, %rdi
	addq	$40, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	jmp	fclose                          # TAILCALL
.LBB5_10:
	.cfi_def_cfa_offset 96
	movl	$.L.str.2, %edi
	xorl	%eax, %eax
	addq	$40, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	jmp	printf                          # TAILCALL
.Lfunc_end5:
	.size	_Z9write_bmpPiii, .Lfunc_end5-_Z9write_bmpPiii
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function main
.LCPI6_0:
	.long	0x3ebd70a4                      # float 0.370000005
.LCPI6_1:
	.long	0x3ed1eb85                      # float 0.409999996
.LCPI6_2:
	.long	0xbe724745                      # float -0.236599997
.LCPI6_3:
	.long	0xbe495183                      # float -0.196600005
.LCPI6_4:
	.long	0x49742400                      # float 1.0E+6
	.text
	.globl	main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$48, %rsp
	.cfi_def_cfa_offset 96
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	cmpl	$2, %edi
	jne	.LBB6_16
# %bb.1:
	movq	8(%rsi), %rbx
	movl	$.L.str.4, %esi
	movq	%rbx, %rdi
	callq	strcmp
	testl	%eax, %eax
	je	.LBB6_4
# %bb.2:
	movl	$.L.str.5, %esi
	movq	%rbx, %rdi
	callq	strcmp
	testl	%eax, %eax
	je	.LBB6_5
# %bb.3:
	movl	$.L.str.6, %esi
	movq	%rbx, %rdi
	callq	strcmp
	testl	%eax, %eax
	je	.LBB6_7
.LBB6_4:                                # %.thread54
	movl	$3145728, %edi                  # imm = 0x300000
	callq	malloc
	movq	%rax, 8(%rsp)
	leaq	16(%rsp), %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	movq	8(%rsp), %rdi
	movss	.LCPI6_0(%rip), %xmm0           # xmm0 = mem[0],zero,zero,zero
	movss	.LCPI6_1(%rip), %xmm1           # xmm1 = mem[0],zero,zero,zero
	movss	.LCPI6_2(%rip), %xmm2           # xmm2 = mem[0],zero,zero,zero
	movss	.LCPI6_3(%rip), %xmm3           # xmm3 = mem[0],zero,zero,zero
	movl	$1024, %esi                     # imm = 0x400
	movl	$768, %edx                      # imm = 0x300
	callq	_Z11mandel_calcPiiiffff
	movb	$1, %bpl
	jmp	.LBB6_10
.LBB6_5:                                # %.critedge
	leaq	8(%rsp), %rdi
	movl	$3145728, %esi                  # imm = 0x300000
	movl	$1, %edx
	callq	hipMallocManaged
	leaq	16(%rsp), %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB6_9
# %bb.6:
	movq	8(%rsp), %rdi
	movss	.LCPI6_0(%rip), %xmm0           # xmm0 = mem[0],zero,zero,zero
	movss	.LCPI6_1(%rip), %xmm1           # xmm1 = mem[0],zero,zero,zero
	movss	.LCPI6_2(%rip), %xmm2           # xmm2 = mem[0],zero,zero,zero
	movss	.LCPI6_3(%rip), %xmm3           # xmm3 = mem[0],zero,zero,zero
	movl	$1024, %esi                     # imm = 0x400
	movl	$768, %edx                      # imm = 0x300
	callq	_Z38__device_stub__mandel_calc_cuda_singlePiiiffff
	jmp	.LBB6_9
.LBB6_7:                                # %select.unfold
	leaq	8(%rsp), %rdi
	movl	$3145728, %esi                  # imm = 0x300000
	movl	$1, %edx
	callq	hipMallocManaged
	leaq	16(%rsp), %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	leaq	127(%rdi), %rdx
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB6_9
# %bb.8:
	movq	8(%rsp), %rdi
	movss	.LCPI6_0(%rip), %xmm0           # xmm0 = mem[0],zero,zero,zero
	movss	.LCPI6_1(%rip), %xmm1           # xmm1 = mem[0],zero,zero,zero
	movss	.LCPI6_2(%rip), %xmm2           # xmm2 = mem[0],zero,zero,zero
	movss	.LCPI6_3(%rip), %xmm3           # xmm3 = mem[0],zero,zero,zero
	movl	$1024, %esi                     # imm = 0x400
	movl	$768, %edx                      # imm = 0x300
	callq	_Z37__device_stub__mandel_calc_cuda_multiPiiiffff
.LBB6_9:
	callq	hipDeviceSynchronize
	xorl	%ebp, %ebp
.LBB6_10:
	leaq	32(%rsp), %r14
	movq	%r14, %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	movq	(%r14), %rbx
	movq	8(%r14), %r15
	movl	$.L.str.7, %edi
	movq	%rbx, %rsi
	movq	%r15, %rdx
	xorl	%eax, %eax
	callq	printf
	movq	16(%rsp), %r14
	movq	24(%rsp), %r12
	movl	$.L.str.7, %edi
	movq	%r14, %rsi
	movq	%r12, %rdx
	xorl	%eax, %eax
	callq	printf
	movq	%r15, %rcx
	subq	%r12, %rcx
	jns	.LBB6_12
# %bb.11:                               # %.lr.ph.preheader
	cmpq	$-999999, %rcx                  # imm = 0xFFF0BDC1
	movq	$-1000000, %rax                 # imm = 0xFFF0BDC0
	cmovgeq	%rcx, %rax
	addq	%r12, %rax
	xorl	%ecx, %ecx
	cmpq	%r15, %rax
	setne	%cl
	subq	%r15, %rax
	subq	%rcx, %rax
	movabsq	$4835703278458516699, %rdx      # imm = 0x431BDE82D7B634DB
	mulq	%rdx
	shrq	$18, %rdx
	addq	%rcx, %rdx
	imulq	$1000000, %rdx, %rcx            # imm = 0xF4240
	addq	%r15, %rcx
	notq	%rdx
	addq	%rdx, %rbx
	subq	%r12, %rcx
	addq	$1000000, %rcx                  # imm = 0xF4240
.LBB6_12:                               # %._crit_edge
	subq	%r14, %rbx
	imulq	$1000000, %rbx, %rax            # imm = 0xF4240
	addq	%rcx, %rax
	xorps	%xmm0, %xmm0
	cvtsi2ss	%rax, %xmm0
	divss	.LCPI6_4(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	movl	$.L.str.8, %edi
	movb	$1, %al
	callq	printf
	testb	%bpl, %bpl
	je	.LBB6_14
# %bb.13:
	movq	8(%rsp), %rdi
	movl	$1024, %esi                     # imm = 0x400
	movl	$768, %edx                      # imm = 0x300
	callq	_Z9write_bmpPiii
	movq	8(%rsp), %rdi
	callq	free
	jmp	.LBB6_15
.LBB6_14:
	movl	$3145728, %edi                  # imm = 0x300000
	callq	malloc
	movq	%rax, %rbx
	movq	8(%rsp), %rsi
	movl	$3145728, %edx                  # imm = 0x300000
	movq	%rax, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movq	%rbx, %rdi
	movl	$1024, %esi                     # imm = 0x400
	movl	$768, %edx                      # imm = 0x300
	callq	_Z9write_bmpPiii
	movq	8(%rsp), %rdi
	callq	hipFree
.LBB6_15:
	xorl	%eax, %eax
	addq	$48, %rsp
	.cfi_def_cfa_offset 48
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.LBB6_16:
	.cfi_def_cfa_offset 96
	movq	(%rsi), %rsi
	movl	$.L.str.3, %edi
	xorl	%eax, %eax
	callq	printf
	xorl	%edi, %edi
	callq	exit
.Lfunc_end6:
	.size	main, .Lfunc_end6-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:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -16
	movq	__hip_gpubin_handle(%rip), %rbx
	testq	%rbx, %rbx
	jne	.LBB7_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB7_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z23mandel_calc_cuda_singlePiiiffff, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z22mandel_calc_cuda_multiPiiiffff, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end7:
	.size	__hip_module_ctor, .Lfunc_end7-__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	.LBB8_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
.LBB8_2:
	retq
.Lfunc_end8:
	.size	__hip_module_dtor, .Lfunc_end8-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_ZL6colors,@object              # @_ZL6colors
	.section	.rodata.str4.16,"aMS",@progbits,4
	.p2align	4, 0x0
_ZL6colors:
	.long	16711680                        # 0xff0000
	.long	15610624                        # 0xee3300
	.long	13391104                        # 0xcc5500
	.long	11162880                        # 0xaa5500
	.long	11154176                        # 0xaa3300
	.long	6710784                         # 0x666600
	.long	10066176                        # 0x999900
	.long	6723840                         # 0x669900
	.long	3381504                         # 0x339900
	.long	39338                           # 0x99aa
	.long	26282                           # 0x66aa
	.long	13226                           # 0x33aa
	.long	170                             # 0xaa
	.long	153                             # 0x99
	.long	102                             # 0x66
	.long	0                               # 0x0
	.size	_ZL6colors, 64

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

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

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"out.bmp"
	.size	.L.str, 8

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

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"Can't open file for writing."
	.size	.L.str.2, 29

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"Usage: %s <normal/cuda/cuda128>\n"
	.size	.L.str.3, 33

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"normal"
	.size	.L.str.4, 7

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

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"cuda128"
	.size	.L.str.6, 8

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	"%ld %ld\n"
	.size	.L.str.7, 9

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	"time=%f\n"
	.size	.L.str.8, 9

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z23mandel_calc_cuda_singlePiiiffff"
	.size	.L__unnamed_1, 36

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z22mandel_calc_cuda_multiPiiiffff"
	.size	.L__unnamed_2, 35

	.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 _Z38__device_stub__mandel_calc_cuda_singlePiiiffff
	.addrsig_sym _Z37__device_stub__mandel_calc_cuda_multiPiiiffff
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z23mandel_calc_cuda_singlePiiiffff
	.addrsig_sym _Z22mandel_calc_cuda_multiPiiiffff
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z23mandel_calc_cuda_singlePiiiffff ; -- Begin function _Z23mandel_calc_cuda_singlePiiiffff
	.globl	_Z23mandel_calc_cuda_singlePiiiffff
	.p2align	8
	.type	_Z23mandel_calc_cuda_singlePiiiffff,@function
_Z23mandel_calc_cuda_singlePiiiffff:    ; @_Z23mandel_calc_cuda_singlePiiiffff
; %bb.0:
	s_load_b128 s[4:7], s[0:1], 0x8
	s_waitcnt lgkmcnt(0)
	s_cmp_lt_i32 s5, 1
	s_cbranch_scc1 .LBB0_12
; %bb.1:                                ; %.preheader45.lr.ph
	s_clause 0x1
	s_load_b64 s[8:9], s[0:1], 0x18
	s_load_b64 s[2:3], s[0:1], 0x0
	v_sub_f32_e64 v14, s7, s6
	v_cvt_f32_i32_e32 v17, s4
	v_cvt_f32_i32_e32 v18, s5
	s_cmp_gt_i32 s4, 0
	v_dual_mov_b32 v15, 0 :: v_dual_mov_b32 v0, 0xff0000
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_div_scale_f32 v13, null, v17, v17, v14
	v_div_scale_f32 v22, vcc_lo, v14, v17, v14
	s_cselect_b32 s1, -1, 0
	v_rcp_f32_e32 v20, v13
	v_mov_b32_e32 v7, 0x669900
	v_mov_b32_e32 v8, 0x339900
	v_mov_b32_e32 v9, 0x99aa
	v_mov_b32_e32 v11, 0x33aa
	s_waitcnt lgkmcnt(0)
	v_sub_f32_e64 v19, s9, s8
	s_waitcnt_depctr 0xfff
	v_fma_f32 v3, -v13, v20, 1.0
	v_mov_b32_e32 v1, 0xee3300
	v_div_scale_f32 v16, null, v18, v18, v19
	v_div_scale_f32 v23, s0, v19, v18, v19
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fmac_f32_e32 v20, v3, v20
	v_rcp_f32_e32 v21, v16
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_mov_b32 v3, 0xaa5500 :: v_dual_mul_f32 v24, v22, v20
	v_fma_f32 v10, -v13, v24, v22
	s_waitcnt_depctr 0xfff
	v_fma_f32 v4, -v16, v21, 1.0
	v_mov_b32_e32 v2, 0xcc5500
	v_dual_mov_b32 v5, 0x666600 :: v_dual_fmac_f32 v24, v10, v20
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_dual_fmac_f32 v21, v4, v21 :: v_dual_mov_b32 v4, 0xaa3300
	v_mov_b32_e32 v10, 0x66aa
	v_fma_f32 v13, -v13, v24, v22
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_mul_f32_e32 v25, v23, v21
	v_cndmask_b32_e64 v22, 0, 1, s1
	s_mov_b32 s1, 0
	v_div_fmas_f32 v20, v13, v20, v24
	s_delay_alu instid0(VALU_DEP_3)
	v_fma_f32 v12, -v16, v25, v23
	v_mov_b32_e32 v6, 0x999900
	s_mov_b32 vcc_lo, s0
	v_mov_b32_e32 v13, 0x99
	v_cmp_ne_u32_e64 s0, 1, v22
	v_dual_fmac_f32 v25, v12, v21 :: v_dual_mov_b32 v12, 0xaa
	v_div_fixup_f32 v17, v20, v17, v14
	v_mov_b32_e32 v14, 0x66
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v16, -v16, v25, v23
	v_div_fmas_f32 v21, v16, v21, v25
	v_mov_b32_e32 v16, s8
	s_mov_b32 s8, 0
	s_delay_alu instid0(VALU_DEP_2)
	v_div_fixup_f32 v18, v21, v18, v19
.LBB0_2:                                ; %.preheader45
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_4 Depth 2
                                        ;       Child Loop BB0_5 Depth 3
	s_and_b32 vcc_lo, exec_lo, s0
	s_cbranch_vccnz .LBB0_11
; %bb.3:                                ; %.preheader.preheader
                                        ;   in Loop: Header=BB0_2 Depth=1
	v_mov_b32_e32 v19, s6
	s_add_i32 s7, s8, s4
.LBB0_4:                                ; %.preheader
                                        ;   Parent Loop BB0_2 Depth=1
                                        ; =>  This Loop Header: Depth=2
                                        ;       Child Loop BB0_5 Depth 3
	v_dual_mov_b32 v20, 0 :: v_dual_mov_b32 v21, 0
	s_mov_b32 s9, 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) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_mul_f32 v22, v20, v20 :: v_dual_add_f32 v23, v21, v21
	v_fma_f32 v22, v21, v21, -v22
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v20, v23, v20, v16
	v_add_f32_e32 v21, v19, v22
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v22, v21, v21
	v_fmac_f32_e32 v22, v20, v20
	s_delay_alu instid0(VALU_DEP_1)
	v_cmp_lt_f32_e32 vcc_lo, 4.0, v22
	s_cbranch_vccnz .LBB0_7
; %bb.6:                                ;   in Loop: Header=BB0_5 Depth=3
	s_add_i32 s9, s9, 1
	s_movk_i32 s10, 0x7f
	s_cmpk_eq_i32 s9, 0x7f
	s_cselect_b32 s11, -1, 0
	s_branch .LBB0_8
.LBB0_7:                                ;   in Loop: Header=BB0_5 Depth=3
	s_mov_b32 s11, -1
	s_mov_b32 s10, s9
                                        ; implicit-def: $vgpr21
                                        ; implicit-def: $vgpr20
.LBB0_8:                                ; %Flow
                                        ;   in Loop: Header=BB0_5 Depth=3
	s_and_not1_b32 vcc_lo, exec_lo, s11
	s_cbranch_vccnz .LBB0_5
; %bb.9:                                ;   in Loop: Header=BB0_4 Depth=2
	s_ashr_i32 s9, s8, 31
	s_lshr_b32 s12, s10, 3
	s_lshl_b64 s[10:11], s[8:9], 2
	s_mov_b32 m0, s12
	s_add_u32 s10, s2, s10
	v_movrels_b32_e32 v20, v0
	v_add_f32_e32 v19, v17, v19
	s_addc_u32 s11, s3, s11
	s_add_i32 s8, s8, 1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_lg_u32 s8, s7
	global_store_b32 v15, v20, s[10:11]
	s_cbranch_scc1 .LBB0_4
; %bb.10:                               ;   in Loop: Header=BB0_2 Depth=1
	s_mov_b32 s8, s7
.LBB0_11:                               ; %._crit_edge
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_add_f32_e32 v16, v18, v16
	s_add_i32 s1, s1, 1
	s_cmp_lg_u32 s1, s5
	s_cbranch_scc1 .LBB0_2
.LBB0_12:                               ; %._crit_edge55
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z23mandel_calc_cuda_singlePiiiffff
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 32
		.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 26
		.amdhsa_next_free_sgpr 13
		.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	_Z23mandel_calc_cuda_singlePiiiffff, .Lfunc_end0-_Z23mandel_calc_cuda_singlePiiiffff
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 632
; NumSgprs: 15
; NumVgprs: 26
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 1
; VGPRBlocks: 3
; NumSGPRsForWavesPerEU: 15
; NumVGPRsForWavesPerEU: 26
; Occupancy: 16
; WaveLimiterHint : 0
; 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
	.protected	_Z22mandel_calc_cuda_multiPiiiffff ; -- Begin function _Z22mandel_calc_cuda_multiPiiiffff
	.globl	_Z22mandel_calc_cuda_multiPiiiffff
	.p2align	8
	.type	_Z22mandel_calc_cuda_multiPiiiffff,@function
_Z22mandel_calc_cuda_multiPiiiffff:     ; @_Z22mandel_calc_cuda_multiPiiiffff
; %bb.0:
	s_load_b128 s[4:7], s[0:1], 0x8
	s_waitcnt lgkmcnt(0)
	s_cmpk_lt_i32 s5, 0x80
	s_cbranch_scc1 .LBB1_10
; %bb.1:                                ; %.preheader52.lr.ph
	s_clause 0x1
	s_load_b64 s[8:9], s[0:1], 0x18
	s_load_b64 s[2:3], s[0:1], 0x0
	v_sub_f32_e64 v1, s7, s6
	v_cvt_f32_i32_e32 v3, s4
	v_cvt_f32_i32_e32 v2, s5
	s_ashr_i32 s1, s5, 31
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	s_lshr_b32 s1, s1, 25
	v_div_scale_f32 v5, null, v3, v3, v1
	v_div_scale_f32 v11, vcc_lo, v1, v3, v1
	s_add_i32 s1, s5, s1
	s_delay_alu instid0(VALU_DEP_2)
	v_rcp_f32_e32 v7, v5
	s_ashr_i32 s1, s1, 7
	s_cmp_gt_i32 s4, 0
	v_mul_lo_u32 v0, s1, v0
	s_mov_b32 s5, 0
	s_waitcnt lgkmcnt(0)
	v_sub_f32_e64 v4, s9, s8
	s_waitcnt_depctr 0xfff
	v_fma_f32 v9, -v5, v7, 1.0
	v_div_scale_f32 v6, null, v2, v2, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_4) | instid1(VALU_DEP_3)
	v_rcp_f32_e32 v8, v6
	s_waitcnt_depctr 0xfff
	v_fma_f32 v10, -v6, v8, 1.0
	v_fmac_f32_e32 v7, v9, v7
	v_div_scale_f32 v9, s0, v4, v2, v4
	v_fmac_f32_e32 v8, v10, v8
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_f32_e32 v10, v11, v7
	v_mul_f32_e32 v12, v9, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v13, -v5, v10, v11
	v_fma_f32 v14, -v6, v12, v9
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v10, v13, v7
	v_fmac_f32_e32 v12, v14, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v5, -v5, v10, v11
	v_fma_f32 v6, -v6, v12, v9
	s_delay_alu instid0(VALU_DEP_2)
	v_div_fmas_f32 v5, v5, v7, v10
	s_mov_b32 vcc_lo, s0
	v_cvt_f32_i32_e32 v7, v0
	v_mul_lo_u32 v0, v0, s4
	v_div_fmas_f32 v6, v6, v8, v12
	v_div_fixup_f32 v3, v5, v3, v1
	v_mov_b32_e32 v5, 0xee3300
	s_cselect_b32 s0, -1, 0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fixup_f32 v2, v6, v2, v4
	v_fma_f32 v4, v2, v7, s8
.LBB1_2:                                ; %.preheader52
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB1_4 Depth 2
                                        ;       Child Loop BB1_5 Depth 3
	s_and_not1_b32 vcc_lo, exec_lo, s0
	s_cbranch_vccnz .LBB1_9
; %bb.3:                                ; %.preheader.preheader
                                        ;   in Loop: Header=BB1_2 Depth=1
	v_mov_b32_e32 v6, s6
	s_mov_b32 s7, 0
.LBB1_4:                                ; %.preheader
                                        ;   Parent Loop BB1_2 Depth=1
                                        ; =>  This Loop Header: Depth=2
                                        ;       Child Loop BB1_5 Depth 3
	v_mov_b32_e32 v1, 0
	v_mov_b32_e32 v7, 0
	s_mov_b32 s8, 0
	s_mov_b32 s9, 0
                                        ; implicit-def: $sgpr10
.LBB1_5:                                ;   Parent Loop BB1_2 Depth=1
                                        ;     Parent Loop BB1_4 Depth=2
                                        ; =>    This Inner Loop Header: Depth=3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_dual_mul_f32 v8, v1, v1 :: v_dual_add_f32 v9, v7, v7
	s_or_b32 s10, s10, exec_lo
	v_fma_f32 v8, v7, v7, -v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v1, v9, v1, v4
	v_add_f32_e32 v7, v6, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v8, v7, v7
	v_fmac_f32_e32 v8, v1, v1
	s_delay_alu instid0(VALU_DEP_1)
	v_cmp_nlt_f32_e32 vcc_lo, 4.0, v8
	v_mov_b32_e32 v8, s9
	s_and_saveexec_b32 s11, vcc_lo
; %bb.6:                                ;   in Loop: Header=BB1_5 Depth=3
	s_add_i32 s9, s9, 1
	v_mov_b32_e32 v8, 0x7f
	s_cmpk_eq_i32 s9, 0x7f
	s_cselect_b32 s12, -1, 0
	s_and_not1_b32 s10, s10, exec_lo
	s_and_b32 s12, s12, exec_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 s10, s10, s12
; %bb.7:                                ; %Flow
                                        ;   in Loop: Header=BB1_5 Depth=3
	s_or_b32 exec_lo, exec_lo, s11
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s11, exec_lo, s10
	s_or_b32 s8, s11, s8
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s8
	s_cbranch_execnz .LBB1_5
; %bb.8:                                ;   in Loop: Header=BB1_4 Depth=2
	s_or_b32 exec_lo, exec_lo, s8
	v_lshrrev_b32_e32 v9, 3, v8
	s_add_i32 s7, s7, 1
	v_add_f32_e32 v6, v3, v6
	s_cmp_lg_u32 s7, s4
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_cmp_eq_u32_e32 vcc_lo, 1, v9
	v_cndmask_b32_e32 v1, 0xff0000, v5, vcc_lo
	v_cmp_ne_u32_e32 vcc_lo, 2, v9
	v_cndmask_b32_e32 v1, 0xcc5500, v1, vcc_lo
	v_cmp_ne_u32_e32 vcc_lo, 3, v9
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v1, 0xaa5500, v1, vcc_lo
	v_cmp_ne_u32_e32 vcc_lo, 4, v9
	v_cndmask_b32_e32 v1, 0xaa3300, v1, vcc_lo
	v_cmp_ne_u32_e32 vcc_lo, 5, v9
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v1, 0x666600, v1, vcc_lo
	v_cmp_ne_u32_e32 vcc_lo, 6, v9
	v_cndmask_b32_e32 v1, 0x999900, v1, vcc_lo
	v_cmp_ne_u32_e32 vcc_lo, 7, v9
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v1, 0x669900, v1, vcc_lo
	v_cmp_ne_u32_e32 vcc_lo, 8, v9
	v_cndmask_b32_e32 v1, 0x339900, v1, vcc_lo
	v_cmp_ne_u32_e32 vcc_lo, 9, v9
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v1, 0x99aa, v1, vcc_lo
	v_cmp_ne_u32_e32 vcc_lo, 10, v9
	v_cndmask_b32_e32 v1, 0x66aa, v1, vcc_lo
	v_cmp_ne_u32_e32 vcc_lo, 11, v9
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v1, 0x33aa, v1, vcc_lo
	v_cmp_ne_u32_e32 vcc_lo, 12, v9
	v_cndmask_b32_e32 v7, 0xaa, v1, vcc_lo
	v_ashrrev_i32_e32 v1, 31, v0
	v_cmp_ne_u32_e32 vcc_lo, 13, v9
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cndmask_b32_e32 v10, 0x99, v7, vcc_lo
	v_lshlrev_b64 v[7:8], 2, v[0:1]
	v_cmp_ne_u32_e32 vcc_lo, 14, v9
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_dual_cndmask_b32 v1, 0x66, v10 :: v_dual_add_nc_u32 v0, 1, v0
	v_add_co_u32 v7, vcc_lo, s2, v7
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_add_co_ci_u32_e32 v8, vcc_lo, s3, v8, vcc_lo
	v_cmp_ne_u32_e32 vcc_lo, 15, v9
	v_cndmask_b32_e32 v1, 0, v1, vcc_lo
	global_store_b32 v[7:8], v1, off
	s_cbranch_scc1 .LBB1_4
.LBB1_9:                                ; %._crit_edge
                                        ;   in Loop: Header=BB1_2 Depth=1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_add_f32_e32 v4, v2, v4
	s_add_i32 s5, s5, 1
	s_cmp_lg_u32 s5, s1
	s_cbranch_scc1 .LBB1_2
.LBB1_10:                               ; %._crit_edge62
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z22mandel_calc_cuda_multiPiiiffff
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 32
		.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 15
		.amdhsa_next_free_sgpr 13
		.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_end1:
	.size	_Z22mandel_calc_cuda_multiPiiiffff, .Lfunc_end1-_Z22mandel_calc_cuda_multiPiiiffff
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 792
; NumSgprs: 15
; NumVgprs: 15
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 1
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 15
; NumVGPRsForWavesPerEU: 15
; Occupancy: 16
; WaveLimiterHint : 0
; 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
      - .offset:         8
        .size:           4
        .value_kind:     by_value
      - .offset:         12
        .size:           4
        .value_kind:     by_value
      - .offset:         16
        .size:           4
        .value_kind:     by_value
      - .offset:         20
        .size:           4
        .value_kind:     by_value
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         28
        .size:           4
        .value_kind:     by_value
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 32
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z23mandel_calc_cuda_singlePiiiffff
    .private_segment_fixed_size: 0
    .sgpr_count:     15
    .sgpr_spill_count: 0
    .symbol:         _Z23mandel_calc_cuda_singlePiiiffff.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     26
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .offset:         8
        .size:           4
        .value_kind:     by_value
      - .offset:         12
        .size:           4
        .value_kind:     by_value
      - .offset:         16
        .size:           4
        .value_kind:     by_value
      - .offset:         20
        .size:           4
        .value_kind:     by_value
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         28
        .size:           4
        .value_kind:     by_value
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 32
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z22mandel_calc_cuda_multiPiiiffff
    .private_segment_fixed_size: 0
    .sgpr_count:     15
    .sgpr_spill_count: 0
    .symbol:         _Z22mandel_calc_cuda_multiPiiiffff.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
