#include <stdio.h>
#include <cuda.h>

#define SIZE 1024
#define START_SCALE 1.5f
#define MAXIT 256
#define C_RE -0.8f
#define C_IM 0.156f
#define ZOOM 200

//======================================================================
/*This function checks whether a point
 belongs to the filled julia set. It 
 returns 0 if the value 'escaped', and
 one if the maximal numer of iterations
 has been reached. It is a function on the 
 device, but it can't be called like a
 kernel from the host. It is used by the
 kernel 'construct_julia_set' below. Note that
 the kernel has the 'global' attribute rather
 than the 'device' attribute.*/
__device__ int julia( int i, int j, float scale)
{

 //rescale grid to actual scale
 float x = scale * (float)(SIZE/2 - i)/(SIZE/2);
 float y = scale * (float)(SIZE/2 - j)/(SIZE/2);

 //real and imaginary part of point in question
 float z_re=x;
 float z_im=y;
 float z_re_old;

 //compute Z(n+1) = Zn^2 + C
 for (int k=0; k<MAXIT;k++)
 {
  z_re_old = z_re;                              //store old real value
  z_re =(z_re*z_re-z_im*z_im) + (C_RE);         //compute Re(Z(n+1))
  z_im = 2.0f*z_re_old*z_im + (C_IM);           //compute Im(Z(n+1))
  if ( sqrt(z_re*z_re+z_im*z_im) > SIZE)        //check if point escaped
  {
   return 0;                                    //point escaped
  }
 }
 return 1;                                      //point in set
}

//======================================================================
/* This function uses the one defined above
 to construct an array called 'set' with ones and zeroes, 
 defining the elements of the julia set. This is the
 kernel function to be called from the host and exectued
 on the device. It thus carries the 'global' attribute.*/
__global__ void construct_julia_set( int *set, float scale)
{
// int i = blockIdx.x*blockDim.x + threadIdx.x;	//determine thread id i
// int j = blockIdx.y*blockDim.y + threadIdx.y;	//determine thread id j
 int i = blockIdx.x;
 int j = blockIdx.y;
 int pos = i + j * SIZE;			//remap to 1-dim array
 set[pos] = julia( i, j, scale);              	//fill the set
}
//======================================================================
int main( void )
{

 using namespace std;

 int set[SIZE*SIZE];                    //the result array on the host
 int *set_d; 		                //the result array on the device
 int written;                           //aux variable for writing file
 int num;                               //numbering of output file
 float x;                               //Re part in complex Z plane
 float y;                               //Im part in complez Z plane
 float scale;                           //initial scale for zoom
 char buffer[32];                       //buffer for filenames
 FILE *out;                             //output file                           
 //dim3 threadsPerBlock(16,16);		//number of threads per block and grid size
 //dim3 numBlocks(SIZE/threadsPerBlock.x,SIZE/threadsPerBlock.y);
 dim3 grid(SIZE,SIZE);

 for(int k=0; k<SIZE*SIZE; k++)
 {
  set[k] = 0;
 }
 cudaMalloc(&set_d, SIZE*SIZE*sizeof(int)); //allocate memory for set on device

 /*The variable 'written' allows us to introduce a newline 
   after each row that has been written to the outpur file. 
   This allows for more freedom in printing the set using
   gnuplot.*/

 written=0;                                             //reset 'written' value
 for( int k=0; k<ZOOM; k++ )                             //k...number of zoom slices to produce
 {                                                      //vary scale for zoom
  scale = START_SCALE *(400.0f-(float)k)/400.0f + 0.01;
  cudaMemcpy(set_d, set, SIZE*SIZE*sizeof(int), cudaMemcpyHostToDevice); //init set on device
//  construct_julia_set<<<numBlocks,threadsPerBlock>>>(set_d, scale); //construct julia set on GPU
  construct_julia_set<<<grid,1>>>(set_d, scale); //construct julia set on GPU
  cudaDeviceSynchronize();
  cudaMemcpy(set, set_d, SIZE*SIZE*sizeof(int), cudaMemcpyDeviceToHost); //copy resultto host
  num = k;                                              //out: 'julia_000.dat', 'julia_001.dat',...
  snprintf(buffer, sizeof(char)*32, "julia_%03i.dat", num);
  out = fopen( buffer, "wt" );                          //write in text mode (wt)
  for (int i=0; i<SIZE; i++)                            //actual grid values x and y
  {
   x = scale * (float)(SIZE/2 - i)/(SIZE/2);
   for (int j=0; j<SIZE; j++)
   {
    y = scale * (float)(SIZE/2 - j)/(SIZE/2);
    int pos = i + j * SIZE;                             //position in  array
    if(set[pos]==1)                                     //write only if part of set
    {
     fprintf(out,"%f %f \n",x,y);
     written = 1;                                       //set written to 1
    }
   }//end inner grid loop (j)
   if( written == 1 )
   {
    fprintf(out,"\n");                                  //add newline if row content not empty
   }
   written=0;                                           //reset written value
  }//end outer grid loop (i)
 fclose(out);                                           //close file
 }//end zoom loop (k)
 cudaFree( set_d );					//deallocate
}

	.file	"tmpxft_0034ea77_00000000-6_julia_set.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2031:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2031:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z5juliaiif
	.type	_Z5juliaiif, @function
_Z5juliaiif:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2027:
	.size	_Z5juliaiif, .-_Z5juliaiif
	.globl	_Z40__device_stub__Z19construct_julia_setPifPif
	.type	_Z40__device_stub__Z19construct_julia_setPifPif, @function
_Z40__device_stub__Z19construct_julia_setPifPif:
.LFB2053:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 8(%rsp)
	leaq	32(%rsp), %rcx
	leaq	24(%rsp), %rdx
	movss	%xmm0, 4(%rsp)
	leaq	52(%rsp), %rsi
	leaq	40(%rsp), %rdi
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 48(%rsp)
	movq	%rax, 88(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 40(%rsp)
	movq	%rax, 52(%rsp)
	movl	$1, 60(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L4
	pushq	32(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z19construct_julia_setPif(%rip), %rdi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 144
	movq	68(%rsp), %rcx
	movl	76(%rsp), %r8d
	movq	56(%rsp), %rsi
	movl	64(%rsp), %edx
	leaq	104(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L4:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L6
	call	__stack_chk_fail@PLT
.L6:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2053:
	.size	_Z40__device_stub__Z19construct_julia_setPifPif, .-_Z40__device_stub__Z19construct_julia_setPifPif
	.globl	_Z19construct_julia_setPif
	.type	_Z19construct_julia_setPif, @function
_Z19construct_julia_setPif:
.LFB2054:
	.cfi_startproc
	endbr64
	jmp	_Z40__device_stub__Z19construct_julia_setPifPif
	.cfi_endproc
.LFE2054:
	.size	_Z19construct_julia_setPif, .-_Z19construct_julia_setPif
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC3:
	.string	"julia_%03i.dat"
.LC4:
	.string	"wt"
.LC6:
	.string	"%f %f \n"
.LC7:
	.string	"\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2028:
	.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
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	leaq	-4194304(%rsp), %r11
	.cfi_def_cfa 11, 4194360
.LPSRL0:
	subq	$4096, %rsp
	orq	$0, (%rsp)
	cmpq	%r11, %rsp
	jne	.LPSRL0
	.cfi_def_cfa_register 7
	subq	$104, %rsp
	.cfi_def_cfa_offset 4194464
	movl	$1048576, %ecx
	movl	$4194304, %esi
	xorl	%ebx, %ebx
	movq	%fs:40, %rax
	movq	%rax, 4194392(%rsp)
	movabsq	$4398046512128, %rax
	leaq	56(%rsp), %rbp
	movl	$1, 40(%rsp)
	movq	%rbp, %rdi
	movq	%rax, 32(%rsp)
	xorl	%eax, %eax
	rep stosl
	leaq	24(%rsp), %rdi
	call	cudaMalloc@PLT
.L16:
	cvtsi2ssl	%ebx, %xmm1
	movss	.LC0(%rip), %xmm0
	movq	24(%rsp), %rdi
	movq	%rbp, %rsi
	movl	$1, %ecx
	movl	$4194304, %edx
	subss	%xmm1, %xmm0
	mulss	.LC1(%rip), %xmm0
	divss	.LC0(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	addsd	.LC2(%rip), %xmm0
	cvtsd2ss	%xmm0, %xmm3
	movss	%xmm3, 8(%rsp)
	call	cudaMemcpy@PLT
	movq	32(%rsp), %rdi
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	40(%rsp), %esi
	movl	$1, %ecx
	movabsq	$4294967297, %rdx
	movl	$1, 52(%rsp)
	movq	%rdx, 44(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L11
	movss	8(%rsp), %xmm0
	movq	24(%rsp), %rdi
	call	_Z40__device_stub__Z19construct_julia_setPifPif
.L11:
	call	cudaDeviceSynchronize@PLT
	movq	24(%rsp), %rsi
	movq	%rbp, %rdi
	xorl	%r13d, %r13d
	movl	$2, %ecx
	movl	$4194304, %edx
	leaq	4194360(%rsp), %r12
	call	cudaMemcpy@PLT
	movl	%ebx, %r9d
	movl	$32, %ecx
	movq	%r12, %rdi
	leaq	.LC3(%rip), %r8
	movl	$2, %edx
	movl	$32, %esi
	xorl	%eax, %eax
	call	__snprintf_chk@PLT
	leaq	.LC4(%rip), %rsi
	movq	%r12, %rdi
	call	fopen@PLT
	movq	%rax, (%rsp)
.L15:
	movl	$512, %eax
	leaq	0(%rbp,%r13,4), %r12
	movl	$512, %r15d
	subl	%r13d, %eax
	leaq	.LC6(%rip), %r14
	cvtsi2ssl	%eax, %xmm2
	mulss	8(%rsp), %xmm2
	xorl	%eax, %eax
	mulss	.LC5(%rip), %xmm2
.L13:
	cmpl	$1, (%r12)
	jne	.L12
	cvtsi2ssl	%r15d, %xmm1
	mulss	8(%rsp), %xmm1
	movq	%r14, %rdx
	movb	$2, %al
	mulss	.LC5(%rip), %xmm1
	movq	(%rsp), %rdi
	movl	$2, %esi
	movss	%xmm2, 12(%rsp)
	cvtss2sd	%xmm2, %xmm0
	cvtss2sd	%xmm1, %xmm1
	call	__fprintf_chk@PLT
	movss	12(%rsp), %xmm2
	movl	$1, %eax
.L12:
	decl	%r15d
	addq	$4096, %r12
	cmpl	$-512, %r15d
	jne	.L13
	decl	%eax
	jne	.L14
	movq	(%rsp), %rdi
	leaq	.LC7(%rip), %rdx
	movl	$2, %esi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
.L14:
	incq	%r13
	cmpq	$1024, %r13
	jne	.L15
	movq	(%rsp), %rdi
	incl	%ebx
	call	fclose@PLT
	cmpl	$200, %ebx
	jne	.L16
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
	movq	4194392(%rsp), %rax
	subq	%fs:40, %rax
	je	.L17
	call	__stack_chk_fail@PLT
.L17:
	addq	$4194408, %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
.LFE2028:
	.size	main, .-main
	.section	.rodata.str1.1
.LC8:
	.string	"_Z19construct_julia_setPif"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2056:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC8(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z19construct_julia_setPif(%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
.LFE2056:
	.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.cst4,"aM",@progbits,4
	.align 4
.LC0:
	.long	1137180672
	.align 4
.LC1:
	.long	1069547520
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC2:
	.long	1202590843
	.long	1065646817
	.section	.rodata.cst4
	.align 4
.LC5:
	.long	989855744
	.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 <stdio.h>
#include <hip/hip_runtime.h>

#define SIZE 1024
#define START_SCALE 1.5f
#define MAXIT 256
#define C_RE -0.8f
#define C_IM 0.156f
#define ZOOM 200

//======================================================================
/*This function checks whether a point
 belongs to the filled julia set. It 
 returns 0 if the value 'escaped', and
 one if the maximal numer of iterations
 has been reached. It is a function on the 
 device, but it can't be called like a
 kernel from the host. It is used by the
 kernel 'construct_julia_set' below. Note that
 the kernel has the 'global' attribute rather
 than the 'device' attribute.*/
__device__ int julia( int i, int j, float scale)
{

 //rescale grid to actual scale
 float x = scale * (float)(SIZE/2 - i)/(SIZE/2);
 float y = scale * (float)(SIZE/2 - j)/(SIZE/2);

 //real and imaginary part of point in question
 float z_re=x;
 float z_im=y;
 float z_re_old;

 //compute Z(n+1) = Zn^2 + C
 for (int k=0; k<MAXIT;k++)
 {
  z_re_old = z_re;                              //store old real value
  z_re =(z_re*z_re-z_im*z_im) + (C_RE);         //compute Re(Z(n+1))
  z_im = 2.0f*z_re_old*z_im + (C_IM);           //compute Im(Z(n+1))
  if ( sqrt(z_re*z_re+z_im*z_im) > SIZE)        //check if point escaped
  {
   return 0;                                    //point escaped
  }
 }
 return 1;                                      //point in set
}

//======================================================================
/* This function uses the one defined above
 to construct an array called 'set' with ones and zeroes, 
 defining the elements of the julia set. This is the
 kernel function to be called from the host and exectued
 on the device. It thus carries the 'global' attribute.*/
__global__ void construct_julia_set( int *set, float scale)
{
// int i = blockIdx.x*blockDim.x + threadIdx.x;	//determine thread id i
// int j = blockIdx.y*blockDim.y + threadIdx.y;	//determine thread id j
 int i = blockIdx.x;
 int j = blockIdx.y;
 int pos = i + j * SIZE;			//remap to 1-dim array
 set[pos] = julia( i, j, scale);              	//fill the set
}
//======================================================================
int main( void )
{

 using namespace std;

 int set[SIZE*SIZE];                    //the result array on the host
 int *set_d; 		                //the result array on the device
 int written;                           //aux variable for writing file
 int num;                               //numbering of output file
 float x;                               //Re part in complex Z plane
 float y;                               //Im part in complez Z plane
 float scale;                           //initial scale for zoom
 char buffer[32];                       //buffer for filenames
 FILE *out;                             //output file                           
 //dim3 threadsPerBlock(16,16);		//number of threads per block and grid size
 //dim3 numBlocks(SIZE/threadsPerBlock.x,SIZE/threadsPerBlock.y);
 dim3 grid(SIZE,SIZE);

 for(int k=0; k<SIZE*SIZE; k++)
 {
  set[k] = 0;
 }
 hipMalloc(&set_d, SIZE*SIZE*sizeof(int)); //allocate memory for set on device

 /*The variable 'written' allows us to introduce a newline 
   after each row that has been written to the outpur file. 
   This allows for more freedom in printing the set using
   gnuplot.*/

 written=0;                                             //reset 'written' value
 for( int k=0; k<ZOOM; k++ )                             //k...number of zoom slices to produce
 {                                                      //vary scale for zoom
  scale = START_SCALE *(400.0f-(float)k)/400.0f + 0.01;
  hipMemcpy(set_d, set, SIZE*SIZE*sizeof(int), hipMemcpyHostToDevice); //init set on device
//  construct_julia_set<<<numBlocks,threadsPerBlock>>>(set_d, scale); //construct julia set on GPU
  construct_julia_set<<<grid,1>>>(set_d, scale); //construct julia set on GPU
  hipDeviceSynchronize();
  hipMemcpy(set, set_d, SIZE*SIZE*sizeof(int), hipMemcpyDeviceToHost); //copy resultto host
  num = k;                                              //out: 'julia_000.dat', 'julia_001.dat',...
  snprintf(buffer, sizeof(char)*32, "julia_%03i.dat", num);
  out = fopen( buffer, "wt" );                          //write in text mode (wt)
  for (int i=0; i<SIZE; i++)                            //actual grid values x and y
  {
   x = scale * (float)(SIZE/2 - i)/(SIZE/2);
   for (int j=0; j<SIZE; j++)
   {
    y = scale * (float)(SIZE/2 - j)/(SIZE/2);
    int pos = i + j * SIZE;                             //position in  array
    if(set[pos]==1)                                     //write only if part of set
    {
     fprintf(out,"%f %f \n",x,y);
     written = 1;                                       //set written to 1
    }
   }//end inner grid loop (j)
   if( written == 1 )
   {
    fprintf(out,"\n");                                  //add newline if row content not empty
   }
   written=0;                                           //reset written value
  }//end outer grid loop (i)
 fclose(out);                                           //close file
 }//end zoom loop (k)
 hipFree( set_d );					//deallocate
}

	.text
	.file	"julia_set.hip"
	.globl	_Z34__device_stub__construct_julia_setPif # -- Begin function _Z34__device_stub__construct_julia_setPif
	.type	_Z34__device_stub__construct_julia_setPif,@function
_Z34__device_stub__construct_julia_setPif: # @_Z34__device_stub__construct_julia_setPif
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$80, %rsp
	.cfi_def_cfa_offset 128
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	24(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	4(%rsp), %rcx
	movss	%xmm0, (%rcx)
	leaq	64(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%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	$_Z19construct_julia_setPif, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$96, %rsp
	.cfi_adjust_cfa_offset -96
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z34__device_stub__construct_julia_setPif, .Lfunc_end0-_Z34__device_stub__construct_julia_setPif
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function main
.LCPI1_0:
	.long	0x43c80000                      # float 400
.LCPI1_1:
	.long	0x3fc00000                      # float 1.5
.LCPI1_3:
	.long	0x3b000000                      # float 0.001953125
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0
.LCPI1_2:
	.quad	0x3f847ae147ae147b              # double 0.01
	.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	$4194376, %rsp                  # imm = 0x400048
	.cfi_def_cfa_offset 4194432
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	leaq	64(%rsp), %r14
	movl	$4194304, %edx                  # imm = 0x400000
	movq	%r14, %rdi
	xorl	%esi, %esi
	callq	memset@PLT
	leaq	16(%rsp), %rdi
	movl	$4194304, %esi                  # imm = 0x400000
	callq	hipMalloc
	xorl	%ebp, %ebp
.LBB1_1:                                # =>This Loop Header: Depth=1
                                        #     Child Loop BB1_4 Depth 2
                                        #       Child Loop BB1_5 Depth 3
                                        #         Child Loop BB1_6 Depth 4
	xorps	%xmm0, %xmm0
	cvtsi2ss	%ebp, %xmm0
	movss	.LCPI1_0(%rip), %xmm2           # xmm2 = mem[0],zero,zero,zero
	movaps	%xmm2, %xmm1
	subss	%xmm0, %xmm1
	mulss	.LCPI1_1(%rip), %xmm1
	divss	%xmm2, %xmm1
	xorps	%xmm0, %xmm0
	cvtss2sd	%xmm1, %xmm0
	addsd	.LCPI1_2(%rip), %xmm0
	cvtsd2ss	%xmm0, %xmm0
	movss	%xmm0, 12(%rsp)                 # 4-byte Spill
	movq	16(%rsp), %rdi
	movl	$4194304, %edx                  # imm = 0x400000
	movq	%r14, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movabsq	$4398046512128, %rdi            # imm = 0x40000000400
	movl	$1, %esi
	movabsq	$4294967297, %rdx               # imm = 0x100000001
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB1_3
# %bb.2:                                #   in Loop: Header=BB1_1 Depth=1
	movq	16(%rsp), %rdi
	movss	12(%rsp), %xmm0                 # 4-byte Reload
                                        # xmm0 = mem[0],zero,zero,zero
	callq	_Z34__device_stub__construct_julia_setPif
.LBB1_3:                                #   in Loop: Header=BB1_1 Depth=1
	callq	hipDeviceSynchronize
	movq	16(%rsp), %rsi
	movl	$4194304, %edx                  # imm = 0x400000
	movq	%r14, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movl	$32, %esi
	movl	$.L.str, %edx
	leaq	32(%rsp), %rbx
	movq	%rbx, %rdi
	movl	%ebp, %ecx
	xorl	%eax, %eax
	callq	snprintf
	movl	$.L.str.1, %esi
	movq	%rbx, %rdi
	callq	fopen
	movq	%rax, %r13
	movq	%r14, %rbx
	xorl	%r15d, %r15d
.LBB1_4:                                #   Parent Loop BB1_1 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB1_5 Depth 3
                                        #         Child Loop BB1_6 Depth 4
	movl	$512, %eax                      # imm = 0x200
	subl	%r15d, %eax
	xorps	%xmm0, %xmm0
	cvtsi2ss	%eax, %xmm0
	mulss	12(%rsp), %xmm0                 # 4-byte Folded Reload
	mulss	.LCPI1_3(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	xorl	%r12d, %r12d
	xorl	%eax, %eax
.LBB1_5:                                # %.outer
                                        #   Parent Loop BB1_1 Depth=1
                                        #     Parent Loop BB1_4 Depth=2
                                        # =>    This Loop Header: Depth=3
                                        #         Child Loop BB1_6 Depth 4
	leaq	1(%r12), %rsi
	leal	-512(%r12), %ecx
	shlq	$12, %r12
	addq	%rbx, %r12
	movq	%r12, %rdx
	movq	%rsi, %r12
.LBB1_6:                                #   Parent Loop BB1_1 Depth=1
                                        #     Parent Loop BB1_4 Depth=2
                                        #       Parent Loop BB1_5 Depth=3
                                        # =>      This Inner Loop Header: Depth=4
	cmpl	$1, (%rdx)
	je	.LBB1_11
# %bb.7:                                #   in Loop: Header=BB1_6 Depth=4
	incq	%r12
	incl	%ecx
	addq	$4096, %rdx                     # imm = 0x1000
	cmpq	$1025, %r12                     # imm = 0x401
	jne	.LBB1_6
	jmp	.LBB1_8
.LBB1_11:                               # %.thread
                                        #   in Loop: Header=BB1_5 Depth=3
	negl	%ecx
	xorps	%xmm1, %xmm1
	cvtsi2ss	%ecx, %xmm1
	mulss	12(%rsp), %xmm1                 # 4-byte Folded Reload
	mulss	.LCPI1_3(%rip), %xmm1
	cvtss2sd	%xmm1, %xmm1
	movl	$.L.str.2, %esi
	movq	%r13, %rdi
	movb	$2, %al
	movsd	%xmm0, 24(%rsp)                 # 8-byte Spill
	callq	fprintf
	movsd	24(%rsp), %xmm0                 # 8-byte Reload
                                        # xmm0 = mem[0],zero
	movb	$1, %al
	cmpq	$1024, %r12                     # imm = 0x400
	jne	.LBB1_5
	jmp	.LBB1_12
.LBB1_8:                                #   in Loop: Header=BB1_4 Depth=2
	testb	$1, %al
	je	.LBB1_13
.LBB1_12:                               # %.thread51
                                        #   in Loop: Header=BB1_4 Depth=2
	movl	$10, %edi
	movq	%r13, %rsi
	callq	fputc@PLT
.LBB1_13:                               #   in Loop: Header=BB1_4 Depth=2
	incq	%r15
	addq	$4, %rbx
	cmpq	$1024, %r15                     # imm = 0x400
	jne	.LBB1_4
# %bb.9:                                #   in Loop: Header=BB1_1 Depth=1
	movq	%r13, %rdi
	callq	fclose
	incl	%ebp
	cmpl	$200, %ebp
	jne	.LBB1_1
# %bb.10:
	movq	16(%rsp), %rdi
	callq	hipFree
	xorl	%eax, %eax
	addq	$4194376, %rsp                  # imm = 0x400048
	.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
.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	$_Z19construct_julia_setPif, %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	_Z19construct_julia_setPif,@object # @_Z19construct_julia_setPif
	.section	.rodata,"a",@progbits
	.globl	_Z19construct_julia_setPif
	.p2align	3, 0x0
_Z19construct_julia_setPif:
	.quad	_Z34__device_stub__construct_julia_setPif
	.size	_Z19construct_julia_setPif, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"julia_%03i.dat"
	.size	.L.str, 15

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

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

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z19construct_julia_setPif"
	.size	.L__unnamed_1, 27

	.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 _Z34__device_stub__construct_julia_setPif
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z19construct_julia_setPif
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z19construct_julia_setPif ; -- Begin function _Z19construct_julia_setPif
	.globl	_Z19construct_julia_setPif
	.p2align	8
	.type	_Z19construct_julia_setPif,@function
_Z19construct_julia_setPif:             ; @_Z19construct_julia_setPif
; %bb.0:
	s_load_b32 s2, s[0:1], 0x8
	s_sub_i32 s3, 0x200, s15
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_cvt_f32_i32_e32 v0, s3
	s_sub_i32 s3, 0x200, s14
	v_cvt_f32_i32_e32 v1, s3
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_mul_f32 v0, s2, v0 :: v_dual_mul_f32 v1, s2, v1
	v_dual_mul_f32 v0, 0x3b000000, v0 :: v_dual_mul_f32 v1, 0x3b000000, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_mul_f32 v2, v0, v0 :: v_dual_add_f32 v3, v1, v1
	v_fma_f32 v2, v1, v1, -v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fmaak_f32 v0, v0, v3, 0x3e1fbe77
	v_add_f32_e32 v1, 0xbf4ccccd, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v2, v1, v1
	v_fmac_f32_e32 v2, v0, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_f32_e32 v3, 0x4f800000, v2
	v_cmp_gt_f32_e32 vcc_lo, 0xf800000, v2
	v_cndmask_b32_e32 v2, v2, v3, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_2)
	v_sqrt_f32_e32 v3, v2
	s_waitcnt_depctr 0xfff
	v_add_nc_u32_e32 v4, -1, v3
	v_add_nc_u32_e32 v5, 1, v3
	v_fma_f32 v6, -v4, v3, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v7, -v5, v3, v2
	v_cmp_ge_f32_e64 s2, 0, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cndmask_b32_e64 v3, v3, v4, s2
	v_cmp_lt_f32_e64 s2, 0, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cndmask_b32_e64 v3, v3, v5, s2
	s_mov_b32 s2, 0
	v_mul_f32_e32 v4, 0x37800000, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v3, v3, v4, vcc_lo
	v_cmp_class_f32_e64 vcc_lo, v2, 0x260
	v_cndmask_b32_e32 v2, v3, v2, vcc_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmp_lt_f32_e32 vcc_lo, 0x44800000, v2
	v_mov_b32_e32 v2, 0
	s_cbranch_vccnz .LBB0_6
.LBB0_1:                                ; %.lr.ph
                                        ; =>This Inner Loop Header: Depth=1
	s_mov_b32 s3, s2
	s_cmpk_eq_i32 s2, 0xff
	s_cbranch_scc1 .LBB0_3
; %bb.2:                                ;   in Loop: Header=BB0_1 Depth=1
	v_mul_f32_e32 v2, v0, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_fma_f32 v2, v1, v1, -v2
	v_add_f32_e32 v1, v1, v1
	v_fmaak_f32 v0, v0, v1, 0x3e1fbe77
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v2, 0xbf4ccccd, v2
	v_mul_f32_e32 v1, v2, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v1, v0, v0
	v_mul_f32_e32 v3, 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, v3, vcc_lo
	v_sqrt_f32_e32 v3, v1
	s_waitcnt_depctr 0xfff
	v_add_nc_u32_e32 v4, -1, v3
	v_add_nc_u32_e32 v5, 1, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v6, -v4, v3, v1
	v_fma_f32 v7, -v5, v3, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_ge_f32_e64 s2, 0, v6
	v_cndmask_b32_e64 v3, v3, v4, s2
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_lt_f32_e64 s2, 0, v7
	v_cndmask_b32_e64 v3, v3, v5, s2
	s_add_i32 s2, s3, 1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v4, 0x37800000, v3
	v_cndmask_b32_e32 v3, v3, v4, 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 v1, v3, v1, vcc_lo
	v_cmp_lt_f32_e64 s4, 0x44800000, v1
	v_mov_b32_e32 v1, v2
	s_branch .LBB0_4
.LBB0_3:                                ;   in Loop: Header=BB0_1 Depth=1
	s_mov_b32 s4, -1
                                        ; implicit-def: $sgpr2
                                        ; implicit-def: $vgpr1
                                        ; implicit-def: $vgpr0
.LBB0_4:                                ; %Flow
                                        ;   in Loop: Header=BB0_1 Depth=1
	s_delay_alu instid0(VALU_DEP_2) | instid1(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s4
	s_cbranch_vccnz .LBB0_1
; %bb.5:                                ; %_Z5juliaiif.exit.loopexit
	s_cmpk_gt_u32 s3, 0xfe
	s_cselect_b32 s2, -1, 0
	s_delay_alu instid0(SALU_CYCLE_1)
	v_cndmask_b32_e64 v2, 0, 1, s2
.LBB0_6:                                ; %_Z5juliaiif.exit
	s_load_b64 s[0:1], s[0:1], 0x0
	s_lshl_b32 s2, s15, 10
	v_mov_b32_e32 v0, 0
	s_add_i32 s2, s2, s14
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_ashr_i32 s3, s2, 31
	s_lshl_b64 s[2:3], s[2:3], 2
	s_waitcnt lgkmcnt(0)
	s_add_u32 s0, s0, s2
	s_addc_u32 s1, s1, s3
	global_store_b32 v0, v2, s[0:1]
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z19construct_julia_setPif
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 12
		.amdhsa_user_sgpr_count 14
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 1
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 8
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end0:
	.size	_Z19construct_julia_setPif, .Lfunc_end0-_Z19construct_julia_setPif
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 600
; NumSgprs: 18
; NumVgprs: 8
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 8
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 14
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 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
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 12
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z19construct_julia_setPif
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z19construct_julia_setPif.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     8
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
