#include <iostream>
#include <cstdlib>
#include <stdio.h>
#include <math.h>
#include <assert.h>
#include <sys/time.h>

__global__ void fPixelGenerator(int width, int frames, unsigned char* pic) {

   int blockID = blockIdx.x;
   int tBlocks = gridDim.x;

   int fromBlocks = (blockID)     * (frames / tBlocks);
   int toBlocks   = (blockID + 1) * (frames / tBlocks);

   int threadID = threadIdx.x;
   int tThreads = blockDim.x;

   int fromThreads = (threadID)     * (width / tThreads);
   int toThreads   = (threadID + 1) * (width / tThreads);

   for (int frame = fromBlocks; frame < toBlocks; frame++) {
      for (int row = fromThreads; row < toThreads; row++) {
         for (int col = 0; col < width; col++) {
            float fx = col - 1024/2;
            float fy = row - 1024/2;
            float d = sqrtf( fx * fx + fy * fy );
            unsigned char color = (unsigned char) (160.0f + 127.0f * cos(d/10.0f - frame/7.0f) / (d/50.0f + 1.0f));

            pic[frame * width * width + row * width + col] = (unsigned char) color;
         }
      }
   }
}

static void writeBMP(const int x, const int y, const unsigned char* const bmp, const char* const name) {

   const unsigned char bmphdr[54] = {66, 77, 255, 255, 255, 255, 0, 0, 0, 0, 54, 4, 0, 0, 40, 0, 0, 0, 255, 255, 255, 255, 255, 255, 255, 255, 1, 0, 8, 0, 0, 0, 0, 0, 255, 255, 255, 255, 196, 14, 0, 0, 196, 14, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
   unsigned char hdr[1078];
   int i, j, c, xcorr, diff;
   FILE* f;

   xcorr = (x + 3) >> 2 << 2;  // BMPs have to be a multiple of 4 pixels wide
   diff = xcorr - x;

   for (i = 0; i < 54; i++) hdr[i] = bmphdr[i];
   *((int*)(&hdr[18])) = xcorr;
   *((int*)(&hdr[22])) = y;
   *((int*)(&hdr[34])) = xcorr * y;
   *((int*)(&hdr[2])) = xcorr * y + 1078;
   for (i = 0; i < 256; i++) { 
      j = i * 4 + 54;
      hdr[j+0] = i;  // blue ColorTable
      hdr[j+1] = 0;  // green
      hdr[j+2] = 0;  // red
      hdr[j+3] = 0;  // dummy
   }

   f = fopen(name, "wb");  assert(f != NULL);
   c = fwrite(hdr, 1, 1078, f);  assert(c == 1078);
   if (diff == 0) {
      c = fwrite(bmp, 1, x * y, f);  assert(c == x * y);
   }
   else {
      *((int*)(&hdr[0])) = 0;  // need up to three zero bytes
      for (j = 0; j < y; j++) {
         c = fwrite(&bmp[j * x], 1, x, f);  assert(c == x);
         c = fwrite(hdr, 1, diff, f);  assert(c == diff);
      }
   }

   fclose(f);
}

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

   // check command line
   if (argc != 5) {
      fprintf(stderr, "usage: %s <frame_width> <num_frames> <num_blocks> <num_threads>\n", argv[0]);
      exit(-1);
   }
   int width = atoi(argv[1]);
   if (width < 100) {
      fprintf(stderr, "error: <frame_width> must be at least 100\n");
      exit(-1);
   }
   int frames = atoi(argv[2]);
   if (frames < 1) {
      fprintf(stderr, "error: <num_frames> must be at least 1\n");
      exit(-1);
   }
   int blocks = atoi(argv[3]);
   if (blocks < 1) {
      fprintf(stderr, "error: <num_blocks> must be at least 1\n");
      exit(-1);
   };
   int threads = atoi(argv[4]);
   if (threads < 1) {
      fprintf(stderr, "error: <num_threads> must be at least 1\n");
      exit(-1);
   };

   if (frames % blocks != 0) {
      fprintf(stderr, "Frames not divisible by number of blocks\n");
      exit(-1);
   }

   if (width % threads != 0) {
      fprintf(stderr, "Width size not divisible by number of threads\n");
      exit(-1);
   }

   printf("computing %d frames of %d by %d picture - %d blocks w/ %d threads.\n", frames, width, width, blocks, threads);

   // allocate picture array
   unsigned char* pic = NULL;
   cudaMallocManaged(&pic, frames * width * width * sizeof(unsigned char));

   // start time
   timeval start, end;
   gettimeofday(&start, NULL);

   // Pixel Generator
   fPixelGenerator<<<blocks, threads>>>(width, frames, pic);
 
   // Wait for GPU to finish before accessing on host
   cudaDeviceSynchronize();

   // end time
   gettimeofday(&end, NULL);
   double runtime = end.tv_sec + end.tv_usec / 1000000.0 - start.tv_sec - start.tv_usec / 1000000.0;
   printf("compute time: %.4f s\n", runtime);

   // verify result by writing frames to BMP files
   if ((width <= 256) && (frames <= 100)) {
      for (int frame = 0; frame < frames; frame++) {
         char name[32];
         sprintf(name, "wave%d.bmp", frame + 1000);
         writeBMP(width, width, &pic[frame * width * width], name);
      }
   }

   // Free memory
   cudaFree(pic);

   return 0;
}
	.file	"tmpxft_0031e595_00000000-6_wavecuda2.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

	.section	.rodata.str1.1,"aMS",@progbits,1
.LC1:
	.string	"wb"
#NO_APP
	.section	.rodata
.LC0:
	.string	"BM\377\377\377\377"
	.string	""
	.string	""
	.string	""
	.string	"6\004"
	.string	""
	.string	"("
	.string	""
	.string	""
	.string	"\377\377\377\377\377\377\377\377\001"
	.string	"\b"
	.string	""
	.string	""
	.string	""
	.string	""
	.string	"\377\377\377\377\304\016"
	.string	""
	.string	"\304\016"
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.text
	.type	_ZL8writeBMPiiPKhPKc, @function
_ZL8writeBMPiiPKhPKc:
.LFB3635:
	.cfi_startproc
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movq	%rcx, %r8
	movl	$54, %ecx
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movl	%esi, %r13d
	leaq	.LC0(%rip), %rsi
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%rdx, %rbp
	xorl	%edx, %edx
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movslq	%edi, %rbx
	subq	$1176, %rsp
	.cfi_def_cfa_offset 1232
	movq	%fs:40, %rax
	movq	%rax, 1160(%rsp)
	xorl	%eax, %eax
	leaq	28(%rsp), %rdi
	leal	3(%rbx), %eax
	rep movsb
	andl	$-4, %eax
	leaq	82(%rsp), %rdi
	leaq	28(%rsp), %rsi
	movl	%eax, %r15d
	movl	$54, %ecx
	rep movsb
	movl	%eax, 100(%rsp)
	imull	%r13d, %eax
	subl	%ebx, %r15d
	movl	%r13d, 104(%rsp)
	movl	%eax, 116(%rsp)
	addl	$1078, %eax
	movl	%eax, 84(%rsp)
	leaq	82(%rsp), %rax
	movq	%rax, %r14
.L2:
	movb	%dl, 54(%rax)
	incl	%edx
	addq	$4, %rax
	movb	$0, 51(%rax)
	movb	$0, 52(%rax)
	movb	$0, 53(%rax)
	cmpl	$256, %edx
	jne	.L2
	leaq	.LC1(%rip), %rsi
	movq	%r8, %rdi
	call	fopen@PLT
	movl	$1078, %edx
	movl	$1, %esi
	movq	%r14, %rdi
	movq	%rax, %rcx
	movq	%rax, 8(%rsp)
	call	fwrite@PLT
	testl	%r15d, %r15d
	jne	.L3
	imull	%r13d, %ebx
	movq	8(%rsp), %rcx
	movl	$1, %esi
	movq	%rbp, %rdi
	movslq	%ebx, %rdx
	call	fwrite@PLT
	jmp	.L4
.L3:
	xorl	%eax, %eax
	xorl	%r12d, %r12d
	movslq	%r15d, %r15
	movl	%eax, 82(%rsp)
.L5:
	cmpl	%r13d, %r12d
	jge	.L4
	movq	8(%rsp), %rcx
	movq	%rbp, %rdi
	movq	%rbx, %rdx
	movl	$1, %esi
	incl	%r12d
	addq	%rbx, %rbp
	call	fwrite@PLT
	movq	8(%rsp), %rcx
	movq	%r15, %rdx
	movq	%r14, %rdi
	movl	$1, %esi
	call	fwrite@PLT
	jmp	.L5
.L4:
	movq	8(%rsp), %rdi
	call	fclose@PLT
	movq	1160(%rsp), %rax
	subq	%fs:40, %rax
	je	.L7
	call	__stack_chk_fail@PLT
.L7:
	addq	$1176, %rsp
	.cfi_def_cfa_offset 56
	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
	ret
	.cfi_endproc
.LFE3635:
	.size	_ZL8writeBMPiiPKhPKc, .-_ZL8writeBMPiiPKhPKc
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB3639:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE3639:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z37__device_stub__Z15fPixelGeneratoriiPhiiPh
	.type	_Z37__device_stub__Z15fPixelGeneratoriiPhiiPh, @function
_Z37__device_stub__Z15fPixelGeneratoriiPhiiPh:
.LFB3661:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movl	%edi, 12(%rsp)
	leaq	24(%rsp), %rcx
	leaq	32(%rsp), %rdi
	movl	%esi, 8(%rsp)
	leaq	44(%rsp), %rsi
	movq	%rdx, (%rsp)
	leaq	16(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	12(%rsp), %rax
	movl	$1, 40(%rsp)
	movq	%rax, 80(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 32(%rsp)
	movq	%rax, 44(%rsp)
	movl	$1, 52(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L12
	pushq	24(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z15fPixelGeneratoriiPh(%rip), %rdi
	pushq	24(%rsp)
	.cfi_def_cfa_offset 144
	movq	60(%rsp), %rcx
	movl	68(%rsp), %r8d
	movq	48(%rsp), %rsi
	movl	56(%rsp), %edx
	leaq	96(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L12:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L14
	call	__stack_chk_fail@PLT
.L14:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3661:
	.size	_Z37__device_stub__Z15fPixelGeneratoriiPhiiPh, .-_Z37__device_stub__Z15fPixelGeneratoriiPhiiPh
	.globl	_Z15fPixelGeneratoriiPh
	.type	_Z15fPixelGeneratoriiPh, @function
_Z15fPixelGeneratoriiPh:
.LFB3662:
	.cfi_startproc
	endbr64
	jmp	_Z37__device_stub__Z15fPixelGeneratoriiPhiiPh
	.cfi_endproc
.LFE3662:
	.size	_Z15fPixelGeneratoriiPh, .-_Z15fPixelGeneratoriiPh
	.section	.rodata.str1.1
.LC2:
	.string	"usage: %s <frame_width> <num_frames> <num_blocks> <num_threads>\n"
.LC3:
	.string	"error: <frame_width> must be at least 100\n"
.LC4:
	.string	"error: <num_frames> must be at least 1\n"
.LC5:
	.string	"error: <num_blocks> must be at least 1\n"
.LC6:
	.string	"error: <num_threads> must be at least 1\n"
.LC7:
	.string	"Frames not divisible by number of blocks\n"
.LC8:
	.string	"Width size not divisible by number of threads\n"
.LC9:
	.string	"computing %d frames of %d by %d picture - %d blocks w/ %d threads.\n"
.LC11:
	.string	"compute time: %.4f s\n"
.LC12:
	.string	"wave%d.bmp"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB3636:
	.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
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%rsi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$120, %rsp
	.cfi_def_cfa_offset 176
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	cmpl	$5, %edi
	je	.L18
	movq	(%rsi), %rcx
	movq	stderr(%rip), %rdi
	leaq	.LC2(%rip), %rdx
	movl	$2, %esi
	call	__fprintf_chk@PLT
	jmp	.L37
.L18:
	movq	8(%rsi), %rdi
	call	atoi@PLT
	movl	%eax, %ebx
	cmpl	$99, %eax
	jg	.L19
	leaq	.LC3(%rip), %rdx
.L38:
	movq	stderr(%rip), %rdi
	movl	$2, %esi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
.L37:
	orl	$-1, %edi
	call	exit@PLT
.L19:
	movq	16(%r12), %rdi
	call	atoi@PLT
	leaq	.LC4(%rip), %rdx
	movl	%eax, %ebp
	testl	%eax, %eax
	jle	.L38
	movq	24(%r12), %rdi
	call	atoi@PLT
	leaq	.LC5(%rip), %rdx
	movl	%eax, %r13d
	testl	%eax, %eax
	jle	.L38
	movq	32(%r12), %rdi
	call	atoi@PLT
	leaq	.LC6(%rip), %rdx
	movl	%eax, %r14d
	testl	%eax, %eax
	jle	.L38
	movl	%ebp, %eax
	cltd
	idivl	%r13d
	testl	%edx, %edx
	je	.L23
	leaq	.LC7(%rip), %rdx
	jmp	.L38
.L23:
	movl	%ebx, %eax
	cltd
	idivl	%r14d
	movl	%edx, %r12d
	leaq	.LC8(%rip), %rdx
	testl	%r12d, %r12d
	jne	.L38
	pushq	%rax
	.cfi_def_cfa_offset 184
	movl	%r13d, %r9d
	movl	%ebx, %r8d
	movl	%ebx, %ecx
	pushq	%r14
	.cfi_def_cfa_offset 192
	movl	%ebp, %edx
	leaq	.LC9(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	xorl	%edx, %edx
	movq	%rdx, 32(%rsp)
	popq	%rcx
	.cfi_def_cfa_offset 184
	movl	$1, %edx
	popq	%rsi
	.cfi_def_cfa_offset 176
	movl	%ebx, %esi
	imull	%ebp, %esi
	leaq	16(%rsp), %rdi
	imull	%ebx, %esi
	movslq	%esi, %rsi
	call	cudaMallocManaged@PLT
	xorl	%esi, %esi
	leaq	40(%rsp), %rdi
	call	gettimeofday@PLT
	movl	%r14d, 56(%rsp)
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movabsq	$4294967297, %rax
	movl	%r13d, 28(%rsp)
	movq	%rax, 60(%rsp)
	movl	64(%rsp), %ecx
	movq	%rax, 32(%rsp)
	movq	56(%rsp), %rdx
	movq	28(%rsp), %rdi
	movl	36(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L25
	movq	16(%rsp), %rdx
	movl	%ebp, %esi
	movl	%ebx, %edi
	call	_Z37__device_stub__Z15fPixelGeneratoriiPhiiPh
.L25:
	call	cudaDeviceSynchronize@PLT
	xorl	%esi, %esi
	leaq	56(%rsp), %rdi
	call	gettimeofday@PLT
	cvtsi2sdq	64(%rsp), %xmm0
	movb	$1, %al
	movsd	.LC10(%rip), %xmm2
	cvtsi2sdq	56(%rsp), %xmm1
	leaq	.LC11(%rip), %rsi
	movl	$2, %edi
	divsd	%xmm2, %xmm0
	addsd	%xmm1, %xmm0
	cvtsi2sdq	40(%rsp), %xmm1
	subsd	%xmm1, %xmm0
	cvtsi2sdq	48(%rsp), %xmm1
	divsd	%xmm2, %xmm1
	subsd	%xmm1, %xmm0
	call	__printf_chk@PLT
	cmpl	$256, %ebx
	jg	.L26
	cmpl	$100, %ebp
	jg	.L26
	movl	%ebx, %r9d
	xorl	%r13d, %r13d
	leaq	.LC12(%rip), %r15
	imull	%ebx, %r9d
	movslq	%r9d, %r14
.L27:
	leaq	72(%rsp), %rax
	leal	1000(%r12), %r8d
	movq	%r15, %rcx
	incl	%r12d
	movq	%rax, %rdi
	movq	%rax, 8(%rsp)
	movl	$32, %edx
	xorl	%eax, %eax
	movl	$2, %esi
	call	__sprintf_chk@PLT
	movq	16(%rsp), %rdx
	movl	%ebx, %esi
	movl	%ebx, %edi
	movq	8(%rsp), %rcx
	addq	%r13, %rdx
	addq	%r14, %r13
	call	_ZL8writeBMPiiPKhPKc
	cmpl	%r12d, %ebp
	jne	.L27
.L26:
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L28
	call	__stack_chk_fail@PLT
.L28:
	addq	$120, %rsp
	.cfi_def_cfa_offset 56
	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
	ret
	.cfi_endproc
.LFE3636:
	.size	main, .-main
	.section	.rodata.str1.1
.LC13:
	.string	"_Z15fPixelGeneratoriiPh"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB3664:
	.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	.LC13(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z15fPixelGeneratoriiPh(%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
.LFE3664:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC10:
	.long	0
	.long	1093567616
	.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 <iostream>
#include <cstdlib>
#include <stdio.h>
#include <math.h>
#include <assert.h>
#include <sys/time.h>

__global__ void fPixelGenerator(int width, int frames, unsigned char* pic) {

   int blockID = blockIdx.x;
   int tBlocks = gridDim.x;

   int fromBlocks = (blockID)     * (frames / tBlocks);
   int toBlocks   = (blockID + 1) * (frames / tBlocks);

   int threadID = threadIdx.x;
   int tThreads = blockDim.x;

   int fromThreads = (threadID)     * (width / tThreads);
   int toThreads   = (threadID + 1) * (width / tThreads);

   for (int frame = fromBlocks; frame < toBlocks; frame++) {
      for (int row = fromThreads; row < toThreads; row++) {
         for (int col = 0; col < width; col++) {
            float fx = col - 1024/2;
            float fy = row - 1024/2;
            float d = sqrtf( fx * fx + fy * fy );
            unsigned char color = (unsigned char) (160.0f + 127.0f * cos(d/10.0f - frame/7.0f) / (d/50.0f + 1.0f));

            pic[frame * width * width + row * width + col] = (unsigned char) color;
         }
      }
   }
}

static void writeBMP(const int x, const int y, const unsigned char* const bmp, const char* const name) {

   const unsigned char bmphdr[54] = {66, 77, 255, 255, 255, 255, 0, 0, 0, 0, 54, 4, 0, 0, 40, 0, 0, 0, 255, 255, 255, 255, 255, 255, 255, 255, 1, 0, 8, 0, 0, 0, 0, 0, 255, 255, 255, 255, 196, 14, 0, 0, 196, 14, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
   unsigned char hdr[1078];
   int i, j, c, xcorr, diff;
   FILE* f;

   xcorr = (x + 3) >> 2 << 2;  // BMPs have to be a multiple of 4 pixels wide
   diff = xcorr - x;

   for (i = 0; i < 54; i++) hdr[i] = bmphdr[i];
   *((int*)(&hdr[18])) = xcorr;
   *((int*)(&hdr[22])) = y;
   *((int*)(&hdr[34])) = xcorr * y;
   *((int*)(&hdr[2])) = xcorr * y + 1078;
   for (i = 0; i < 256; i++) { 
      j = i * 4 + 54;
      hdr[j+0] = i;  // blue ColorTable
      hdr[j+1] = 0;  // green
      hdr[j+2] = 0;  // red
      hdr[j+3] = 0;  // dummy
   }

   f = fopen(name, "wb");  assert(f != NULL);
   c = fwrite(hdr, 1, 1078, f);  assert(c == 1078);
   if (diff == 0) {
      c = fwrite(bmp, 1, x * y, f);  assert(c == x * y);
   }
   else {
      *((int*)(&hdr[0])) = 0;  // need up to three zero bytes
      for (j = 0; j < y; j++) {
         c = fwrite(&bmp[j * x], 1, x, f);  assert(c == x);
         c = fwrite(hdr, 1, diff, f);  assert(c == diff);
      }
   }

   fclose(f);
}

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

   // check command line
   if (argc != 5) {
      fprintf(stderr, "usage: %s <frame_width> <num_frames> <num_blocks> <num_threads>\n", argv[0]);
      exit(-1);
   }
   int width = atoi(argv[1]);
   if (width < 100) {
      fprintf(stderr, "error: <frame_width> must be at least 100\n");
      exit(-1);
   }
   int frames = atoi(argv[2]);
   if (frames < 1) {
      fprintf(stderr, "error: <num_frames> must be at least 1\n");
      exit(-1);
   }
   int blocks = atoi(argv[3]);
   if (blocks < 1) {
      fprintf(stderr, "error: <num_blocks> must be at least 1\n");
      exit(-1);
   };
   int threads = atoi(argv[4]);
   if (threads < 1) {
      fprintf(stderr, "error: <num_threads> must be at least 1\n");
      exit(-1);
   };

   if (frames % blocks != 0) {
      fprintf(stderr, "Frames not divisible by number of blocks\n");
      exit(-1);
   }

   if (width % threads != 0) {
      fprintf(stderr, "Width size not divisible by number of threads\n");
      exit(-1);
   }

   printf("computing %d frames of %d by %d picture - %d blocks w/ %d threads.\n", frames, width, width, blocks, threads);

   // allocate picture array
   unsigned char* pic = NULL;
   hipMallocManaged(&pic, frames * width * width * sizeof(unsigned char));

   // start time
   timeval start, end;
   gettimeofday(&start, NULL);

   // Pixel Generator
   fPixelGenerator<<<blocks, threads>>>(width, frames, pic);
 
   // Wait for GPU to finish before accessing on host
   hipDeviceSynchronize();

   // end time
   gettimeofday(&end, NULL);
   double runtime = end.tv_sec + end.tv_usec / 1000000.0 - start.tv_sec - start.tv_usec / 1000000.0;
   printf("compute time: %.4f s\n", runtime);

   // verify result by writing frames to BMP files
   if ((width <= 256) && (frames <= 100)) {
      for (int frame = 0; frame < frames; frame++) {
         char name[32];
         sprintf(name, "wave%d.bmp", frame + 1000);
         writeBMP(width, width, &pic[frame * width * width], name);
      }
   }

   // Free memory
   hipFree(pic);

   return 0;
}
	.text
	.file	"wavecuda2.hip"
                                        # Start of file scope inline assembly
	.globl	_ZSt21ios_base_library_initv

                                        # End of file scope inline assembly
	.globl	_Z30__device_stub__fPixelGeneratoriiPh # -- Begin function _Z30__device_stub__fPixelGeneratoriiPh
	.type	_Z30__device_stub__fPixelGeneratoriiPh,@function
_Z30__device_stub__fPixelGeneratoriiPh: # @_Z30__device_stub__fPixelGeneratoriiPh
	.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	$96, %rsp
	.cfi_def_cfa_offset 144
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	4(%rsp), %rax
	movl	%edi, (%rax)
	movq	%rsp, %rcx
	movl	%esi, (%rcx)
	leaq	24(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	64(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	movq	%rsi, 16(%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z15fPixelGeneratoriiPh, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$112, %rsp
	.cfi_adjust_cfa_offset -112
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z30__device_stub__fPixelGeneratoriiPh, .Lfunc_end0-_Z30__device_stub__fPixelGeneratoriiPh
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function main
.LCPI1_0:
	.quad	0x412e848000000000              # double 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	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$1224, %rsp                     # imm = 0x4C8
	.cfi_def_cfa_offset 1280
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%rsi, %r15
	cmpl	$5, %edi
	jne	.LBB1_20
# %bb.1:
	movq	8(%r15), %rdi
	callq	atoi
                                        # kill: def $eax killed $eax def $rax
	cmpl	$99, %eax
	jle	.LBB1_21
# %bb.2:
	movq	%rax, 24(%rsp)                  # 8-byte Spill
	movq	16(%r15), %rdi
	callq	atoi
	testl	%eax, %eax
	jle	.LBB1_22
# %bb.3:
	movl	%eax, %ebp
	movq	24(%r15), %rdi
	callq	atoi
	testl	%eax, %eax
	jle	.LBB1_23
# %bb.4:
	movl	%eax, %r14d
	movq	32(%r15), %rdi
	callq	atoi
	testl	%eax, %eax
	jle	.LBB1_25
# %bb.5:
	movl	%eax, %r15d
	movl	%ebp, %eax
	xorl	%edx, %edx
	divl	%r14d
	testl	%edx, %edx
	jne	.LBB1_26
# %bb.6:
	movq	24(%rsp), %rbx                  # 8-byte Reload
	movl	%ebx, %eax
	xorl	%edx, %edx
	divl	%r15d
	testl	%edx, %edx
	jne	.LBB1_27
# %bb.7:
	movl	$.L.str.7, %edi
	movl	%ebp, %esi
	movl	%ebx, %edx
	movl	%ebx, %ecx
	movl	%r14d, %r8d
	movl	%r15d, %r9d
	xorl	%eax, %eax
	callq	printf
	leaq	8(%rsp), %rdi
	movq	$0, (%rdi)
	movl	%ebx, %r12d
	imull	%ebx, %r12d
	movl	%r12d, %esi
	imull	%ebp, %esi
	movl	$1, %edx
	callq	hipMallocManaged
	leaq	80(%rsp), %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	movl	%r14d, %edi
	btsq	$32, %rdi
	movl	%r15d, %edx
	btsq	$32, %rdx
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB1_9
# %bb.8:
	movq	8(%rsp), %rdx
	movl	%ebx, %edi
	movl	%ebp, %esi
	callq	_Z30__device_stub__fPixelGeneratoriiPh
.LBB1_9:
	callq	hipDeviceSynchronize
	leaq	128(%rsp), %r14
	movq	%r14, %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	cvtsi2sdq	(%r14), %xmm1
	cvtsi2sdq	8(%r14), %xmm0
	movsd	.LCPI1_0(%rip), %xmm2           # xmm2 = mem[0],zero
	divsd	%xmm2, %xmm0
	addsd	%xmm1, %xmm0
	xorps	%xmm1, %xmm1
	cvtsi2sdq	80(%rsp), %xmm1
	subsd	%xmm1, %xmm0
	xorps	%xmm1, %xmm1
	cvtsi2sdq	88(%rsp), %xmm1
	divsd	%xmm2, %xmm1
	subsd	%xmm1, %xmm0
	movl	$.L.str.8, %edi
	movb	$1, %al
	callq	printf
	cmpl	$256, %ebx                      # imm = 0x100
	ja	.LBB1_19
# %bb.10:
	cmpl	$100, %ebp
	ja	.LBB1_19
# %bb.11:                               # %.lr.ph
	movq	%rbx, %r14
	leal	3(%rbx), %eax
	andl	$2147483644, %eax               # imm = 0x7FFFFFFC
	movl	%eax, %ecx
	imull	%r14d, %ecx
	movq	%rcx, 48(%rsp)                  # 8-byte Spill
	addl	$1078, %ecx                     # imm = 0x436
	movl	%ecx, 16(%rsp)                  # 4-byte Spill
	movl	%eax, 4(%rsp)                   # 4-byte Spill
	subl	%r14d, %eax
	movl	%r14d, %ebx
	movslq	%eax, %r15
	movl	%r12d, %eax
	movq	%rax, 32(%rsp)                  # 8-byte Spill
	movl	%ebp, %eax
	movq	%rax, 40(%rsp)                  # 8-byte Spill
	xorl	%ebp, %ebp
	xorl	%r13d, %r13d
	movl	%r12d, 20(%rsp)                 # 4-byte Spill
.LBB1_12:                               # =>This Loop Header: Depth=1
                                        #     Child Loop BB1_13 Depth 2
                                        #     Child Loop BB1_17 Depth 2
	movl	%ebp, %ebp
	leal	1000(%r13), %edx
	movl	$.L.str.9, %esi
	leaq	96(%rsp), %rdi
	xorl	%eax, %eax
	callq	sprintf
	movl	%r12d, %eax
	movq	8(%rsp), %r12
                                        # kill: def $eax killed $eax def $rax
	movq	%r13, 64(%rsp)                  # 8-byte Spill
	imull	%r13d, %eax
	movq	%rax, 56(%rsp)                  # 8-byte Spill
	leaq	188(%rsp), %rax
	movw	$0, 8(%rax)
	movq	$0, (%rax)
	movw	$19778, 144(%rsp)               # imm = 0x4D42
	xorl	%ecx, %ecx
	movl	%ecx, 150(%rsp)
	movabsq	$171798692918, %rdx             # imm = 0x2800000436
	movq	%rdx, 154(%rsp)
	movw	$1, 170(%rsp)
	movb	$8, 172(%rsp)
	movl	%ecx, -15(%rax)
	movb	$0, -11(%rax)
	movl	$3780, 182(%rsp)                # imm = 0xEC4
	movw	$3780, 186(%rsp)                # imm = 0xEC4
	movl	4(%rsp), %eax                   # 4-byte Reload
	movl	%eax, 162(%rsp)
	movl	%r14d, 166(%rsp)
	movq	48(%rsp), %rax                  # 8-byte Reload
	movl	%eax, 178(%rsp)
	movl	16(%rsp), %eax                  # 4-byte Reload
	movl	%eax, 146(%rsp)
	xorl	%eax, %eax
.LBB1_13:                               #   Parent Loop BB1_12 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movb	%al, 198(%rsp,%rax,4)
	movw	$0, 199(%rsp,%rax,4)
	movb	$0, 201(%rsp,%rax,4)
	incq	%rax
	cmpq	$256, %rax                      # imm = 0x100
	jne	.LBB1_13
# %bb.14:                               #   in Loop: Header=BB1_12 Depth=1
	movl	$.L.str.10, %esi
	leaq	96(%rsp), %rdi
	callq	fopen
	movq	%rax, %r13
	movl	$1, %esi
	movl	$1078, %edx                     # imm = 0x436
	leaq	144(%rsp), %rdi
	movq	%rax, %rcx
	callq	fwrite
	cmpl	%r14d, 4(%rsp)                  # 4-byte Folded Reload
	movq	%rbp, 72(%rsp)                  # 8-byte Spill
	jne	.LBB1_16
# %bb.15:                               #   in Loop: Header=BB1_12 Depth=1
	addq	56(%rsp), %r12                  # 8-byte Folded Reload
	movl	$1, %esi
	movq	%r12, %rdi
	movq	32(%rsp), %rdx                  # 8-byte Reload
	movq	%r13, %rcx
	callq	fwrite
	jmp	.LBB1_18
.LBB1_16:                               # %.lr.ph.i
                                        #   in Loop: Header=BB1_12 Depth=1
	movl	$0, 144(%rsp)
	addq	%rbp, %r12
	movq	%rbx, %r14
	leaq	144(%rsp), %rbp
.LBB1_17:                               #   Parent Loop BB1_12 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movl	$1, %esi
	movq	%r12, %rdi
	movq	%rbx, %rdx
	movq	%r13, %rcx
	callq	fwrite
	movl	$1, %esi
	movq	%rbp, %rdi
	movq	%r15, %rdx
	movq	%r13, %rcx
	callq	fwrite
	addq	%rbx, %r12
	decq	%r14
	jne	.LBB1_17
.LBB1_18:                               # %_ZL8writeBMPiiPKhPKc.exit
                                        #   in Loop: Header=BB1_12 Depth=1
	movq	%r13, %rdi
	callq	fclose
	movq	64(%rsp), %r13                  # 8-byte Reload
	incq	%r13
	movl	20(%rsp), %r12d                 # 4-byte Reload
	movq	72(%rsp), %rbp                  # 8-byte Reload
	addl	%r12d, %ebp
	cmpq	40(%rsp), %r13                  # 8-byte Folded Reload
	movq	24(%rsp), %r14                  # 8-byte Reload
	jne	.LBB1_12
.LBB1_19:                               # %.loopexit
	movq	8(%rsp), %rdi
	callq	hipFree
	xorl	%eax, %eax
	addq	$1224, %rsp                     # imm = 0x4C8
	.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
	retq
.LBB1_20:
	.cfi_def_cfa_offset 1280
	movq	stderr(%rip), %rdi
	movq	(%r15), %rdx
	movl	$.L.str, %esi
	xorl	%eax, %eax
	callq	fprintf
	jmp	.LBB1_29
.LBB1_21:
	movq	stderr(%rip), %rcx
	movl	$.L.str.1, %edi
	movl	$42, %esi
	jmp	.LBB1_28
.LBB1_22:
	movq	stderr(%rip), %rcx
	movl	$.L.str.2, %edi
	jmp	.LBB1_24
.LBB1_23:
	movq	stderr(%rip), %rcx
	movl	$.L.str.3, %edi
.LBB1_24:
	movl	$39, %esi
	jmp	.LBB1_28
.LBB1_25:
	movq	stderr(%rip), %rcx
	movl	$.L.str.4, %edi
	movl	$40, %esi
	jmp	.LBB1_28
.LBB1_26:
	movq	stderr(%rip), %rcx
	movl	$.L.str.5, %edi
	movl	$41, %esi
	jmp	.LBB1_28
.LBB1_27:
	movq	stderr(%rip), %rcx
	movl	$.L.str.6, %edi
	movl	$46, %esi
.LBB1_28:
	movl	$1, %edx
	callq	fwrite
.LBB1_29:
	movl	$-1, %edi
	callq	exit
.Lfunc_end1:
	.size	main, .Lfunc_end1-main
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	jne	.LBB2_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB2_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z15fPixelGeneratoriiPh, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end2:
	.size	__hip_module_ctor, .Lfunc_end2-__hip_module_ctor
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_dtor
	.type	__hip_module_dtor,@function
__hip_module_dtor:                      # @__hip_module_dtor
	.cfi_startproc
# %bb.0:
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB3_2
# %bb.1:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__hipUnregisterFatBinary
	movq	$0, __hip_gpubin_handle(%rip)
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB3_2:
	retq
.Lfunc_end3:
	.size	__hip_module_dtor, .Lfunc_end3-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z15fPixelGeneratoriiPh,@object # @_Z15fPixelGeneratoriiPh
	.section	.rodata,"a",@progbits
	.globl	_Z15fPixelGeneratoriiPh
	.p2align	3, 0x0
_Z15fPixelGeneratoriiPh:
	.quad	_Z30__device_stub__fPixelGeneratoriiPh
	.size	_Z15fPixelGeneratoriiPh, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"usage: %s <frame_width> <num_frames> <num_blocks> <num_threads>\n"
	.size	.L.str, 65

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"error: <frame_width> must be at least 100\n"
	.size	.L.str.1, 43

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"error: <num_frames> must be at least 1\n"
	.size	.L.str.2, 40

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"error: <num_blocks> must be at least 1\n"
	.size	.L.str.3, 40

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"error: <num_threads> must be at least 1\n"
	.size	.L.str.4, 41

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"Frames not divisible by number of blocks\n"
	.size	.L.str.5, 42

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"Width size not divisible by number of threads\n"
	.size	.L.str.6, 47

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	"computing %d frames of %d by %d picture - %d blocks w/ %d threads.\n"
	.size	.L.str.7, 68

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	"compute time: %.4f s\n"
	.size	.L.str.8, 22

	.type	.L.str.9,@object                # @.str.9
.L.str.9:
	.asciz	"wave%d.bmp"
	.size	.L.str.9, 11

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

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z15fPixelGeneratoriiPh"
	.size	.L__unnamed_1, 24

	.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 _Z30__device_stub__fPixelGeneratoriiPh
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z15fPixelGeneratoriiPh
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z15fPixelGeneratoriiPh ; -- Begin function _Z15fPixelGeneratoriiPh
	.globl	_Z15fPixelGeneratoriiPh
	.p2align	8
	.type	_Z15fPixelGeneratoriiPh,@function
_Z15fPixelGeneratoriiPh:                ; @_Z15fPixelGeneratoriiPh
; %bb.0:
	s_load_b32 s4, s[0:1], 0x10
	s_add_u32 s2, s0, 16
	s_addc_u32 s3, s1, 0
	s_waitcnt lgkmcnt(0)
	s_ashr_i32 s6, s4, 31
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_add_i32 s4, s4, s6
	s_xor_b32 s7, s4, s6
	s_load_b64 s[4:5], s[0:1], 0x0
	v_cvt_f32_u32_e32 v1, s7
	s_sub_i32 s9, 0, s7
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_4) | instid1(VALU_DEP_1)
	v_rcp_iflag_f32_e32 v1, v1
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v1, 0x4f7ffffe, v1
	s_waitcnt lgkmcnt(0)
	s_ashr_i32 s10, s5, 31
	v_cvt_u32_f32_e32 v1, v1
	s_add_i32 s5, s5, s10
	s_xor_b32 s6, s10, s6
	s_xor_b32 s5, s5, s10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_readfirstlane_b32 s8, v1
	s_mul_i32 s9, s9, s8
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_mul_hi_u32 s9, s8, s9
	s_add_i32 s8, s8, s9
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_mul_hi_u32 s8, s5, s8
	s_mul_i32 s9, s8, s7
	s_delay_alu instid0(SALU_CYCLE_1)
	s_sub_i32 s5, s5, s9
	s_add_i32 s9, s8, 1
	s_sub_i32 s10, s5, s7
	s_cmp_ge_u32 s5, s7
	s_cselect_b32 s8, s9, s8
	s_cselect_b32 s5, s10, s5
	s_add_i32 s9, s8, 1
	s_cmp_ge_u32 s5, s7
	s_cselect_b32 s5, s9, s8
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_xor_b32 s5, s5, s6
	s_sub_i32 s9, s5, s6
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_mul_i32 s5, s9, s15
	s_add_i32 s8, s5, s9
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_ge_i32 s5, s8
	s_cbranch_scc1 .LBB0_13
; %bb.1:                                ; %.preheader44.lr.ph
	s_load_b32 s2, s[2:3], 0xc
	s_ashr_i32 s7, s4, 31
	s_mul_i32 s15, s15, s4
	s_add_i32 s10, s4, s7
	s_mov_b32 s12, 0x37d75334
	s_xor_b32 s10, s10, s7
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cvt_f32_u32_e32 v1, s2
	s_sub_i32 s6, 0, s2
	v_rcp_iflag_f32_e32 v1, v1
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v1, 0x4f7ffffe, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_u32_f32_e32 v1, v1
	v_readfirstlane_b32 s3, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_mul_i32 s6, s6, s3
	s_mul_hi_u32 s6, s3, s6
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_add_i32 s3, s3, s6
	s_mul_hi_u32 s3, s10, s3
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_mul_i32 s6, s3, s2
	s_sub_i32 s6, s10, s6
	s_add_i32 s10, s3, 1
	s_sub_i32 s11, s6, s2
	s_cmp_ge_u32 s6, s2
	s_cselect_b32 s3, s10, s3
	s_cselect_b32 s6, s11, s6
	s_add_i32 s10, s3, 1
	s_cmp_ge_u32 s6, s2
	s_mov_b32 s11, 0xb94c1982
	s_cselect_b32 s2, s10, s3
	s_mul_i32 s3, s4, s4
	s_xor_b32 s2, s2, s7
	s_mov_b32 s10, 0x7fffff
	s_sub_i32 s2, s2, s7
	s_load_b64 s[6:7], s[0:1], 0x8
	v_mul_lo_u32 v0, s2, v0
	s_cmp_gt_i32 s4, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s9, v[0:1]
	v_dual_mov_b32 v2, 0 :: v_dual_add_nc_u32 v3, s2, v0
	s_cselect_b32 s9, -1, 0
	v_cmp_lt_i32_e64 s0, v0, v3
	s_delay_alu instid0(VALU_DEP_3)
	v_mul_lo_u32 v4, s4, v1
.LBB0_2:                                ; %.preheader44
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_4 Depth 2
                                        ;       Child Loop BB0_6 Depth 3
	s_delay_alu instid0(VALU_DEP_2)
	s_and_saveexec_b32 s13, s0
	s_cbranch_execz .LBB0_12
; %bb.3:                                ; %.preheader.lr.ph
                                        ;   in Loop: Header=BB0_2 Depth=1
	v_cvt_f32_i32_e32 v1, s5
	s_mov_b32 s14, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_div_scale_f32 v5, null, 0x40e00000, 0x40e00000, v1
	v_div_scale_f32 v8, vcc_lo, v1, 0x40e00000, v1
	v_rcp_f32_e32 v6, v5
	s_waitcnt_depctr 0xfff
	v_fma_f32 v7, -v5, v6, 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v6, v7, v6
	v_mul_f32_e32 v7, v8, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v9, -v5, v7, v8
	v_fmac_f32_e32 v7, v9, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v5, -v5, v7, v8
	v_div_fmas_f32 v6, v5, v6, v7
	v_mov_b32_e32 v5, v4
	v_mov_b32_e32 v7, v0
	s_delay_alu instid0(VALU_DEP_3)
	v_div_fixup_f32 v6, v6, 0x40e00000, v1
.LBB0_4:                                ; %.preheader
                                        ;   Parent Loop BB0_2 Depth=1
                                        ; =>  This Loop Header: Depth=2
                                        ;       Child Loop BB0_6 Depth 3
	s_and_not1_b32 vcc_lo, exec_lo, s9
	s_cbranch_vccnz .LBB0_11
; %bb.5:                                ; %.lr.ph
                                        ;   in Loop: Header=BB0_4 Depth=2
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v1, 0xfffffe00, v7
	s_mov_b32 s15, 0
	v_cvt_f32_i32_e32 v1, v1
	s_delay_alu instid0(VALU_DEP_1)
	v_mul_f32_e32 v8, v1, v1
.LBB0_6:                                ;   Parent Loop BB0_2 Depth=1
                                        ;     Parent Loop BB0_4 Depth=2
                                        ; =>    This Inner Loop Header: Depth=3
	s_add_i32 s1, s15, 0xfffffe00
	s_mov_b32 s2, exec_lo
	v_cvt_f32_i32_e32 v1, s1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v1, v1, v1, v8
	v_mul_f32_e32 v9, 0x4f800000, v1
	v_cmp_gt_f32_e32 vcc_lo, 0xf800000, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v1, v1, v9, vcc_lo
	v_sqrt_f32_e32 v9, v1
	s_waitcnt_depctr 0xfff
	v_add_nc_u32_e32 v10, -1, v9
	v_add_nc_u32_e32 v11, 1, v9
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v12, -v10, v9, v1
	v_fma_f32 v13, -v11, v9, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_ge_f32_e64 s1, 0, v12
	v_cndmask_b32_e64 v9, v9, v10, s1
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_lt_f32_e64 s1, 0, v13
	v_cndmask_b32_e64 v9, v9, v11, s1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v10, 0x37800000, v9
	v_cndmask_b32_e32 v9, v9, v10, vcc_lo
	v_cmp_class_f32_e64 vcc_lo, v1, 0x260
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v9, v9, v1, vcc_lo
	v_div_scale_f32 v1, null, 0x41200000, 0x41200000, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f32_e32 v10, v1
	s_waitcnt_depctr 0xfff
	v_fma_f32 v11, -v1, v10, 1.0
	v_fmac_f32_e32 v10, v11, v10
	v_div_scale_f32 v11, vcc_lo, v9, 0x41200000, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v12, v11, v10
	v_fma_f32 v13, -v1, v12, v11
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v12, v13, v10
	v_fma_f32 v1, -v1, v12, v11
                                        ; implicit-def: $vgpr11
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fmas_f32 v1, v1, v10, v12
	v_div_fixup_f32 v1, v1, 0x41200000, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_f32_e32 v10, v1, v6
                                        ; implicit-def: $vgpr1
	v_cmpx_ngt_f32_e64 0x48000000, |v10|
	s_xor_b32 s16, exec_lo, s2
	s_cbranch_execz .LBB0_8
; %bb.7:                                ;   in Loop: Header=BB0_6 Depth=3
	v_and_b32_e32 v16, 0x7fffffff, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_and_or_b32 v19, v16, s10, 0x800000
	v_mad_u64_u32 v[11:12], null, 0xfe5163ab, v19, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mov_b32_e32 v1, v12
	v_lshrrev_b32_e32 v16, 23, v16
	v_mad_u64_u32 v[12:13], null, 0x3c439041, v19, v[1:2]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v17, 0xffffff88, v16
	v_cmp_lt_u32_e32 vcc_lo, 63, v17
	v_cndmask_b32_e64 v18, 0, 0xffffffc0, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_mov_b32 v1, v13 :: v_dual_add_nc_u32 v18, v18, v17
	v_mad_u64_u32 v[13:14], null, 0xdb629599, v19, v[1:2]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_lt_u32_e64 s1, 31, v18
	v_cndmask_b32_e32 v11, v13, v11, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e64 v20, 0, 0xffffffe0, s1
	v_dual_mov_b32 v1, v14 :: v_dual_add_nc_u32 v20, v20, v18
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mad_u64_u32 v[14:15], null, 0xf534ddc0, v19, v[1:2]
	v_cmp_lt_u32_e64 s2, 31, v20
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_mov_b32 v1, v15 :: v_dual_cndmask_b32 v12, v14, v12
	v_mad_u64_u32 v[15:16], null, 0xfc2757d1, v19, v[1:2]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cndmask_b32_e64 v11, v12, v11, s1
	v_mov_b32_e32 v1, v16
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[16:17], null, 0x4e441529, v19, v[1:2]
	v_mov_b32_e32 v1, v17
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_mad_u64_u32 v[17:18], null, 0xa2f9836e, v19, v[1:2]
	v_cndmask_b32_e64 v1, 0, 0xffffffe0, s2
	v_cndmask_b32_e32 v19, v16, v14, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_nc_u32_e32 v1, v1, v20
	v_dual_cndmask_b32 v17, v17, v15 :: v_dual_cndmask_b32 v16, v18, v16
	v_cndmask_b32_e32 v15, v15, v13, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_sub_nc_u32_e32 v18, 32, v1
	v_cndmask_b32_e64 v14, v17, v19, s1
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cndmask_b32_e64 v16, v16, v17, s1
	v_cndmask_b32_e64 v17, v19, v15, s1
	v_cndmask_b32_e64 v15, v15, v12, s1
	v_cmp_eq_u32_e32 vcc_lo, 0, v1
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cndmask_b32_e64 v16, v16, v14, s2
	v_cndmask_b32_e64 v14, v14, v17, s2
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_cndmask_b32_e64 v17, v17, v15, s2
	v_cndmask_b32_e64 v11, v15, v11, s2
	v_alignbit_b32 v19, v16, v14, v18
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_alignbit_b32 v13, v14, v17, v18
	v_dual_cndmask_b32 v1, v19, v16 :: v_dual_cndmask_b32 v12, v13, v14
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_alignbit_b32 v16, v17, v11, v18
	v_bfe_u32 v13, v1, 29, 1
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_alignbit_b32 v14, v1, v12, 30
	v_cndmask_b32_e32 v16, v16, v17, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v15, 0, v13
	v_alignbit_b32 v12, v12, v16, 30
	v_alignbit_b32 v11, v16, v11, 30
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_xor_b32_e32 v14, v14, v15
	v_xor_b32_e32 v12, v12, v15
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_xor_b32_e32 v11, v11, v15
	v_lshrrev_b32_e32 v15, 29, v1
	v_clz_i32_u32_e32 v17, v14
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_lshlrev_b32_e32 v15, 31, v15
	v_min_u32_e32 v17, 32, v17
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_or_b32_e32 v18, 0.5, v15
	v_sub_nc_u32_e32 v16, 31, v17
	v_lshlrev_b32_e32 v19, 23, v17
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_alignbit_b32 v14, v14, v12, v16
	v_alignbit_b32 v11, v12, v11, v16
	v_sub_nc_u32_e32 v18, v18, v19
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_alignbit_b32 v12, v14, v11, 9
	v_lshrrev_b32_e32 v14, 9, v14
	v_clz_i32_u32_e32 v16, v12
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_min_u32_e32 v16, 32, v16
	v_sub_nc_u32_e32 v20, 31, v16
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_alignbit_b32 v11, v12, v11, v20
	v_or_b32_e32 v12, v14, v18
	v_add_lshl_u32 v14, v16, v17, 23
	v_lshrrev_b32_e32 v11, 9, v11
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_f32_e32 v16, 0x3fc90fda, v12
	v_sub_nc_u32_e32 v11, v11, v14
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v14, 0x3fc90fda, v12, -v16
	v_add_nc_u32_e32 v11, 0x33000000, v11
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v14, 0x33a22168, v12
	v_or_b32_e32 v11, v11, v15
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v14, 0x3fc90fda, v11
	v_lshrrev_b32_e32 v11, 30, v1
	v_add_f32_e32 v1, v16, v14
	s_delay_alu instid0(VALU_DEP_2)
	v_add_nc_u32_e32 v11, v13, v11
.LBB0_8:                                ; %Flow
                                        ;   in Loop: Header=BB0_6 Depth=3
	s_and_not1_saveexec_b32 s1, s16
; %bb.9:                                ;   in Loop: Header=BB0_6 Depth=3
	v_mul_f32_e64 v1, 0x3f22f983, |v10|
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_rndne_f32_e32 v11, v1
	v_fma_f32 v1, 0xbfc90fda, v11, |v10|
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v1, 0xb3a22168, v11
	v_fmac_f32_e32 v1, 0xa7c234c4, v11
	v_cvt_i32_f32_e32 v11, v11
; %bb.10:                               ; %_ZL3cosf.exit
                                        ;   in Loop: Header=BB0_6 Depth=3
	s_or_b32 exec_lo, exec_lo, s1
	v_div_scale_f32 v12, null, 0x42480000, 0x42480000, v9
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_dual_mul_f32 v14, v1, v1 :: v_dual_and_b32 v19, 1, v11
	v_div_scale_f32 v18, vcc_lo, v9, 0x42480000, v9
	v_rcp_f32_e32 v13, v12
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fmaak_f32 v17, s12, v14, 0xbab64f3b
	v_cmp_eq_u32_e64 s1, 0, v19
	v_lshlrev_b32_e32 v11, 30, v11
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_and_b32_e32 v11, 0x80000000, v11
	s_waitcnt_depctr 0xfff
	v_fma_f32 v15, -v12, v13, 1.0
	v_dual_fmaak_f32 v16, s11, v14, 0x3c0881c4 :: v_dual_fmac_f32 v13, v15, v13
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_fmaak_f32 v15, v14, v16, 0xbe2aaa9d
	v_fmaak_f32 v16, v14, v17, 0x3d2aabf7
	v_mul_f32_e32 v17, v18, v13
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_f32_e32 v15, v14, v15
	v_fmaak_f32 v16, v14, v16, 0xbf000004
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f32 v20, -v12, v17, v18
	v_fmac_f32_e32 v1, v1, v15
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f32 v14, v14, v16, 1.0
	v_fmac_f32_e32 v17, v20, v13
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cndmask_b32_e64 v1, -v1, v14, s1
	v_fma_f32 v12, -v12, v17, v18
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_xor_b32_e32 v1, v11, v1
	v_div_fmas_f32 v11, v12, v13, v17
	v_cmp_class_f32_e64 vcc_lo, v10, 0x1f8
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_f32_e32 v1, 0x42fe0000, v1
	v_div_fixup_f32 v9, v11, 0x42480000, v9
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v1, 0x7fc00000, v1, vcc_lo
	v_add_f32_e32 v9, 1.0, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_div_scale_f32 v10, null, v9, v9, v1
	v_div_scale_f32 v13, vcc_lo, v1, v9, v1
	v_rcp_f32_e32 v11, v10
	s_waitcnt_depctr 0xfff
	v_fma_f32 v12, -v10, v11, 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v11, v12, v11
	v_mul_f32_e32 v12, v13, v11
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v14, -v10, v12, v13
	v_fmac_f32_e32 v12, v14, v11
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v10, -v10, v12, v13
	v_div_fmas_f32 v10, v10, v11, v12
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(SALU_CYCLE_1)
	v_div_fixup_f32 v1, v10, v9, v1
	v_add_nc_u32_e32 v9, s15, v5
	s_add_i32 s15, s15, 1
	s_cmp_eq_u32 s4, s15
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f32_e32 v1, 0x43200000, v1
	v_ashrrev_i32_e32 v10, 31, v9
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v9, vcc_lo, s6, v9
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cvt_i32_f32_e32 v1, v1
	v_add_co_ci_u32_e32 v10, vcc_lo, s7, v10, vcc_lo
	global_store_b8 v[9:10], v1, off
	s_cbranch_scc0 .LBB0_6
.LBB0_11:                               ; %._crit_edge
                                        ;   in Loop: Header=BB0_4 Depth=2
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v7, 1, v7
	v_add_nc_u32_e32 v5, s4, v5
	v_cmp_eq_u32_e32 vcc_lo, v7, v3
	s_or_b32 s14, vcc_lo, s14
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s14
	s_cbranch_execnz .LBB0_4
.LBB0_12:                               ; %Flow65
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s13
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_add_nc_u32_e32 v4, s3, v4
	s_add_i32 s5, s5, 1
	s_cmp_eq_u32 s5, s8
	s_cbranch_scc0 .LBB0_2
.LBB0_13:                               ; %._crit_edge49
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z15fPixelGeneratoriiPh
		.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 17
		.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	_Z15fPixelGeneratoriiPh, .Lfunc_end0-_Z15fPixelGeneratoriiPh
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 2160
; NumSgprs: 19
; NumVgprs: 21
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 2
; NumSGPRsForWavesPerEU: 19
; NumVGPRsForWavesPerEU: 21
; 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:
      - .offset:         0
        .size:           4
        .value_kind:     by_value
      - .offset:         4
        .size:           4
        .value_kind:     by_value
      - .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:           _Z15fPixelGeneratoriiPh
    .private_segment_fixed_size: 0
    .sgpr_count:     19
    .sgpr_spill_count: 0
    .symbol:         _Z15fPixelGeneratoriiPh.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
