

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

/* To index element (i,j) of a 2D array stored as 1D */
#define index(i, j, N)  ((i)*(N)) + (j)

/*****************************************************************/

// Function declarations: Feel free to add any functions you want.
void  seq_heat_dist(float *, unsigned int, unsigned int);
void  gpu_heat_dist(float *, unsigned int, unsigned int);


/*****************************************************************/
/**** Do NOT CHANGE ANYTHING in main() function ******/

int main(int argc, char * argv[])
{
  unsigned int N; /* Dimention of NxN matrix */
  int type_of_device = 0; // CPU or GPU
  int iterations = 0;
  int i;
  
  /* The 2D array of points will be treated as 1D array of NxN elements */
  float * playground; 
  
  // to measure time taken by a specific part of the code 
  double time_taken;
  clock_t start, end;
  
  if(argc != 4)
  {
    fprintf(stderr, "usage: heatdist num  iterations  who\n");
    fprintf(stderr, "num = dimension of the square matrix (50 and up)\n");
    fprintf(stderr, "iterations = number of iterations till stopping (1 and up)\n");
    fprintf(stderr, "who = 0: sequential code on CPU, 1: GPU execution\n");
    exit(1);
  }
  
  type_of_device = atoi(argv[3]);
  N = (unsigned int) atoi(argv[1]);
  iterations = (unsigned int) atoi(argv[2]);
 
  
  /* Dynamically allocate NxN array of floats */
  playground = (float *)calloc(N*N, sizeof(float));
  if( !playground )
  {
   fprintf(stderr, " Cannot allocate the %u x %u array\n", N, N);
   exit(1);
  }
  
  /* Initialize it: calloc already initalized everything to 0 */
  // Edge elements to 70F
  for(i = 0; i < N; i++)
    playground[index(0,i,N)] = 70;
    
  for(i = 0; i < N; i++)
    playground[index(i,0,N)] = 70;
  
  for(i = 0; i < N; i++)
    playground[index(i,N-1, N)] = 70;
  
  for(i = 0; i < N; i++)
    playground[index(N-1,i,N)] = 70;
  
  // from (0,10) to (0,30) inclusive are 100F
  for(i = 10; i <= 30; i++)
    playground[index(0,i,N)] = 100;
  
   // from (n-1,10) to (n-1,30) inclusive are 150F
  for(i = 10; i <= 30; i++)
    playground[index(N-1,i,N)] = 150;
  
  if( !type_of_device ) // The CPU sequential version
  { 
    start = clock();
    seq_heat_dist(playground, N, iterations);
    end = clock();
  }
  else  // The GPU version
  {   
     start = clock();
     gpu_heat_dist(playground, N, iterations); 
     end = clock();    
  }
  
  
  time_taken = ((double)(end - start))/ CLOCKS_PER_SEC;
  
  printf("Time taken for %s is %lf\n", type_of_device == 0? "CPU" : "GPU", time_taken);
  
  free(playground);
  
  return 0;

}


/*****************  The CPU sequential version (DO NOT CHANGE THIS) **************/
void  seq_heat_dist(float * playground, unsigned int N, unsigned int iterations)
{
  // Loop indices
  int i, j, k;
  int upper = N-1;
  
  // number of bytes to be copied between array temp and array playground
  unsigned int num_bytes = 0;
  
  float * temp; 
  /* Dynamically allocate another array for temp values */
  /* Dynamically allocate NxN array of floats */
  temp = (float *)calloc(N*N, sizeof(float));
  if( !temp )
  {
   fprintf(stderr, " Cannot allocate temp %u x %u array\n", N, N);
   exit(1);
  }
  
  num_bytes = N*N*sizeof(float);
  
  /* Copy initial array in temp */
  memcpy((void *)temp, (void *) playground, num_bytes);
  
  for( k = 0; k < iterations; k++)
  {
    /* Calculate new values and store them in temp */
    for(i = 1; i < upper; i++)
      for(j = 1; j < upper; j++)
	temp[index(i,j,N)] = (playground[index(i-1,j,N)] + 
	                      playground[index(i+1,j,N)] + 
			      playground[index(i,j-1,N)] + 
			      playground[index(i,j+1,N)])/4.0;
  
			      
   			      
    /* Move new values into old values */ 
    memcpy((void *)playground, (void *) temp, num_bytes);
  }
  
}

/***************** The GPU version: Write your code here *********************/
/* This function can call one or more kernels if you want ********************/

// There will be two main functions that can be parallelized: one to average individual points around each point and one to update the current matrix's points for the next iteration to work with. 

__global__ void spread_to_point(float * current, unsigned int N, unsigned int iter, float * fresh) 
{ // Averages the four surrounding points to update a single point.

  // Let's make a grid-stride with a 2D grid, to fit the problem.

  unsigned int ind_i = blockDim.x * blockIdx.x + threadIdx.x; // Current block and current thread for the i-coord.
  unsigned int ind_j = blockDim.y * blockIdx.y + threadIdx.y; // Current block and current thread for the j-coord.

  int stride_i = blockDim.x * gridDim.x; // Next in line, if needed.
  int stride_j = blockDim.y * gridDim.y;  // Next in line, if needed. 

  for ( int j = ind_j+1 ; j < N-1 ; j += stride_j )
    for ( int i = ind_i+1 ; i < N-1 ; i += stride_i )
    {

    // Uncomment here if you want to take advantage of shared memory. 
    // __shared__ 
    // current[(i-1) * Nm2 + j], 
    // current[(i+1) * Nm2 + j], 
    // current[i * Nm2 + (j-1)], 
    // current[i * Nm2 + (j+1)];

    fresh[i * N + j] = ( // Multiply N by i (the row #) since the input data still represents the matrix as a 1D structure. 
      current[(i-1) * N + j] + 
      current[(i+1) * N + j] + 
      current[i * N + (j-1)] + 
      current[i * N + (j+1)]
      ) / 4;
    }
}

__global__ void overwrite_current_iteration(float * current, unsigned int N, unsigned int iter, float * fresh) 
{ // After computing all values using the old iteration, this function will make the new values take the old values' places. 

  // Again let's make a grid-stride with a 2D grid to fit the problem.

  unsigned int ind_i = blockDim.x * blockIdx.x + threadIdx.x; // Current block and current thread for the i-coord.
  unsigned int ind_j = blockDim.y * blockIdx.y + threadIdx.y; // Current block and current thread for the j-coord.

  int stride_i = blockDim.x * gridDim.x;
  int stride_j = blockDim.y * gridDim.y;   

  int index = ind_i * N + ind_j; 

  for ( int j = ind_j+1 ; j < N-1 ; j += stride_j )
    for ( int i = ind_i+1 ; i < N-1 ; i += stride_i )
      current[index] = fresh[index];
}

// The two commented functions below were anticipating writing parallel code to initialize the mesh temperatures (parallelized instead of sequential), however that is already done sequentially in main(), so the two below are not needed, but may be conceptually useful.  

// __global__ void initialize_edges_or_not(int N, float * grid, int iterations) 
// { // Checks if a point is on the edge or not, setting to 70 and 0 respectively.  This is used only once, to efficiently initialize values.

//   int i = blockDim.x * blockIdx.x +threadIdx.x; // Current block and current thread for the i-coord.
//   int j = blockDim.y * blockIdx.y +threadIdx.y; // Current block and current thread for the j-coord.

//   int index = i * N + j;

//   if ( index == 0 || index == N - 1 ) 
//   {
//     h[index] = 70;
//   }
//   else if ( index%N == 0 ) 
//   {
//     h[index] = 70;
//   } 
//   else {
//     h[index] = 0;
//   }
// }

// __global__ void initialize_special_sections(int N, float * grid, int iterations) 
// { // Can specificy swaths of points to have special values. This is used only once, to efficiently initialize values. Separated into another function to prevent too many redundant condition checks with high N in the init..edge() func. 

//   int i = blockDim.x * blockIdx.x +threadIdx.x; // Current block and current thread for the i-coord.
//   int j = blockDim.y * blockIdx.y +threadIdx.y; // Current block and current thread for the j-coord.

//   int index = i * N + j;
//   if (i==0 || i==N) 
//   {
//     // So on...
//   }
// }

int calcBlocks(unsigned int Nm2, int tpb) 
{
  int remain = Nm2 % tpb;
  int sheared = Nm2 - remain;
  int result = (sheared / tpb) + 1;

  return result; 
}

void  gpu_heat_dist(float * playground, unsigned int N, unsigned int iterations)
{
  float *current, *fresh; 
  cudaMallocManaged(&current, N*N*sizeof(float));
  cudaMallocManaged(&fresh, N*N*sizeof(float));
  
  cudaMemcpy(current, playground, N*N*sizeof(float), cudaMemcpyHostToDevice); 
  cudaMemset(fresh, 0, N*N*sizeof(int));
  // Tiling setup - how many threads per block, and how many blocks in the grid. 

  unsigned int blkcount = (N + 16 - 1) / 16;
  dim3 threadsPerBlock(16, 16);

  // If you want to control the block calculation, then uncomment here and comment the other numBlocks(...) declaration. 
  // dim3 numBlocks
  // (
  //   calcBlocks(Nm2, threadsPerBlock.x), 
  //   calcBlocks(Nm2, threadsPerBlock.y)
  // );
  dim3 numBlocks(blkcount,blkcount);

  for ( int i = 0 ; i < iterations ; i++ ) 
  {
    spread_to_point<<<numBlocks,threadsPerBlock>>>(current, N, iterations, fresh);
    cudaDeviceSynchronize();
    overwrite_current_iteration<<<numBlocks,threadsPerBlock>>>(current, N, iterations, fresh);
    cudaDeviceSynchronize();

    int lessAmt = 10; 
    for ( int j = N ; j < N+lessAmt ; j++ )
      printf("%.5f", current[j]);
    printf("\n");
  }
	
  cudaMemcpy(playground, current, N*sizeof(float), cudaMemcpyDeviceToHost);
  cudaFree(&current);
  cudaFree(&fresh);


}




	.file	"tmpxft_00372be0_00000000-6_heatdist.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2033:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2033:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	" Cannot allocate temp %u x %u array\n"
	.text
	.globl	_Z13seq_heat_distPfjj
	.type	_Z13seq_heat_distPfjj, @function
_Z13seq_heat_distPfjj:
.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
	movl	%edx, %r14d
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	leal	-1(%rsi), %r13d
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%rdi, %rbp
	movl	%esi, %edi
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	imull	%esi, %edi
	movl	%esi, %ebx
	movl	$4, %esi
	pushq	%rcx
	.cfi_def_cfa_offset 64
	leaq	0(,%rdi,4), %r15
	call	calloc@PLT
	testq	%rax, %rax
	jne	.L3
	movq	stderr(%rip), %rdi
	movl	%ebx, %r8d
	movl	%ebx, %ecx
	xorl	%eax, %eax
	leaq	.LC0(%rip), %rdx
	movl	$2, %esi
	call	__fprintf_chk@PLT
	movl	$1, %edi
	call	exit@PLT
.L3:
	movl	%r15d, %r12d
	movq	%r15, %rcx
	movq	%rbp, %rsi
	movq	%rax, %rdi
	movq	%r12, %rdx
	call	__memcpy_chk@PLT
	movss	.LC1(%rip), %xmm1
	movl	%r14d, %edx
	xorl	%r9d, %r9d
	movq	%rax, %r8
.L4:
	cmpq	%r9, %rdx
	je	.L12
	movl	%ebx, %eax
	movl	$1, %esi
.L7:
	cmpl	%r13d, %esi
	jge	.L13
	movl	%eax, %r10d
	leal	(%rbx,%rax), %edi
	incl	%esi
	incl	%eax
	subl	%ebx, %r10d
	movl	$1, %ecx
.L5:
	leal	(%r10,%rcx), %r14d
	leal	(%rdi,%rcx), %r11d
	incl	%ecx
	movss	0(%rbp,%r14,4), %xmm0
	leal	1(%rax), %r14d
	addss	0(%rbp,%r11,4), %xmm0
	leal	-1(%rax), %r11d
	addss	0(%rbp,%r11,4), %xmm0
	addss	0(%rbp,%r14,4), %xmm0
	movl	%eax, %r11d
	movq	%r14, %rax
	mulss	%xmm1, %xmm0
	movss	%xmm0, (%r8,%r11,4)
	cmpl	%ecx, %r13d
	jg	.L5
	movl	%edi, %eax
	jmp	.L7
.L13:
	movq	%rbp, %rdi
	movq	%r8, %rsi
	movq	%r12, %rcx
	incq	%r9
	rep movsb
	jmp	.L4
.L12:
	popq	%rax
	.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
.LFE2028:
	.size	_Z13seq_heat_distPfjj, .-_Z13seq_heat_distPfjj
	.globl	_Z10calcBlocksji
	.type	_Z10calcBlocksji, @function
_Z10calcBlocksji:
.LFB2029:
	.cfi_startproc
	endbr64
	movl	%edi, %eax
	xorl	%edx, %edx
	divl	%esi
	movl	%edi, %eax
	subl	%edx, %eax
	cltd
	idivl	%esi
	incl	%eax
	ret
	.cfi_endproc
.LFE2029:
	.size	_Z10calcBlocksji, .-_Z10calcBlocksji
	.globl	_Z39__device_stub__Z15spread_to_pointPfjjS_PfjjS_
	.type	_Z39__device_stub__Z15spread_to_pointPfjjS_PfjjS_, @function
_Z39__device_stub__Z15spread_to_pointPfjjS_PfjjS_:
.LFB2055:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	56(%rsp), %rdi
	movl	%esi, 20(%rsp)
	leaq	68(%rsp), %rsi
	movl	%edx, 16(%rsp)
	leaq	40(%rsp), %rdx
	movq	%rcx, 8(%rsp)
	leaq	48(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 104(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L15
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z15spread_to_pointPfjjS_(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 176
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	120(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L15:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L17
	call	__stack_chk_fail@PLT
.L17:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2055:
	.size	_Z39__device_stub__Z15spread_to_pointPfjjS_PfjjS_, .-_Z39__device_stub__Z15spread_to_pointPfjjS_PfjjS_
	.globl	_Z15spread_to_pointPfjjS_
	.type	_Z15spread_to_pointPfjjS_, @function
_Z15spread_to_pointPfjjS_:
.LFB2056:
	.cfi_startproc
	endbr64
	jmp	_Z39__device_stub__Z15spread_to_pointPfjjS_PfjjS_
	.cfi_endproc
.LFE2056:
	.size	_Z15spread_to_pointPfjjS_, .-_Z15spread_to_pointPfjjS_
	.globl	_Z51__device_stub__Z27overwrite_current_iterationPfjjS_PfjjS_
	.type	_Z51__device_stub__Z27overwrite_current_iterationPfjjS_PfjjS_, @function
_Z51__device_stub__Z27overwrite_current_iterationPfjjS_PfjjS_:
.LFB2057:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	56(%rsp), %rdi
	movl	%esi, 20(%rsp)
	leaq	68(%rsp), %rsi
	movl	%edx, 16(%rsp)
	leaq	40(%rsp), %rdx
	movq	%rcx, 8(%rsp)
	leaq	48(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 104(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L20
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z27overwrite_current_iterationPfjjS_(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 176
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	120(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L20:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L22
	call	__stack_chk_fail@PLT
.L22:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2057:
	.size	_Z51__device_stub__Z27overwrite_current_iterationPfjjS_PfjjS_, .-_Z51__device_stub__Z27overwrite_current_iterationPfjjS_PfjjS_
	.globl	_Z27overwrite_current_iterationPfjjS_
	.type	_Z27overwrite_current_iterationPfjjS_, @function
_Z27overwrite_current_iterationPfjjS_:
.LFB2058:
	.cfi_startproc
	endbr64
	jmp	_Z51__device_stub__Z27overwrite_current_iterationPfjjS_PfjjS_
	.cfi_endproc
.LFE2058:
	.size	_Z27overwrite_current_iterationPfjjS_, .-_Z27overwrite_current_iterationPfjjS_
	.section	.rodata.str1.1
.LC2:
	.string	"%.5f"
.LC3:
	.string	"\n"
	.text
	.globl	_Z13gpu_heat_distPfjj
	.type	_Z13gpu_heat_distPfjj, @function
_Z13gpu_heat_distPfjj:
.LFB2030:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movl	%edx, %r14d
	movl	$1, %edx
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	xorl	%r13d, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%r14, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movl	%esi, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	imull	%esi, %ebp
	movl	%esi, %ebx
	subq	$88, %rsp
	.cfi_def_cfa_offset 144
	salq	$2, %rbp
	movq	%rdi, (%rsp)
	movq	%rbp, %rsi
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	xorl	%eax, %eax
	leaq	32(%rsp), %rax
	movq	%rax, %rdi
	movq	%rax, 8(%rsp)
	call	cudaMallocManaged@PLT
	leaq	40(%rsp), %rax
	movl	$1, %edx
	movq	%rbp, %rsi
	movq	%rax, %rdi
	movq	%rax, 16(%rsp)
	call	cudaMallocManaged@PLT
	movq	(%rsp), %rsi
	movq	32(%rsp), %rdi
	movq	%rbp, %rdx
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movq	40(%rsp), %rdi
	movq	%rbp, %rdx
	xorl	%esi, %esi
	call	cudaMemset@PLT
	leal	15(%rbx), %eax
	movabsq	$68719476752, %rcx
	movl	$1, 56(%rsp)
	shrl	$4, %eax
	movq	%rcx, 48(%rsp)
	movl	%eax, 60(%rsp)
	movl	%eax, 64(%rsp)
	movslq	%ebx, %rax
	salq	$2, %rax
	movl	$1, 68(%rsp)
	movq	%rax, 24(%rsp)
.L26:
	cmpq	%r14, %r13
	je	.L34
	movl	56(%rsp), %ecx
	movq	48(%rsp), %rdx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	60(%rsp), %rdi
	movl	68(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L27
	movq	40(%rsp), %rcx
	movq	32(%rsp), %rdi
	movl	%r12d, %edx
	movl	%ebx, %esi
	call	_Z39__device_stub__Z15spread_to_pointPfjjS_PfjjS_
.L27:
	call	cudaDeviceSynchronize@PLT
	movl	56(%rsp), %ecx
	movl	68(%rsp), %esi
	xorl	%r9d, %r9d
	movq	48(%rsp), %rdx
	movq	60(%rsp), %rdi
	xorl	%r8d, %r8d
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L28
	movq	40(%rsp), %rcx
	movq	32(%rsp), %rdi
	movl	%r12d, %edx
	movl	%ebx, %esi
	call	_Z51__device_stub__Z27overwrite_current_iterationPfjjS_PfjjS_
.L28:
	call	cudaDeviceSynchronize@PLT
	cmpl	$-10, %ebx
	movq	24(%rsp), %r15
	sbbl	%ebp, %ebp
	andl	$10, %ebp
	incl	%ebp
.L29:
	decl	%ebp
	je	.L35
	movq	32(%rsp), %rax
	leaq	.LC2(%rip), %rsi
	movl	$2, %edi
	cvtss2sd	(%rax,%r15), %xmm0
	movb	$1, %al
	addq	$4, %r15
	call	__printf_chk@PLT
	jmp	.L29
.L35:
	leaq	.LC3(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	incq	%r13
	call	__printf_chk@PLT
	jmp	.L26
.L34:
	movq	32(%rsp), %rsi
	movq	(%rsp), %rdi
	movl	%ebx, %edx
	movl	$2, %ecx
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L32
	call	__stack_chk_fail@PLT
.L32:
	addq	$88, %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
.LFE2030:
	.size	_Z13gpu_heat_distPfjj, .-_Z13gpu_heat_distPfjj
	.section	.rodata.str1.1
.LC4:
	.string	"CPU"
.LC5:
	.string	"GPU"
.LC6:
	.string	"usage: heatdist num  iterations  who\n"
.LC7:
	.string	"num = dimension of the square matrix (50 and up)\n"
.LC8:
	.string	"iterations = number of iterations till stopping (1 and up)\n"
.LC9:
	.string	"who = 0: sequential code on CPU, 1: GPU execution\n"
.LC10:
	.string	" Cannot allocate the %u x %u array\n"
.LC15:
	.string	"Time taken for %s is %lf\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	cmpl	$4, %edi
	je	.L37
	movq	stderr(%rip), %rdi
	leaq	.LC6(%rip), %rdx
	movl	$2, %esi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
	movq	stderr(%rip), %rdi
	leaq	.LC7(%rip), %rdx
	xorl	%eax, %eax
	movl	$2, %esi
	call	__fprintf_chk@PLT
	movq	stderr(%rip), %rdi
	leaq	.LC8(%rip), %rdx
	xorl	%eax, %eax
	movl	$2, %esi
	call	__fprintf_chk@PLT
	movq	stderr(%rip), %rdi
	leaq	.LC9(%rip), %rdx
	xorl	%eax, %eax
	movl	$2, %esi
	call	__fprintf_chk@PLT
	jmp	.L58
.L37:
	movq	24(%rsi), %rdi
	movq	%rsi, %rbp
	call	atoi@PLT
	movq	8(%rbp), %rdi
	movl	%eax, %r13d
	call	atoi@PLT
	movq	16(%rbp), %rdi
	movl	%eax, %ebx
	call	atoi@PLT
	movl	%ebx, %edi
	movl	$4, %esi
	imull	%ebx, %edi
	movl	%eax, %r14d
	call	calloc@PLT
	movss	.LC11(%rip), %xmm0
	movq	%rax, %rbp
	xorl	%eax, %eax
	testq	%rbp, %rbp
	jne	.L38
	movq	stderr(%rip), %rdi
	movl	%ebx, %r8d
	movl	%ebx, %ecx
	leaq	.LC10(%rip), %rdx
	movl	$2, %esi
	call	__fprintf_chk@PLT
.L58:
	movl	$1, %edi
	call	exit@PLT
.L38:
	cmpl	%ebx, %eax
	jnb	.L59
	movss	%xmm0, 0(%rbp,%rax,4)
	incq	%rax
	jmp	.L38
.L59:
	movss	.LC11(%rip), %xmm0
	xorl	%eax, %eax
	xorl	%edx, %edx
.L40:
	cmpl	%edx, %ebx
	je	.L60
	movl	%eax, %ecx
	incl	%edx
	addl	%ebx, %eax
	movss	%xmm0, 0(%rbp,%rcx,4)
	jmp	.L40
.L60:
	leal	-1(%rbx), %eax
	movss	.LC11(%rip), %xmm0
	xorl	%ecx, %ecx
	movl	%eax, %esi
.L42:
	cmpl	%ecx, %edx
	je	.L61
	movl	%esi, %edi
	incl	%ecx
	addl	%ebx, %esi
	movss	%xmm0, 0(%rbp,%rdi,4)
	jmp	.L42
.L61:
	imull	%ebx, %eax
	movss	.LC11(%rip), %xmm0
	xorl	%ecx, %ecx
.L44:
	cmpl	%ecx, %edx
	je	.L62
	leal	(%rcx,%rax), %esi
	incl	%ecx
	movss	%xmm0, 0(%rbp,%rsi,4)
	jmp	.L44
.L62:
	movss	.LC12(%rip), %xmm0
	movl	$10, %edx
.L46:
	movss	%xmm0, 0(%rbp,%rdx,4)
	incq	%rdx
	cmpq	$31, %rdx
	jne	.L46
	movss	.LC13(%rip), %xmm0
	movl	$10, %edx
.L47:
	leal	(%rdx,%rax), %ecx
	incl	%edx
	movss	%xmm0, 0(%rbp,%rcx,4)
	cmpl	$31, %edx
	jne	.L47
	testl	%r13d, %r13d
	jne	.L48
	call	clock@PLT
	movl	%r14d, %edx
	movl	%ebx, %esi
	movq	%rbp, %rdi
	movq	%rax, %r12
	call	_Z13seq_heat_distPfjj
	jmp	.L57
.L48:
	call	clock@PLT
	movl	%r14d, %edx
	movl	%ebx, %esi
	movq	%rbp, %rdi
	movq	%rax, %r12
	call	_Z13gpu_heat_distPfjj
.L57:
	call	clock@PLT
	leaq	.LC5(%rip), %rdx
	movl	$2, %edi
	leaq	.LC15(%rip), %rsi
	subq	%r12, %rax
	testl	%r13d, %r13d
	cvtsi2sdq	%rax, %xmm0
	leaq	.LC4(%rip), %rax
	divsd	.LC14(%rip), %xmm0
	cmove	%rax, %rdx
	movb	$1, %al
	call	__printf_chk@PLT
	movq	%rbp, %rdi
	call	free@PLT
	popq	%rbx
	.cfi_def_cfa_offset 40
	xorl	%eax, %eax
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2027:
	.size	main, .-main
	.section	.rodata.str1.1
.LC16:
	.string	"_Z27overwrite_current_iterationPfjjS_"
.LC17:
	.string	"_Z15spread_to_pointPfjjS_"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2060:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC16(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z27overwrite_current_iterationPfjjS_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC17(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z15spread_to_pointPfjjS_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2060:
	.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
.LC1:
	.long	1048576000
	.align 4
.LC11:
	.long	1116471296
	.align 4
.LC12:
	.long	1120403456
	.align 4
.LC13:
	.long	1125515264
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC14:
	.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 <stdlib.h>
#include <stdio.h>
#include <time.h> 

/* To index element (i,j) of a 2D array stored as 1D */
#define index(i, j, N)  ((i)*(N)) + (j)

/*****************************************************************/

// Function declarations: Feel free to add any functions you want.
void  seq_heat_dist(float *, unsigned int, unsigned int);
void  gpu_heat_dist(float *, unsigned int, unsigned int);


/*****************************************************************/
/**** Do NOT CHANGE ANYTHING in main() function ******/

int main(int argc, char * argv[])
{
  unsigned int N; /* Dimention of NxN matrix */
  int type_of_device = 0; // CPU or GPU
  int iterations = 0;
  int i;
  
  /* The 2D array of points will be treated as 1D array of NxN elements */
  float * playground; 
  
  // to measure time taken by a specific part of the code 
  double time_taken;
  clock_t start, end;
  
  if(argc != 4)
  {
    fprintf(stderr, "usage: heatdist num  iterations  who\n");
    fprintf(stderr, "num = dimension of the square matrix (50 and up)\n");
    fprintf(stderr, "iterations = number of iterations till stopping (1 and up)\n");
    fprintf(stderr, "who = 0: sequential code on CPU, 1: GPU execution\n");
    exit(1);
  }
  
  type_of_device = atoi(argv[3]);
  N = (unsigned int) atoi(argv[1]);
  iterations = (unsigned int) atoi(argv[2]);
 
  
  /* Dynamically allocate NxN array of floats */
  playground = (float *)calloc(N*N, sizeof(float));
  if( !playground )
  {
   fprintf(stderr, " Cannot allocate the %u x %u array\n", N, N);
   exit(1);
  }
  
  /* Initialize it: calloc already initalized everything to 0 */
  // Edge elements to 70F
  for(i = 0; i < N; i++)
    playground[index(0,i,N)] = 70;
    
  for(i = 0; i < N; i++)
    playground[index(i,0,N)] = 70;
  
  for(i = 0; i < N; i++)
    playground[index(i,N-1, N)] = 70;
  
  for(i = 0; i < N; i++)
    playground[index(N-1,i,N)] = 70;
  
  // from (0,10) to (0,30) inclusive are 100F
  for(i = 10; i <= 30; i++)
    playground[index(0,i,N)] = 100;
  
   // from (n-1,10) to (n-1,30) inclusive are 150F
  for(i = 10; i <= 30; i++)
    playground[index(N-1,i,N)] = 150;
  
  if( !type_of_device ) // The CPU sequential version
  { 
    start = clock();
    seq_heat_dist(playground, N, iterations);
    end = clock();
  }
  else  // The GPU version
  {   
     start = clock();
     gpu_heat_dist(playground, N, iterations); 
     end = clock();    
  }
  
  
  time_taken = ((double)(end - start))/ CLOCKS_PER_SEC;
  
  printf("Time taken for %s is %lf\n", type_of_device == 0? "CPU" : "GPU", time_taken);
  
  free(playground);
  
  return 0;

}


/*****************  The CPU sequential version (DO NOT CHANGE THIS) **************/
void  seq_heat_dist(float * playground, unsigned int N, unsigned int iterations)
{
  // Loop indices
  int i, j, k;
  int upper = N-1;
  
  // number of bytes to be copied between array temp and array playground
  unsigned int num_bytes = 0;
  
  float * temp; 
  /* Dynamically allocate another array for temp values */
  /* Dynamically allocate NxN array of floats */
  temp = (float *)calloc(N*N, sizeof(float));
  if( !temp )
  {
   fprintf(stderr, " Cannot allocate temp %u x %u array\n", N, N);
   exit(1);
  }
  
  num_bytes = N*N*sizeof(float);
  
  /* Copy initial array in temp */
  memcpy((void *)temp, (void *) playground, num_bytes);
  
  for( k = 0; k < iterations; k++)
  {
    /* Calculate new values and store them in temp */
    for(i = 1; i < upper; i++)
      for(j = 1; j < upper; j++)
	temp[index(i,j,N)] = (playground[index(i-1,j,N)] + 
	                      playground[index(i+1,j,N)] + 
			      playground[index(i,j-1,N)] + 
			      playground[index(i,j+1,N)])/4.0;
  
			      
   			      
    /* Move new values into old values */ 
    memcpy((void *)playground, (void *) temp, num_bytes);
  }
  
}

/***************** The GPU version: Write your code here *********************/
/* This function can call one or more kernels if you want ********************/

// There will be two main functions that can be parallelized: one to average individual points around each point and one to update the current matrix's points for the next iteration to work with. 

__global__ void spread_to_point(float * current, unsigned int N, unsigned int iter, float * fresh) 
{ // Averages the four surrounding points to update a single point.

  // Let's make a grid-stride with a 2D grid, to fit the problem.

  unsigned int ind_i = blockDim.x * blockIdx.x + threadIdx.x; // Current block and current thread for the i-coord.
  unsigned int ind_j = blockDim.y * blockIdx.y + threadIdx.y; // Current block and current thread for the j-coord.

  int stride_i = blockDim.x * gridDim.x; // Next in line, if needed.
  int stride_j = blockDim.y * gridDim.y;  // Next in line, if needed. 

  for ( int j = ind_j+1 ; j < N-1 ; j += stride_j )
    for ( int i = ind_i+1 ; i < N-1 ; i += stride_i )
    {

    // Uncomment here if you want to take advantage of shared memory. 
    // __shared__ 
    // current[(i-1) * Nm2 + j], 
    // current[(i+1) * Nm2 + j], 
    // current[i * Nm2 + (j-1)], 
    // current[i * Nm2 + (j+1)];

    fresh[i * N + j] = ( // Multiply N by i (the row #) since the input data still represents the matrix as a 1D structure. 
      current[(i-1) * N + j] + 
      current[(i+1) * N + j] + 
      current[i * N + (j-1)] + 
      current[i * N + (j+1)]
      ) / 4;
    }
}

__global__ void overwrite_current_iteration(float * current, unsigned int N, unsigned int iter, float * fresh) 
{ // After computing all values using the old iteration, this function will make the new values take the old values' places. 

  // Again let's make a grid-stride with a 2D grid to fit the problem.

  unsigned int ind_i = blockDim.x * blockIdx.x + threadIdx.x; // Current block and current thread for the i-coord.
  unsigned int ind_j = blockDim.y * blockIdx.y + threadIdx.y; // Current block and current thread for the j-coord.

  int stride_i = blockDim.x * gridDim.x;
  int stride_j = blockDim.y * gridDim.y;   

  int index = ind_i * N + ind_j; 

  for ( int j = ind_j+1 ; j < N-1 ; j += stride_j )
    for ( int i = ind_i+1 ; i < N-1 ; i += stride_i )
      current[index] = fresh[index];
}

// The two commented functions below were anticipating writing parallel code to initialize the mesh temperatures (parallelized instead of sequential), however that is already done sequentially in main(), so the two below are not needed, but may be conceptually useful.  

// __global__ void initialize_edges_or_not(int N, float * grid, int iterations) 
// { // Checks if a point is on the edge or not, setting to 70 and 0 respectively.  This is used only once, to efficiently initialize values.

//   int i = blockDim.x * blockIdx.x +threadIdx.x; // Current block and current thread for the i-coord.
//   int j = blockDim.y * blockIdx.y +threadIdx.y; // Current block and current thread for the j-coord.

//   int index = i * N + j;

//   if ( index == 0 || index == N - 1 ) 
//   {
//     h[index] = 70;
//   }
//   else if ( index%N == 0 ) 
//   {
//     h[index] = 70;
//   } 
//   else {
//     h[index] = 0;
//   }
// }

// __global__ void initialize_special_sections(int N, float * grid, int iterations) 
// { // Can specificy swaths of points to have special values. This is used only once, to efficiently initialize values. Separated into another function to prevent too many redundant condition checks with high N in the init..edge() func. 

//   int i = blockDim.x * blockIdx.x +threadIdx.x; // Current block and current thread for the i-coord.
//   int j = blockDim.y * blockIdx.y +threadIdx.y; // Current block and current thread for the j-coord.

//   int index = i * N + j;
//   if (i==0 || i==N) 
//   {
//     // So on...
//   }
// }

int calcBlocks(unsigned int Nm2, int tpb) 
{
  int remain = Nm2 % tpb;
  int sheared = Nm2 - remain;
  int result = (sheared / tpb) + 1;

  return result; 
}

void  gpu_heat_dist(float * playground, unsigned int N, unsigned int iterations)
{
  float *current, *fresh; 
  hipMallocManaged(&current, N*N*sizeof(float));
  hipMallocManaged(&fresh, N*N*sizeof(float));
  
  hipMemcpy(current, playground, N*N*sizeof(float), hipMemcpyHostToDevice); 
  hipMemset(fresh, 0, N*N*sizeof(int));
  // Tiling setup - how many threads per block, and how many blocks in the grid. 

  unsigned int blkcount = (N + 16 - 1) / 16;
  dim3 threadsPerBlock(16, 16);

  // If you want to control the block calculation, then uncomment here and comment the other numBlocks(...) declaration. 
  // dim3 numBlocks
  // (
  //   calcBlocks(Nm2, threadsPerBlock.x), 
  //   calcBlocks(Nm2, threadsPerBlock.y)
  // );
  dim3 numBlocks(blkcount,blkcount);

  for ( int i = 0 ; i < iterations ; i++ ) 
  {
    spread_to_point<<<numBlocks,threadsPerBlock>>>(current, N, iterations, fresh);
    hipDeviceSynchronize();
    overwrite_current_iteration<<<numBlocks,threadsPerBlock>>>(current, N, iterations, fresh);
    hipDeviceSynchronize();

    int lessAmt = 10; 
    for ( int j = N ; j < N+lessAmt ; j++ )
      printf("%.5f", current[j]);
    printf("\n");
  }
	
  hipMemcpy(playground, current, N*sizeof(float), hipMemcpyDeviceToHost);
  hipFree(&current);
  hipFree(&fresh);


}




	.text
	.file	"heatdist.hip"
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function main
.LCPI0_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	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	cmpl	$4, %edi
	jne	.LBB0_22
# %bb.1:
	movq	%rsi, %rbx
	movq	24(%rsi), %rdi
	callq	atoi
	movl	%eax, %ebp
	movq	8(%rbx), %rdi
	callq	atoi
	movl	%eax, %r14d
	movq	16(%rbx), %rdi
	callq	atoi
	movl	%eax, %r15d
	movl	%r14d, %edi
	imull	%edi, %edi
	movl	$4, %esi
	callq	calloc
	testq	%rax, %rax
	je	.LBB0_23
# %bb.2:                                # %.preheader67
	movq	%rax, %rbx
	testl	%r14d, %r14d
	je	.LBB0_14
# %bb.3:                                # %.lr.ph.preheader
	movl	%r14d, %eax
	xorl	%ecx, %ecx
.LBB0_4:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movl	$1116471296, (%rbx,%rcx,4)      # imm = 0x428C0000
	incq	%rcx
	cmpq	%rcx, %rax
	jne	.LBB0_4
# %bb.5:                                # %.preheader66
	testl	%r14d, %r14d
	je	.LBB0_14
# %bb.6:                                # %.lr.ph70.preheader
	xorl	%ecx, %ecx
	movq	%rax, %rdx
.LBB0_7:                                # %.lr.ph70
                                        # =>This Inner Loop Header: Depth=1
	movl	%ecx, %esi
	movl	$1116471296, (%rbx,%rsi,4)      # imm = 0x428C0000
	addl	%r14d, %ecx
	decq	%rdx
	jne	.LBB0_7
# %bb.8:                                # %.preheader65
	testl	%r14d, %r14d
	je	.LBB0_14
# %bb.9:                                # %.lr.ph72
	leal	-1(%r14), %ecx
	movq	%rax, %rdx
.LBB0_10:                               # =>This Inner Loop Header: Depth=1
	movl	%ecx, %esi
	movl	$1116471296, (%rbx,%rsi,4)      # imm = 0x428C0000
	addl	%r14d, %ecx
	decq	%rdx
	jne	.LBB0_10
# %bb.11:                               # %.preheader64
	testl	%r14d, %r14d
	je	.LBB0_14
# %bb.12:                               # %.lr.ph74
	leal	-1(%r14), %ecx
	imull	%r14d, %ecx
.LBB0_13:                               # =>This Inner Loop Header: Depth=1
	movl	%ecx, %edx
	movl	$1116471296, (%rbx,%rdx,4)      # imm = 0x428C0000
	incl	%ecx
	decq	%rax
	jne	.LBB0_13
.LBB0_14:                               # %.preheader63.preheader
	movl	$10, %eax
.LBB0_15:                               # %.preheader63
                                        # =>This Inner Loop Header: Depth=1
	movl	$1120403456, (%rbx,%rax,4)      # imm = 0x42C80000
	incq	%rax
	cmpq	$31, %rax
	jne	.LBB0_15
# %bb.16:                               # %.preheader
	leal	-1(%r14), %eax
	imull	%r14d, %eax
	addl	$10, %eax
	xorl	%ecx, %ecx
.LBB0_17:                               # =>This Inner Loop Header: Depth=1
	leal	(%rax,%rcx), %edx
	movl	$1125515264, (%rbx,%rdx,4)      # imm = 0x43160000
	incq	%rcx
	cmpq	$21, %rcx
	jne	.LBB0_17
# %bb.18:
	callq	clock
	movq	%rax, %r12
	movq	%rbx, %rdi
	movl	%r14d, %esi
	movl	%r15d, %edx
	testl	%ebp, %ebp
	je	.LBB0_20
# %bb.19:
	callq	_Z13gpu_heat_distPfjj
	movl	$.L.str.7, %r14d
	jmp	.LBB0_21
.LBB0_20:
	callq	_Z13seq_heat_distPfjj
	movl	$.L.str.6, %r14d
.LBB0_21:
	callq	clock
	subq	%r12, %rax
	cvtsi2sd	%rax, %xmm0
	divsd	.LCPI0_0(%rip), %xmm0
	movl	$.L.str.5, %edi
	movq	%r14, %rsi
	movb	$1, %al
	callq	printf
	movq	%rbx, %rdi
	callq	free
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.LBB0_22:
	.cfi_def_cfa_offset 48
	movq	stderr(%rip), %rcx
	movl	$.L.str, %edi
	movl	$37, %esi
	movl	$1, %edx
	callq	fwrite@PLT
	movq	stderr(%rip), %rcx
	movl	$.L.str.1, %edi
	movl	$49, %esi
	movl	$1, %edx
	callq	fwrite@PLT
	movq	stderr(%rip), %rcx
	movl	$.L.str.2, %edi
	movl	$59, %esi
	movl	$1, %edx
	callq	fwrite@PLT
	movq	stderr(%rip), %rcx
	movl	$.L.str.3, %edi
	movl	$50, %esi
	movl	$1, %edx
	callq	fwrite@PLT
	jmp	.LBB0_24
.LBB0_23:
	movq	stderr(%rip), %rdi
	movl	$.L.str.4, %esi
	movl	%r14d, %edx
	movl	%r14d, %ecx
	xorl	%eax, %eax
	callq	fprintf
.LBB0_24:
	movl	$1, %edi
	callq	exit
.Lfunc_end0:
	.size	main, .Lfunc_end0-main
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function _Z13seq_heat_distPfjj
.LCPI1_0:
	.long	0x3e800000                      # float 0.25
	.text
	.globl	_Z13seq_heat_distPfjj
	.type	_Z13seq_heat_distPfjj,@function
_Z13seq_heat_distPfjj:                  # @_Z13seq_heat_distPfjj
	.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	$24, %rsp
	.cfi_def_cfa_offset 80
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	%edx, %ebp
                                        # kill: def $esi killed $esi def $rsi
	movq	%rdi, %r14
	movq	%rsi, %rbx
	movl	%esi, %r15d
	imull	%r15d, %r15d
	movl	$4, %esi
	movq	%r15, %rdi
	callq	calloc
	testq	%rax, %rax
	je	.LBB1_10
# %bb.1:
	movq	%rax, %r12
	shll	$2, %r15d
	movq	%rax, %rdi
	movq	%r14, %rsi
	movq	%r15, 16(%rsp)                  # 8-byte Spill
	movq	%r15, %rdx
	callq	memcpy@PLT
	movl	%ebp, 4(%rsp)                   # 4-byte Spill
	testl	%ebp, %ebp
	je	.LBB1_9
# %bb.2:                                # %.preheader46.lr.ph
	movq	%rbx, %rax
	leal	-1(%rbx), %r13d
	movl	%r13d, %ebx
	movl	%eax, %ebp
	addl	%eax, %eax
	orq	$1, %rax
	movq	%rax, 8(%rsp)                   # 8-byte Spill
	decq	%rbx
	xorl	%r15d, %r15d
	movss	.LCPI1_0(%rip), %xmm1           # xmm1 = mem[0],zero,zero,zero
.LBB1_3:                                # %.preheader46
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB1_5 Depth 2
                                        #       Child Loop BB1_6 Depth 3
	cmpl	$2, %r13d
	jl	.LBB1_8
# %bb.4:                                # %.preheader.preheader
                                        #   in Loop: Header=BB1_3 Depth=1
	movl	$1, %eax
	xorl	%ecx, %ecx
	movq	8(%rsp), %rdx                   # 8-byte Reload
	movq	%rbp, %rsi
.LBB1_5:                                # %.preheader
                                        #   Parent Loop BB1_3 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB1_6 Depth 3
	incl	%eax
	xorl	%edi, %edi
.LBB1_6:                                #   Parent Loop BB1_3 Depth=1
                                        #     Parent Loop BB1_5 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	leal	(%rcx,%rdi), %r8d
	incl	%r8d
	movss	(%r14,%r8,4), %xmm0             # xmm0 = mem[0],zero,zero,zero
	leal	(%rdx,%rdi), %r8d
	addss	(%r14,%r8,4), %xmm0
	leal	(%rsi,%rdi), %r8d
	addss	(%r14,%r8,4), %xmm0
	leal	2(%rsi,%rdi), %r8d
	addss	(%r14,%r8,4), %xmm0
	mulss	%xmm1, %xmm0
	leal	(%rsi,%rdi), %r8d
	incl	%r8d
	movss	%xmm0, (%r12,%r8,4)
	incq	%rdi
	cmpq	%rdi, %rbx
	jne	.LBB1_6
# %bb.7:                                # %._crit_edge
                                        #   in Loop: Header=BB1_5 Depth=2
	addq	%rbp, %rsi
	addq	%rbp, %rdx
	addq	%rbp, %rcx
	cmpl	%r13d, %eax
	jne	.LBB1_5
.LBB1_8:                                # %._crit_edge49
                                        #   in Loop: Header=BB1_3 Depth=1
	movq	%r14, %rdi
	movq	%r12, %rsi
	movq	16(%rsp), %rdx                  # 8-byte Reload
	callq	memcpy@PLT
	movss	.LCPI1_0(%rip), %xmm1           # xmm1 = mem[0],zero,zero,zero
	incl	%r15d
	cmpl	4(%rsp), %r15d                  # 4-byte Folded Reload
	jne	.LBB1_3
.LBB1_9:                                # %._crit_edge51
	addq	$24, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.LBB1_10:
	.cfi_def_cfa_offset 80
	movq	stderr(%rip), %rdi
	movl	$.L.str.8, %esi
	movq	%rbx, %rcx
	movl	%ecx, %edx
                                        # kill: def $ecx killed $ecx killed $rcx
	xorl	%eax, %eax
	callq	fprintf
	movl	$1, %edi
	callq	exit
.Lfunc_end1:
	.size	_Z13seq_heat_distPfjj, .Lfunc_end1-_Z13seq_heat_distPfjj
	.cfi_endproc
                                        # -- End function
	.globl	_Z13gpu_heat_distPfjj           # -- Begin function _Z13gpu_heat_distPfjj
	.type	_Z13gpu_heat_distPfjj,@function
_Z13gpu_heat_distPfjj:                  # @_Z13gpu_heat_distPfjj
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$40, %rsp
	.cfi_def_cfa_offset 96
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	%edx, %ebp
                                        # kill: def $esi killed $esi def $rsi
	movq	%rdi, %rbx
	movq	%rsi, (%rsp)                    # 8-byte Spill
	movl	%esi, %r15d
	imull	%r15d, %r15d
	shlq	$2, %r15
	leaq	8(%rsp), %r12
	movq	%r12, %rdi
	movq	%r15, %rsi
	movl	$1, %edx
	callq	hipMallocManaged
	leaq	16(%rsp), %r13
	movq	%r13, %rdi
	movq	%r15, %rsi
	movl	$1, %edx
	callq	hipMallocManaged
	movq	(%r12), %rdi
	movq	%rbx, 24(%rsp)                  # 8-byte Spill
	movq	%rbx, %rsi
	movq	%r15, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%r13), %rdi
	xorl	%r13d, %r13d
	xorl	%esi, %esi
	movq	%r15, %rdx
	callq	hipMemset
	testl	%ebp, %ebp
	je	.LBB2_8
# %bb.1:                                # %.lr.ph52
	movl	%ebp, %r14d
	movq	(%rsp), %rcx                    # 8-byte Reload
	leal	15(%rcx), %eax
	shrl	$4, %eax
	movq	%rax, %r15
	shlq	$32, %r15
	orq	%rax, %r15
	movabsq	$68719476752, %r12              # imm = 0x1000000010
	leal	10(%rcx), %ebx
	movslq	%ecx, %rax
	movq	%rax, 32(%rsp)                  # 8-byte Spill
.LBB2_2:                                # =>This Loop Header: Depth=1
                                        #     Child Loop BB2_9 Depth 2
	movq	%r15, %rdi
	movl	$1, %esi
	movq	%r12, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB2_4
# %bb.3:                                #   in Loop: Header=BB2_2 Depth=1
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rcx
	movq	(%rsp), %rsi                    # 8-byte Reload
                                        # kill: def $esi killed $esi killed $rsi
	movl	%r14d, %edx
	callq	_Z30__device_stub__spread_to_pointPfjjS_
.LBB2_4:                                #   in Loop: Header=BB2_2 Depth=1
	callq	hipDeviceSynchronize
	movq	%r15, %rdi
	movl	$1, %esi
	movq	%r12, %rbp
	movq	%r12, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB2_6
# %bb.5:                                #   in Loop: Header=BB2_2 Depth=1
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rcx
	movq	(%rsp), %rsi                    # 8-byte Reload
                                        # kill: def $esi killed $esi killed $rsi
	movl	%r14d, %edx
	callq	_Z42__device_stub__overwrite_current_iterationPfjjS_
.LBB2_6:                                #   in Loop: Header=BB2_2 Depth=1
	callq	hipDeviceSynchronize
	movq	32(%rsp), %r12                  # 8-byte Reload
	cmpl	$-11, (%rsp)                    # 4-byte Folded Reload
	ja	.LBB2_7
.LBB2_9:                                # %.lr.ph
                                        #   Parent Loop BB2_2 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movq	8(%rsp), %rax
	xorps	%xmm0, %xmm0
	cvtss2sd	(%rax,%r12,4), %xmm0
	movl	$.L.str.9, %edi
	movb	$1, %al
	callq	printf
	incq	%r12
	cmpl	%r12d, %ebx
	ja	.LBB2_9
.LBB2_7:                                # %._crit_edge
                                        #   in Loop: Header=BB2_2 Depth=1
	movl	$10, %edi
	callq	putchar@PLT
	incl	%r13d
	cmpl	%r14d, %r13d
	movq	%rbp, %r12
	jne	.LBB2_2
.LBB2_8:                                # %._crit_edge53
	leaq	8(%rsp), %r15
	movq	(%r15), %rsi
	movl	(%rsp), %edx                    # 4-byte Reload
	shlq	$2, %rdx
	movq	24(%rsp), %rdi                  # 8-byte Reload
	movl	$2, %ecx
	callq	hipMemcpy
	movq	%r15, %rdi
	callq	hipFree
	leaq	16(%rsp), %rdi
	callq	hipFree
	addq	$40, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end2:
	.size	_Z13gpu_heat_distPfjj, .Lfunc_end2-_Z13gpu_heat_distPfjj
	.cfi_endproc
                                        # -- End function
	.globl	_Z30__device_stub__spread_to_pointPfjjS_ # -- Begin function _Z30__device_stub__spread_to_pointPfjjS_
	.type	_Z30__device_stub__spread_to_pointPfjjS_,@function
_Z30__device_stub__spread_to_pointPfjjS_: # @_Z30__device_stub__spread_to_pointPfjjS_
	.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	$112, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	12(%rsp), %rdi
	movl	%esi, (%rdi)
	leaq	8(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	24(%rsp), %r12
	leaq	16(%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	$_Z15spread_to_pointPfjjS_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	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_end3:
	.size	_Z30__device_stub__spread_to_pointPfjjS_, .Lfunc_end3-_Z30__device_stub__spread_to_pointPfjjS_
	.cfi_endproc
                                        # -- End function
	.globl	_Z42__device_stub__overwrite_current_iterationPfjjS_ # -- Begin function _Z42__device_stub__overwrite_current_iterationPfjjS_
	.type	_Z42__device_stub__overwrite_current_iterationPfjjS_,@function
_Z42__device_stub__overwrite_current_iterationPfjjS_: # @_Z42__device_stub__overwrite_current_iterationPfjjS_
	.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	$112, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	12(%rsp), %rdi
	movl	%esi, (%rdi)
	leaq	8(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	24(%rsp), %r12
	leaq	16(%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	$_Z27overwrite_current_iterationPfjjS_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end4:
	.size	_Z42__device_stub__overwrite_current_iterationPfjjS_, .Lfunc_end4-_Z42__device_stub__overwrite_current_iterationPfjjS_
	.cfi_endproc
                                        # -- End function
	.globl	_Z10calcBlocksji                # -- Begin function _Z10calcBlocksji
	.type	_Z10calcBlocksji,@function
_Z10calcBlocksji:                       # @_Z10calcBlocksji
	.cfi_startproc
# %bb.0:
	movl	%edi, %eax
	xorl	%edx, %edx
	divl	%esi
	subl	%edx, %edi
	movl	%edi, %eax
	cltd
	idivl	%esi
	incl	%eax
	retq
.Lfunc_end5:
	.size	_Z10calcBlocksji, .Lfunc_end5-_Z10calcBlocksji
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -16
	movq	__hip_gpubin_handle(%rip), %rbx
	testq	%rbx, %rbx
	jne	.LBB6_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB6_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z15spread_to_pointPfjjS_, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z27overwrite_current_iterationPfjjS_, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end6:
	.size	__hip_module_ctor, .Lfunc_end6-__hip_module_ctor
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_dtor
	.type	__hip_module_dtor,@function
__hip_module_dtor:                      # @__hip_module_dtor
	.cfi_startproc
# %bb.0:
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB7_2
# %bb.1:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__hipUnregisterFatBinary
	movq	$0, __hip_gpubin_handle(%rip)
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB7_2:
	retq
.Lfunc_end7:
	.size	__hip_module_dtor, .Lfunc_end7-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"usage: heatdist num  iterations  who\n"
	.size	.L.str, 38

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"num = dimension of the square matrix (50 and up)\n"
	.size	.L.str.1, 50

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"iterations = number of iterations till stopping (1 and up)\n"
	.size	.L.str.2, 60

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"who = 0: sequential code on CPU, 1: GPU execution\n"
	.size	.L.str.3, 51

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	" Cannot allocate the %u x %u array\n"
	.size	.L.str.4, 36

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"Time taken for %s is %lf\n"
	.size	.L.str.5, 26

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"CPU"
	.size	.L.str.6, 4

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

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	" Cannot allocate temp %u x %u array\n"
	.size	.L.str.8, 37

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

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

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

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z15spread_to_pointPfjjS_"
	.size	.L__unnamed_1, 26

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z27overwrite_current_iterationPfjjS_"
	.size	.L__unnamed_2, 38

	.type	__hip_fatbin_wrapper,@object    # @__hip_fatbin_wrapper
	.section	.hipFatBinSegment,"a",@progbits
	.p2align	3, 0x0
__hip_fatbin_wrapper:
	.long	1212764230                      # 0x48495046
	.long	1                               # 0x1
	.quad	__hip_fatbin
	.quad	0
	.size	__hip_fatbin_wrapper, 24

	.type	__hip_gpubin_handle,@object     # @__hip_gpubin_handle
	.local	__hip_gpubin_handle
	.comm	__hip_gpubin_handle,8,8
	.section	.init_array,"aw",@init_array
	.p2align	3, 0x0
	.quad	__hip_module_ctor
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.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__spread_to_pointPfjjS_
	.addrsig_sym _Z42__device_stub__overwrite_current_iterationPfjjS_
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z15spread_to_pointPfjjS_
	.addrsig_sym _Z27overwrite_current_iterationPfjjS_
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z15spread_to_pointPfjjS_ ; -- Begin function _Z15spread_to_pointPfjjS_
	.globl	_Z15spread_to_pointPfjjS_
	.p2align	8
	.type	_Z15spread_to_pointPfjjS_,@function
_Z15spread_to_pointPfjjS_:              ; @_Z15spread_to_pointPfjjS_
; %bb.0:
	s_clause 0x1
	s_load_b32 s4, s[0:1], 0x24
	s_load_b32 s9, s[0:1], 0x8
	s_add_u32 s2, s0, 24
	v_bfe_u32 v2, v0, 10, 10
	s_addc_u32 s3, s1, 0
	s_mov_b32 s5, exec_lo
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s8, s4, 16
	s_add_i32 s6, s9, -1
	s_mul_i32 s15, s15, s8
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add3_u32 v3, v2, s15, 1
	v_cmpx_gt_u32_e64 s6, v3
	s_cbranch_execz .LBB0_6
; %bb.1:                                ; %.lr.ph41
	v_and_b32_e32 v0, 0x3ff, v0
	s_and_b32 s7, s4, 0xffff
	s_load_b64 s[10:11], s[2:3], 0x0
	s_clause 0x1
	s_load_b64 s[2:3], s[0:1], 0x0
	s_load_b64 s[4:5], s[0:1], 0x10
	v_mad_u64_u32 v[5:6], null, s14, s7, v[0:1]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_add_nc_u32_e32 v0, 2, v5
	v_add_nc_u32_e32 v4, 1, v5
	v_mul_lo_u32 v5, s9, v5
	v_mul_lo_u32 v6, s9, v0
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_3) | instid1(VALU_DEP_4)
	v_mad_u64_u32 v[0:1], null, s9, v4, v[2:3]
	v_cmp_gt_u32_e32 vcc_lo, s6, v4
	s_waitcnt lgkmcnt(0)
	s_mul_i32 s7, s10, s7
	v_add3_u32 v5, v2, v5, 1
	s_mul_i32 s8, s11, s8
	s_mul_i32 s9, s7, s9
	v_add3_u32 v6, v2, v6, 1
	v_mov_b32_e32 v2, 0
	s_mov_b32 s10, 0
.LBB0_2:                                ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_4 Depth 2
	s_and_saveexec_b32 s11, vcc_lo
	s_cbranch_execz .LBB0_5
; %bb.3:                                ; %.lr.ph
                                        ;   in Loop: Header=BB0_2 Depth=1
	v_dual_mov_b32 v7, v5 :: v_dual_mov_b32 v8, v6
	v_mov_b32_e32 v9, v0
	v_mov_b32_e32 v10, v4
	s_mov_b32 s12, 0
.LBB0_4:                                ;   Parent Loop BB0_2 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_dual_mov_b32 v12, v2 :: v_dual_add_nc_u32 v1, s15, v7
	v_add_nc_u32_e32 v7, s9, v7
	v_lshlrev_b64 v[13:14], 2, v[1:2]
	v_add_nc_u32_e32 v1, s15, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_lshlrev_b64 v[15:16], 2, v[1:2]
	v_add_nc_u32_e32 v1, s15, v9
	v_add_co_u32 v13, s0, s2, v13
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e64 v14, s0, s3, v14, s0
	v_add_nc_u32_e32 v11, 2, v1
	v_add_co_u32 v15, s0, s2, v15
	v_lshlrev_b64 v[17:18], 2, v[1:2]
	v_add_co_ci_u32_e64 v16, s0, s3, v16, s0
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_4) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[11:12], 2, v[11:12]
	global_load_b32 v19, v[13:14], off
	v_add_nc_u32_e32 v1, 1, v1
	global_load_b32 v15, v[15:16], off
	v_add_co_u32 v13, s0, s2, v17
	v_add_co_ci_u32_e64 v14, s0, s3, v18, s0
	v_add_co_u32 v11, s0, s2, v11
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e64 v12, s0, s3, v12, s0
	s_clause 0x1
	global_load_b32 v13, v[13:14], off
	global_load_b32 v14, v[11:12], off
	v_add_nc_u32_e32 v10, s7, v10
	v_add_nc_u32_e32 v9, s9, v9
	s_waitcnt vmcnt(2)
	v_dual_add_f32 v11, v19, v15 :: v_dual_add_nc_u32 v8, s9, v8
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_le_u32_e64 s0, s6, v10
	s_or_b32 s12, s0, s12
	s_waitcnt vmcnt(1)
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_add_f32_e32 v13, v11, v13
	v_lshlrev_b64 v[11:12], 2, v[1:2]
	s_waitcnt vmcnt(0)
	v_add_f32_e32 v1, v13, v14
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_co_u32 v11, s1, s4, v11
	v_add_co_ci_u32_e64 v12, s1, s5, v12, s1
	s_delay_alu instid0(VALU_DEP_3)
	v_mul_f32_e32 v1, 0x3e800000, v1
	global_store_b32 v[11:12], v1, off
	s_and_not1_b32 exec_lo, exec_lo, s12
	s_cbranch_execnz .LBB0_4
.LBB0_5:                                ; %Flow60
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s11
	v_add_nc_u32_e32 v3, s8, v3
	v_add_nc_u32_e32 v0, s8, v0
	v_add_nc_u32_e32 v6, s8, v6
	v_add_nc_u32_e32 v5, s8, v5
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_le_u32_e64 s0, s6, v3
	s_or_b32 s10, s0, s10
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s10
	s_cbranch_execnz .LBB0_2
.LBB0_6:                                ; %Flow62
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z15spread_to_pointPfjjS_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 280
		.amdhsa_user_sgpr_count 14
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 1
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 1
		.amdhsa_next_free_vgpr 20
		.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	_Z15spread_to_pointPfjjS_, .Lfunc_end0-_Z15spread_to_pointPfjjS_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 596
; NumSgprs: 18
; NumVgprs: 20
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 2
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 20
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 14
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 1
	.text
	.protected	_Z27overwrite_current_iterationPfjjS_ ; -- Begin function _Z27overwrite_current_iterationPfjjS_
	.globl	_Z27overwrite_current_iterationPfjjS_
	.p2align	8
	.type	_Z27overwrite_current_iterationPfjjS_,@function
_Z27overwrite_current_iterationPfjjS_:  ; @_Z27overwrite_current_iterationPfjjS_
; %bb.0:
	s_clause 0x1
	s_load_b32 s4, s[0:1], 0x24
	s_load_b32 s6, s[0:1], 0x8
	v_bfe_u32 v3, v0, 10, 10
	s_add_u32 s2, s0, 24
	s_addc_u32 s3, s1, 0
	s_mov_b32 s7, exec_lo
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s5, s4, 16
	s_add_i32 s4, s6, -1
	v_mad_u64_u32 v[1:2], null, s15, s5, v[3:4]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v4, 1, v1
	v_cmpx_gt_u32_e64 s4, v4
	s_cbranch_execz .LBB1_5
; %bb.1:                                ; %.lr.ph28
	s_load_b32 s7, s[2:3], 0xc
	v_and_b32_e32 v0, 0x3ff, v0
	s_load_b64 s[8:9], s[0:1], 0x10
	s_waitcnt lgkmcnt(0)
	s_and_b32 s7, s7, 0xffff
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_mad_u64_u32 v[2:3], null, s14, s7, v[0:1]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_4) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[5:6], null, v2, s6, v[1:2]
	s_load_b64 s[6:7], s[0:1], 0x0
	s_load_b32 s1, s[2:3], 0x4
	v_add_nc_u32_e32 v0, 1, v2
	s_mov_b32 s2, 0
	v_cmp_gt_u32_e32 vcc_lo, s4, v0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v6, 31, v5
	v_lshlrev_b64 v[2:3], 2, v[5:6]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_co_u32 v0, s0, s8, v2
	v_add_co_ci_u32_e64 v1, s0, s9, v3, s0
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v2, s0, s6, v2
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e64 v3, s0, s7, v3, s0
	s_mul_i32 s1, s1, s5
.LBB1_2:                                ; =>This Inner Loop Header: Depth=1
	s_and_saveexec_b32 s0, vcc_lo
	s_cbranch_execz .LBB1_4
; %bb.3:                                ; %.lr.ph.preheader
                                        ;   in Loop: Header=BB1_2 Depth=1
	global_load_b32 v5, v[0:1], off
	s_waitcnt vmcnt(0)
	global_store_b32 v[2:3], v5, off
.LBB1_4:                                ; %._crit_edge
                                        ;   in Loop: Header=BB1_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s0
	v_add_nc_u32_e32 v4, s1, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_le_u32_e64 s0, s4, v4
	s_or_b32 s2, s0, s2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s2
	s_cbranch_execnz .LBB1_2
.LBB1_5:                                ; %Flow35
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z27overwrite_current_iterationPfjjS_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 280
		.amdhsa_user_sgpr_count 14
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 1
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 1
		.amdhsa_next_free_vgpr 7
		.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_end1:
	.size	_Z27overwrite_current_iterationPfjjS_, .Lfunc_end1-_Z27overwrite_current_iterationPfjjS_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 308
; NumSgprs: 18
; NumVgprs: 7
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 7
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 14
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 1
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .offset:         8
        .size:           4
        .value_kind:     by_value
      - .offset:         12
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         28
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         36
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         38
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         40
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         42
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         44
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         46
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         88
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 280
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z15spread_to_pointPfjjS_
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z15spread_to_pointPfjjS_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     20
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .offset:         8
        .size:           4
        .value_kind:     by_value
      - .offset:         12
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         28
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         36
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         38
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         40
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         42
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         44
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         46
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         88
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 280
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z27overwrite_current_iterationPfjjS_
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z27overwrite_current_iterationPfjjS_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     7
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
