/* Matrix normalization.
 * Compile with "gcc matrixNorm.c"
 */

/* ****** ADD YOUR CODE AT THE END OF THIS FILE. ******
 * You need not submit the provided code.
 */

#include <stdio.h>
#include <stdlib.h>
#include <unistd.h>
#include <math.h>
#include <sys/types.h>
#include <sys/times.h>
#include <sys/time.h>
#include <time.h>

/* Program Parameters */
#define MAXN 8000  /* Max value of N */
int N;  /* Matrix size */

/* Matrices */
volatile float A[MAXN][MAXN], B[MAXN][MAXN];

/* junk */
#define randm() 4|2[uid]&3

/* Prototype */
void matrixNorm_GPU();

__global__ void blank_function();

/* returns a seed for srand based on the time */
unsigned int time_seed() {
  struct timeval t;
  struct timezone tzdummy;

  gettimeofday(&t, &tzdummy);
  return (unsigned int)(t.tv_usec);
}

/* Set the program parameters from the command-line arguments */
void parameters(int argc, char **argv) {
  int seed = 0;  /* Random seed */
  char uid[32]; /*User name */

  /* Read command-line arguments */
  srand(time_seed());  /* Randomize */

  if (argc == 3) {
    seed = atoi(argv[2]);
    srand(seed);
    printf("Random seed = %i\n", seed);
  }
  if (argc >= 2) {
    N = atoi(argv[1]);
    if (N < 1 || N > MAXN) {
      printf("N = %i is out of range.\n", N);
      exit(0);
    }
  }
  else {
    printf("Usage: %s <matrix_dimension> [random seed]\n",
           argv[0]);
    exit(0);
  }

  /* Print parameters */
  printf("\nMatrix dimension N = %i.\n", N);
}

/* Initialize A and B*/
void initialize_inputs() {
  int row, col;

  printf("\nInitializing...\n");
  for (col = 0; col < N; col++) {
    for (row = 0; row < N; row++) {
      A[row][col] = (float)rand() / 32768.0;
      B[row][col] = 0.0;
    }
  }

}

/* Print input matrices */
void print_inputs() {
  int row, col;

  if (N < 10) {
    printf("\nA =\n\t");
    for (row = 0; row < N; row++) {
      for (col = 0; col < N; col++) {
	    printf("%5.2f%s", A[row][col], (col < N-1) ? ", " : ";\n\t");
      }
    }
  }
}

void print_B() {
    int row, col;

    if (N < 10) {
        printf("\nB =\n\t");
        for (row = 0; row < N; row++) {
            for (col = 0; col < N; col++) {
                printf("%1.10f%s", B[row][col], (col < N-1) ? ", " : ";\n\t");
            }
        }
    }
}

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

  blank_function<<<1, 1>>>(); /* Blank Function (defined below) having no content which is simply used to initialize CUDA environment as it is implicitly initialized on first call and this delay is not caused by parallelized algorithm.
                                 For reference, follow this link - https://devtalk.nvidia.com/default/topic/392429/first-cudamalloc-takes-long-time-/ */
  /* Timing variables */
  struct timeval etstart, etstop;  /* Elapsed times using gettimeofday() */
  struct timezone tzdummy;
  clock_t etstart2, etstop2;  /* Elapsed times using times() */
  unsigned long long usecstart, usecstop;
  struct tms cputstart, cputstop;  /* CPU times for my processes */

  /* Process program parameters */
  parameters(argc, argv);

  /* Initialize A and B */
  initialize_inputs();

  /* Print input matrices */
  print_inputs();

  /* Start Clock */
  printf("\nStarting clock.\n");
  gettimeofday(&etstart, &tzdummy);
  etstart2 = times(&cputstart);

  /* Gaussian Elimination */
  matrixNorm_GPU();

  /* Stop Clock */
  gettimeofday(&etstop, &tzdummy);
  etstop2 = times(&cputstop);
  printf("Stopped clock.\n");
  usecstart = (unsigned long long)etstart.tv_sec * 1000000 + etstart.tv_usec;
  usecstop = (unsigned long long)etstop.tv_sec * 1000000 + etstop.tv_usec;

  /* Display output */
  print_B();

  /* Display timing results */
  printf("\nElapsed time = %g ms.\n",
	 (float)(usecstop - usecstart)/(float)1000);

  printf("(CPU times are accurate to the nearest %g ms)\n",
	 1.0/(float)CLOCKS_PER_SEC * 1000.0);
  printf("My total CPU time for parent = %g ms.\n",
	 (float)( (cputstop.tms_utime + cputstop.tms_stime) -
		  (cputstart.tms_utime + cputstart.tms_stime) ) /
	 (float)CLOCKS_PER_SEC * 1000);
  printf("My system CPU time for parent = %g ms.\n",
	 (float)(cputstop.tms_stime - cputstart.tms_stime) /
	 (float)CLOCKS_PER_SEC * 1000);
  printf("My total CPU time for child processes = %g ms.\n",
	 (float)( (cputstop.tms_cutime + cputstop.tms_cstime) -
		  (cputstart.tms_cutime + cputstart.tms_cstime) ) /
	 (float)CLOCKS_PER_SEC * 1000);
      /* Contrary to the man pages, this appears not to include the parent */
  printf("--------------------------------------------\n");

  exit(0);
}

/* ------------------ Above Was Provided --------------------- */

/****** You will replace this routine with your own parallel version *******/
/* Provided global variables are MAXN, N, A[][] and B[][],
 * defined in the beginning of this code.  B[][] is initialized to zeros.
 */

/* Blank function which is called at first as the first call to any CUDA kernel or function implicitly initializes the CUDA environment.
For reference, follow this link - https://devtalk.nvidia.com/default/topic/392429/first-cudamalloc-takes-long-time-/ */
__global__ void blank_function()
{

}

/* Computes the sums of each group of numbers whose size is denoted by fragSize. */
__global__ void computeSums(float *d_A, float *d_B, size_t pitch_A, size_t pitch_B, int n, int fullBlock, int fragSize, int lastBlockStartRow)
{
    /* Initializing shared memory which will hold the numbers each thread block will be summing. */
    extern __shared__ float chunk[];

    /* Variables to keep track of which number the given is operating on in a 2D perspective. */
    int tx, ty_base, i;
    tx = blockIdx.x;

    /* Calculating the y index of the first element of the numbers being handled by the thread block. */
    if (fullBlock == 0)
    {
        ty_base = threadIdx.x * fragSize + lastBlockStartRow;
    }
    else
    {
        ty_base = (threadIdx.x + blockIdx.y * blockDim.x) * fragSize;
    }

    float *bElem, *aElem;

    /* Getting the starting index of the number the given thread begins operating from on its fragment. This is relative to the chunk[] array. */
    int localStartIndex = threadIdx.x * fragSize;

    /* Copying elements which the block will be summing from global to shared memory of the thread block. */
    for (i = 0; i < fragSize; i++)
    {
        aElem = (float*)((char*)d_A + (pitch_A * (ty_base + i)));
        chunk[localStartIndex + i] = aElem[tx];
    }

    /* Summing up the elements in the block. */
    if (ty_base < (n - 1))
    {
        int inc = 1;
        while(inc < n)
        {
            if ((localStartIndex % (inc * fragSize)) == 0)
            {
                for (i = localStartIndex + inc; (i < localStartIndex + fragSize * inc) && (i < (blockDim.x * fragSize)) && (i < n); i += inc)
                {
                    chunk[localStartIndex] += chunk[i];
                }
                inc *= fragSize;
            }
            else
                break;
            __syncthreads();
        }
    }

    /* Copying result to array d_B on the GPU. */
    for (i = 0; i < fragSize; i++)
    {
        bElem = (float*)((char*)d_B + (pitch_B * (ty_base + i)));
        bElem[tx] = chunk[localStartIndex + i];
    }
}

/* Computes Average Values of the columns in the input 2D array flattened as d_B. */
__global__ void computeFinalSums(float *d_B, size_t pitch_B, int n, int ifLastBlock, int prevBlockSize, int lastBlockStartCol, float *d_Avgs)
{
    int tx, i;
    float *bElem, *bElemBase;

    /* Getting address of starting element of the column which the thread is handling. */
    bElemBase = (float *)((char *)d_B);

    /* Getting x position of the starting element of the column on which the thread begins work. */
    if (ifLastBlock == 0)
    {
        tx = blockIdx.x * blockDim.x + threadIdx.x;
    }
    else
    {
        tx = lastBlockStartCol + threadIdx.x;
    }

    /* Summing up to get final sums for each column and subsequently calculating average. */
    for (i = prevBlockSize; i < n; i += prevBlockSize)
    {
        bElem = (float *)((char *)d_B + (pitch_B * i));
        bElemBase[tx] += bElem[tx];
    }
    bElemBase[tx] /= n;
    d_Avgs[tx] = bElemBase[tx];
}

/* Computes the Variances of each column of input flattened 2D array d_A. */
__global__ void computeVarianceSquares(float *d_A, size_t pitch_A, float *d_Avgs, int n, int ifLastBlock, int lastBlockStartRow)
{
    int tx, ty;
    tx = blockIdx.x;

    /* Getting y position of the starting element of the column on which the thread begins work. */
    if (ifLastBlock == 1)
    {
        ty = threadIdx.x + lastBlockStartRow;
    }
    else
    {
        ty = threadIdx.x + blockIdx.y * blockDim.x;
    }

    /* Getting address of starting element of the column which the thread is handling. */
    float *aElem = (float *)((char *)d_A + (pitch_A * ty));

    /* Calculating and storing the variances in global memory. */
    aElem[tx] = pow((aElem[tx] - d_Avgs[tx]), 2);
}

/* Finally computes the matrix norms for each element in the input flattened 2D matrix d_A. */
__global__ void computeNorms(float *d_A, size_t pitch_A, float *d_Avgs, float *d_Vars, int n, int ifLastBlock, int lastBlockStartRow)
{
    int tx, ty;
    tx = blockIdx.x;

    /* Getting y position of the element on which the thread begins work. */
    if (ifLastBlock == 1)
    {
        ty = threadIdx.x + lastBlockStartRow;
    }
    else
    {
        ty = threadIdx.x + blockIdx.y * blockDim.x;
    }

    float *aElem = (float *)((char *)d_A + (pitch_A * ty));

    /* Calculating and finally storing the norm value of each element. */
    if (d_Vars[tx] != 0)
        aElem[tx] = (aElem[tx] - d_Avgs[tx]) / sqrt(d_Vars[tx]);
    else
        aElem[tx] = 0;
}

/* Self Implemented ceiling function. */
__host__ __device__ int ceil_h_d(float f)
{
    int tmp = (int) f;
	if (f > tmp)
		tmp++;
	return tmp;
}

/* Function to calculate the normalized values of a matrix, implemented on GPU. */
void matrixNorm_GPU()
{
    /* Getting Device Properties of GPU to scale performance according to GPU specs. */
    cudaDeviceProp prop;
    cudaGetDeviceProperties(&prop, 0);

    int numThreadsPerMP = prop.maxThreadsPerMultiProcessor; /* Getting max. number of concurrent threads each Streaming Multiprocessor in the GPU can run. This is a multiple of the warp size which is 32.*/
    int fragSize = 10; /* The number of elements which are summed by a thread. */
    int BLOCKS_PER_MP = 8; /* Number of threadBlocks each Streaming Multiprocessor supports. */

    int i, j;
    int fullblockSize = numThreadsPerMP / BLOCKS_PER_MP; /* Computing number of threads per thread Block shall run. It should be a multiple of the warp size for maximum efficiency and memory coalescing. */

    int numElemsCol = ceil_h_d((float) N / (float) fragSize); /* Getting number of summing blocks per column */

    int blocksReqdPerCol = ceil_h_d((float) numElemsCol / (float) fullblockSize); /* Getting number of thread blocks required to sum each column. */
    int lastBlockSize = numElemsCol - (blocksReqdPerCol - 1) * fullblockSize; /* Since thread block sizes are multiples of warp size, therefore the last blocksize if not a multiple of warp size, is handled by this thread block. */

    int lastBlockStartRow = (blocksReqdPerCol - 1) * fullblockSize * fragSize; /* Getting the starting row index of the chunk of array being handled by the final thread block */

    float *d_A, *d_Avgs, *d_Vars; /* Declaring device pointers */

    size_t dev_pitch_A;
    size_t host_pitch = N * sizeof(float);

    /* Allocating memory for 2D arrays on GPU */
    cudaMallocPitch(&d_A, &dev_pitch_A, N * sizeof(float), N * sizeof(float));

    /* Allocating memory for 1D arrays on GPU */
    cudaMalloc((void **)&d_Avgs, N * sizeof(float));
    cudaMalloc((void **)&d_Vars, N * sizeof(float));

    /* Allocating memory for flattened 1D array on Host */
    float *A_flat = (float *)malloc(N * N * sizeof(float));

    /* Flattening 2D array on host to a 1D array for transfer to GPU */
    for (i = 0; i < N; i++)
    {
        for (j = 0; j < N; j++)
        {
            A_flat[j + i * N] = A[i][j];
        }
    }

    /* Copying flattened 2D array A from host to GPU */
    if (cudaMemcpy2D(d_A, dev_pitch_A, A_flat, host_pitch, N * sizeof(float), N, cudaMemcpyHostToDevice)!= cudaSuccess)
        printf("ERROR");

    dim3 numFullBlocks(N, (blocksReqdPerCol - 1));

    /* Parallel computation of sums of each sum fragment in each column. */
    computeSums<<<numFullBlocks, fullblockSize, (fullblockSize * sizeof(float) * fragSize)>>>(d_A, d_A, dev_pitch_A, dev_pitch_A, N, 1, fragSize, lastBlockStartRow);
    computeSums<<<N, lastBlockSize, (lastBlockSize * sizeof(float) * fragSize)>>>(d_A, d_A, dev_pitch_A, dev_pitch_A, N, 0, fragSize, lastBlockStartRow);


    int numFinalSumBlocksReqd = ceil_h_d((float) N / (float) fullblockSize);
    int lastFinalSumBlockSize = N - (numFinalSumBlocksReqd - 1) * fullblockSize;
    int lastBlockStartCol = (numFinalSumBlocksReqd - 1) * fullblockSize;

    /* Parallel computation of averages of each column which is stored in array d_Avgs. */
    computeFinalSums<<<(numFinalSumBlocksReqd - 1), fullblockSize>>>(d_A, dev_pitch_A, N, 0, (fullblockSize * fragSize), lastBlockStartCol, d_Avgs);
    computeFinalSums<<<1, lastFinalSumBlockSize>>>(d_A, dev_pitch_A, N, 1, (fullblockSize * fragSize), lastBlockStartCol, d_Avgs);

    int numVarBlocksPerCol = ceil_h_d((float) N / (float) fullblockSize);
    int lastVarBlockSize = N - (numVarBlocksPerCol - 1) * fullblockSize;
    int lastVarBlockStartRow = (numVarBlocksPerCol - 1) * fullblockSize;

    dim3 numFullVarBlocks(N, (numVarBlocksPerCol - 1));

    /* Refreshing GPU's copy of array A with original copy from host as it has changed. */
    if (cudaMemcpy2D(d_A, dev_pitch_A, A_flat, host_pitch, N * sizeof(float), N, cudaMemcpyHostToDevice)!= cudaSuccess)
        printf("ERROR");

    /* Parallel computation of variances of each column which is  stored in array d_Vars */
    computeVarianceSquares<<<numFullVarBlocks, fullblockSize>>>(d_A, dev_pitch_A, d_Avgs, N, 0, lastVarBlockStartRow);
    computeVarianceSquares<<<N, lastVarBlockSize>>>(d_A, dev_pitch_A, d_Avgs, N, 1, lastVarBlockStartRow);

    /* Parallel computation of sums of variances in each column. */
    computeSums<<<numFullBlocks, fullblockSize, (fullblockSize * sizeof(float) * fragSize)>>>(d_A, d_A, dev_pitch_A, dev_pitch_A, N, 1, fragSize, lastBlockStartRow);
    computeSums<<<N, lastBlockSize, (lastBlockSize * sizeof(float) * fragSize)>>>(d_A, d_A, dev_pitch_A, dev_pitch_A, N, 0, fragSize, lastBlockStartRow);

    /* Parallel computation of averages of variances in each column */
    computeFinalSums<<<(numFinalSumBlocksReqd - 1), fullblockSize>>>(d_A, dev_pitch_A, N, 0, (fullblockSize * fragSize), lastBlockStartCol, d_Vars);
    computeFinalSums<<<1, lastFinalSumBlockSize>>>(d_A, dev_pitch_A, N, 1, (fullblockSize * fragSize), lastBlockStartCol, d_Vars);

    /* Refreshing GPU's copy of array A with original copy from host as it has changed. */
    if (cudaMemcpy2D(d_A, dev_pitch_A, A_flat, host_pitch, N * sizeof(float), N, cudaMemcpyHostToDevice)!= cudaSuccess)
        printf("ERROR");

    /* Calculating Final Norm values for each column. */
    computeNorms<<<numFullVarBlocks, fullblockSize>>>(d_A, dev_pitch_A, d_Avgs, d_Vars, N, 0, lastVarBlockStartRow);
    computeNorms<<<N, lastVarBlockSize>>>(d_A, dev_pitch_A, d_Avgs, d_Vars, N, 1, lastVarBlockStartRow);

    /* Copying result back to host. */
    cudaMemcpy2D(A_flat, host_pitch, d_A, dev_pitch_A, N * sizeof(float), N, cudaMemcpyDeviceToHost);

    /* Freeing Memory on GPU */
    cudaFree(d_A);
    cudaFree(d_Avgs);
    cudaFree(d_Vars);

    /* Unflattening returned array from 1D to 2D. */
    for (i = 0; i < N; i++) ///Unflattening array returned from GPU
    {
        for (j = 0; j < N; j++)
        {
            B[i][j] = A_flat[j + i * N];
        }
    }
}

	.file	"tmpxft_0025bd70_00000000-6_matrixNorm_GPU.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2050:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2050:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z9time_seedv
	.type	_Z9time_seedv, @function
_Z9time_seedv:
.LFB2040:
	.cfi_startproc
	endbr64
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movq	%fs:40, %rax
	movq	%rax, 24(%rsp)
	xorl	%eax, %eax
	movq	%rsp, %rsi
	leaq	8(%rsp), %rdi
	call	gettimeofday@PLT
	movl	16(%rsp), %eax
	movq	24(%rsp), %rdx
	subq	%fs:40, %rdx
	je	.L3
	call	__stack_chk_fail@PLT
.L3:
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2040:
	.size	_Z9time_seedv, .-_Z9time_seedv
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"Random seed = %i\n"
.LC1:
	.string	"N = %i is out of range.\n"
.LC2:
	.string	"Usage: %s <matrix_dimension> [random seed]\n"
.LC3:
	.string	"\nMatrix dimension N = %i.\n"
	.text
	.globl	_Z10parametersiPPc
	.type	_Z10parametersiPPc, @function
_Z10parametersiPPc:
.LFB2041:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movl	%edi, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movq	%rsi, %rbx
	pushq	%r8
	.cfi_def_cfa_offset 32
	call	_Z9time_seedv
	movl	%eax, %edi
	call	srand@PLT
	cmpl	$3, %ebp
	jne	.L7
	movq	16(%rbx), %rdi
	call	atoi@PLT
	movl	%eax, %ebp
	movl	%eax, %edi
	call	srand@PLT
	movl	%ebp, %edx
	leaq	.LC0(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	jmp	.L8
.L7:
	decl	%ebp
	jle	.L9
.L8:
	movq	8(%rbx), %rdi
	call	atoi@PLT
	movl	%eax, N(%rip)
	movl	%eax, %edx
	leal	-1(%rax), %eax
	cmpl	$7999, %eax
	jbe	.L10
	leaq	.LC1(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	jmp	.L12
.L9:
	movq	(%rbx), %rdx
	leaq	.LC2(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L12:
	xorl	%edi, %edi
	call	exit@PLT
.L10:
	popq	%rcx
	.cfi_def_cfa_offset 24
	leaq	.LC3(%rip), %rsi
	popq	%rbx
	.cfi_def_cfa_offset 16
	movl	$2, %edi
	xorl	%eax, %eax
	popq	%rbp
	.cfi_def_cfa_offset 8
	jmp	__printf_chk@PLT
	.cfi_endproc
.LFE2041:
	.size	_Z10parametersiPPc, .-_Z10parametersiPPc
	.section	.rodata.str1.1
.LC4:
	.string	"\nInitializing...\n"
	.text
	.globl	_Z17initialize_inputsv
	.type	_Z17initialize_inputsv, @function
_Z17initialize_inputsv:
.LFB2042:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	leaq	.LC4(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	xorl	%ebx, %ebx
	call	__printf_chk@PLT
.L14:
	cmpl	%ebx, N(%rip)
	jle	.L19
	xorl	%ebp, %ebp
	movslq	%ebx, %r12
.L16:
	cmpl	%ebp, N(%rip)
	jle	.L20
	call	rand@PLT
	leaq	A(%rip), %rdx
	cvtsi2ssl	%eax, %xmm0
	mulss	.LC5(%rip), %xmm0
	movslq	%ebp, %rax
	incl	%ebp
	imulq	$8000, %rax, %rax
	addq	%r12, %rax
	movss	%xmm0, (%rdx,%rax,4)
	leaq	B(%rip), %rdx
	movl	$0x00000000, (%rdx,%rax,4)
	jmp	.L16
.L20:
	incl	%ebx
	jmp	.L14
.L19:
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2042:
	.size	_Z17initialize_inputsv, .-_Z17initialize_inputsv
	.section	.rodata.str1.1
.LC7:
	.string	", "
.LC8:
	.string	";\n\t"
.LC9:
	.string	"\nA =\n\t"
.LC10:
	.string	"%5.2f%s"
	.text
	.globl	_Z12print_inputsv
	.type	_Z12print_inputsv, @function
_Z12print_inputsv:
.LFB2043:
	.cfi_startproc
	endbr64
	cmpl	$9, N(%rip)
	jg	.L30
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	leaq	.LC9(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	leaq	.LC8(%rip), %r12
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	xorl	%ebp, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	pushq	%rdx
	.cfi_def_cfa_offset 48
	call	__printf_chk@PLT
.L23:
	cmpl	N(%rip), %ebp
	jge	.L33
	xorl	%ebx, %ebx
	movslq	%ebp, %r13
.L26:
	movl	N(%rip), %eax
	cmpl	%ebx, %eax
	jle	.L34
	decl	%eax
	movslq	%ebx, %rcx
	leaq	.LC7(%rip), %rdx
	movl	$2, %edi
	cmpl	%ebx, %eax
	leaq	.LC10(%rip), %rsi
	cmovle	%r12, %rdx
	imulq	$8000, %r13, %rax
	incl	%ebx
	addq	%rcx, %rax
	leaq	A(%rip), %rcx
	movss	(%rcx,%rax,4), %xmm0
	movb	$1, %al
	cvtss2sd	%xmm0, %xmm0
	call	__printf_chk@PLT
	jmp	.L26
.L34:
	incl	%ebp
	jmp	.L23
.L33:
	popq	%rax
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	ret
.L30:
	.cfi_restore 3
	.cfi_restore 6
	.cfi_restore 12
	.cfi_restore 13
	ret
	.cfi_endproc
.LFE2043:
	.size	_Z12print_inputsv, .-_Z12print_inputsv
	.section	.rodata.str1.1
.LC11:
	.string	"\nB =\n\t"
.LC12:
	.string	"%1.10f%s"
	.text
	.globl	_Z7print_Bv
	.type	_Z7print_Bv, @function
_Z7print_Bv:
.LFB2044:
	.cfi_startproc
	endbr64
	cmpl	$9, N(%rip)
	jg	.L44
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	leaq	.LC11(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	leaq	.LC8(%rip), %r12
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	xorl	%ebp, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	pushq	%rdx
	.cfi_def_cfa_offset 48
	call	__printf_chk@PLT
.L37:
	cmpl	N(%rip), %ebp
	jge	.L47
	xorl	%ebx, %ebx
	movslq	%ebp, %r13
.L40:
	movl	N(%rip), %eax
	cmpl	%ebx, %eax
	jle	.L48
	decl	%eax
	movslq	%ebx, %rcx
	leaq	.LC7(%rip), %rdx
	movl	$2, %edi
	cmpl	%ebx, %eax
	leaq	.LC12(%rip), %rsi
	cmovle	%r12, %rdx
	imulq	$8000, %r13, %rax
	incl	%ebx
	addq	%rcx, %rax
	leaq	B(%rip), %rcx
	movss	(%rcx,%rax,4), %xmm0
	movb	$1, %al
	cvtss2sd	%xmm0, %xmm0
	call	__printf_chk@PLT
	jmp	.L40
.L48:
	incl	%ebp
	jmp	.L37
.L47:
	popq	%rax
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	ret
.L44:
	.cfi_restore 3
	.cfi_restore 6
	.cfi_restore 12
	.cfi_restore 13
	ret
	.cfi_endproc
.LFE2044:
	.size	_Z7print_Bv, .-_Z7print_Bv
	.globl	_Z8ceil_h_df
	.type	_Z8ceil_h_df, @function
_Z8ceil_h_df:
.LFB2046:
	.cfi_startproc
	endbr64
	cvttss2sil	%xmm0, %eax
	cvtsi2ssl	%eax, %xmm1
	comiss	%xmm1, %xmm0
	jbe	.L49
	incl	%eax
.L49:
	ret
	.cfi_endproc
.LFE2046:
	.size	_Z8ceil_h_df, .-_Z8ceil_h_df
	.globl	_Z33__device_stub__Z14blank_functionvv
	.type	_Z33__device_stub__Z14blank_functionvv, @function
_Z33__device_stub__Z14blank_functionvv:
.LFB2072:
	.cfi_startproc
	endbr64
	subq	$88, %rsp
	.cfi_def_cfa_offset 96
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	movabsq	$4294967297, %rax
	leaq	8(%rsp), %rcx
	movq	%rsp, %rdx
	leaq	28(%rsp), %rsi
	leaq	16(%rsp), %rdi
	movl	$1, 24(%rsp)
	movl	$1, 36(%rsp)
	movq	%rax, 16(%rsp)
	movq	%rax, 28(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L53
	pushq	8(%rsp)
	.cfi_def_cfa_offset 104
	leaq	_Z14blank_functionv(%rip), %rdi
	pushq	8(%rsp)
	.cfi_def_cfa_offset 112
	movq	44(%rsp), %rcx
	movl	52(%rsp), %r8d
	movq	32(%rsp), %rsi
	movl	40(%rsp), %edx
	leaq	80(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 104
	popq	%rdx
	.cfi_def_cfa_offset 96
.L53:
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L55
	call	__stack_chk_fail@PLT
.L55:
	addq	$88, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2072:
	.size	_Z33__device_stub__Z14blank_functionvv, .-_Z33__device_stub__Z14blank_functionvv
	.globl	_Z14blank_functionv
	.type	_Z14blank_functionv, @function
_Z14blank_functionv:
.LFB2073:
	.cfi_startproc
	endbr64
	jmp	_Z33__device_stub__Z14blank_functionvv
	.cfi_endproc
.LFE2073:
	.size	_Z14blank_functionv, .-_Z14blank_functionv
	.globl	_Z39__device_stub__Z11computeSumsPfS_mmiiiiPfS_mmiiii
	.type	_Z39__device_stub__Z11computeSumsPfS_mmiiiiPfS_mmiiii, @function
_Z39__device_stub__Z11computeSumsPfS_mmiiiiPfS_mmiiii:
.LFB2074:
	.cfi_startproc
	endbr64
	subq	$200, %rsp
	.cfi_def_cfa_offset 208
	movq	%rdi, 40(%rsp)
	leaq	72(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	84(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	56(%rsp), %rdx
	movq	%rcx, 16(%rsp)
	leaq	64(%rsp), %rcx
	movl	%r8d, 12(%rsp)
	movl	%r9d, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 184(%rsp)
	xorl	%eax, %eax
	leaq	40(%rsp), %rax
	movl	$1, 80(%rsp)
	movq	%rax, 120(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	208(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	216(%rsp), %rax
	movq	%rax, 176(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 72(%rsp)
	movq	%rax, 84(%rsp)
	movl	$1, 92(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L58
	pushq	64(%rsp)
	.cfi_def_cfa_offset 216
	leaq	_Z11computeSumsPfS_mmiiii(%rip), %rdi
	pushq	64(%rsp)
	.cfi_def_cfa_offset 224
	movq	100(%rsp), %rcx
	movl	108(%rsp), %r8d
	movq	88(%rsp), %rsi
	movl	96(%rsp), %edx
	leaq	136(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 216
	popq	%rdx
	.cfi_def_cfa_offset 208
.L58:
	movq	184(%rsp), %rax
	subq	%fs:40, %rax
	je	.L60
	call	__stack_chk_fail@PLT
.L60:
	addq	$200, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2074:
	.size	_Z39__device_stub__Z11computeSumsPfS_mmiiiiPfS_mmiiii, .-_Z39__device_stub__Z11computeSumsPfS_mmiiiiPfS_mmiiii
	.globl	_Z11computeSumsPfS_mmiiii
	.type	_Z11computeSumsPfS_mmiiii, @function
_Z11computeSumsPfS_mmiiii:
.LFB2075:
	.cfi_startproc
	endbr64
	jmp	_Z39__device_stub__Z11computeSumsPfS_mmiiiiPfS_mmiiii
	.cfi_endproc
.LFE2075:
	.size	_Z11computeSumsPfS_mmiiii, .-_Z11computeSumsPfS_mmiiii
	.globl	_Z43__device_stub__Z16computeFinalSumsPfmiiiiS_PfmiiiiS_
	.type	_Z43__device_stub__Z16computeFinalSumsPfmiiiiS_PfmiiiiS_, @function
_Z43__device_stub__Z16computeFinalSumsPfmiiiiS_PfmiiiiS_:
.LFB2076:
	.cfi_startproc
	endbr64
	subq	$184, %rsp
	.cfi_def_cfa_offset 192
	movq	192(%rsp), %rax
	movq	%rdi, 40(%rsp)
	leaq	64(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	76(%rsp), %rsi
	movl	%edx, 28(%rsp)
	leaq	48(%rsp), %rdx
	movl	%ecx, 24(%rsp)
	leaq	56(%rsp), %rcx
	movl	%r8d, 20(%rsp)
	movl	%r9d, 16(%rsp)
	movq	%rax, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 168(%rsp)
	xorl	%eax, %eax
	leaq	40(%rsp), %rax
	movl	$1, 72(%rsp)
	movq	%rax, 112(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	28(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 160(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 64(%rsp)
	movq	%rax, 76(%rsp)
	movl	$1, 84(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L63
	pushq	56(%rsp)
	.cfi_def_cfa_offset 200
	leaq	_Z16computeFinalSumsPfmiiiiS_(%rip), %rdi
	pushq	56(%rsp)
	.cfi_def_cfa_offset 208
	movq	92(%rsp), %rcx
	movl	100(%rsp), %r8d
	movq	80(%rsp), %rsi
	movl	88(%rsp), %edx
	leaq	128(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 200
	popq	%rdx
	.cfi_def_cfa_offset 192
.L63:
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	je	.L65
	call	__stack_chk_fail@PLT
.L65:
	addq	$184, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2076:
	.size	_Z43__device_stub__Z16computeFinalSumsPfmiiiiS_PfmiiiiS_, .-_Z43__device_stub__Z16computeFinalSumsPfmiiiiS_PfmiiiiS_
	.globl	_Z16computeFinalSumsPfmiiiiS_
	.type	_Z16computeFinalSumsPfmiiiiS_, @function
_Z16computeFinalSumsPfmiiiiS_:
.LFB2077:
	.cfi_startproc
	endbr64
	jmp	_Z43__device_stub__Z16computeFinalSumsPfmiiiiS_PfmiiiiS_
	.cfi_endproc
.LFE2077:
	.size	_Z16computeFinalSumsPfmiiiiS_, .-_Z16computeFinalSumsPfmiiiiS_
	.globl	_Z48__device_stub__Z22computeVarianceSquaresPfmS_iiiPfmS_iii
	.type	_Z48__device_stub__Z22computeVarianceSquaresPfmS_iiiPfmS_iii, @function
_Z48__device_stub__Z22computeVarianceSquaresPfmS_iiiPfmS_iii:
.LFB2078:
	.cfi_startproc
	endbr64
	subq	$184, %rsp
	.cfi_def_cfa_offset 192
	movq	%rdi, 40(%rsp)
	leaq	72(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	84(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	56(%rsp), %rdx
	movl	%ecx, 20(%rsp)
	leaq	64(%rsp), %rcx
	movl	%r8d, 16(%rsp)
	movl	%r9d, 12(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 168(%rsp)
	xorl	%eax, %eax
	leaq	40(%rsp), %rax
	movl	$1, 80(%rsp)
	movq	%rax, 120(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 160(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 72(%rsp)
	movq	%rax, 84(%rsp)
	movl	$1, 92(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L68
	pushq	64(%rsp)
	.cfi_def_cfa_offset 200
	leaq	_Z22computeVarianceSquaresPfmS_iii(%rip), %rdi
	pushq	64(%rsp)
	.cfi_def_cfa_offset 208
	movq	100(%rsp), %rcx
	movl	108(%rsp), %r8d
	movq	88(%rsp), %rsi
	movl	96(%rsp), %edx
	leaq	136(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 200
	popq	%rdx
	.cfi_def_cfa_offset 192
.L68:
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	je	.L70
	call	__stack_chk_fail@PLT
.L70:
	addq	$184, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2078:
	.size	_Z48__device_stub__Z22computeVarianceSquaresPfmS_iiiPfmS_iii, .-_Z48__device_stub__Z22computeVarianceSquaresPfmS_iiiPfmS_iii
	.globl	_Z22computeVarianceSquaresPfmS_iii
	.type	_Z22computeVarianceSquaresPfmS_iii, @function
_Z22computeVarianceSquaresPfmS_iii:
.LFB2079:
	.cfi_startproc
	endbr64
	jmp	_Z48__device_stub__Z22computeVarianceSquaresPfmS_iiiPfmS_iii
	.cfi_endproc
.LFE2079:
	.size	_Z22computeVarianceSquaresPfmS_iii, .-_Z22computeVarianceSquaresPfmS_iii
	.globl	_Z40__device_stub__Z12computeNormsPfmS_S_iiiPfmS_S_iii
	.type	_Z40__device_stub__Z12computeNormsPfmS_S_iiiPfmS_S_iii, @function
_Z40__device_stub__Z12computeNormsPfmS_S_iiiPfmS_S_iii:
.LFB2080:
	.cfi_startproc
	endbr64
	subq	$184, %rsp
	.cfi_def_cfa_offset 192
	movq	%rdi, 40(%rsp)
	leaq	64(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	76(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	48(%rsp), %rdx
	movq	%rcx, 16(%rsp)
	leaq	56(%rsp), %rcx
	movl	%r8d, 12(%rsp)
	movl	%r9d, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 168(%rsp)
	xorl	%eax, %eax
	leaq	40(%rsp), %rax
	movl	$1, 72(%rsp)
	movq	%rax, 112(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	192(%rsp), %rax
	movq	%rax, 160(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 64(%rsp)
	movq	%rax, 76(%rsp)
	movl	$1, 84(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L73
	pushq	56(%rsp)
	.cfi_def_cfa_offset 200
	leaq	_Z12computeNormsPfmS_S_iii(%rip), %rdi
	pushq	56(%rsp)
	.cfi_def_cfa_offset 208
	movq	92(%rsp), %rcx
	movl	100(%rsp), %r8d
	movq	80(%rsp), %rsi
	movl	88(%rsp), %edx
	leaq	128(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 200
	popq	%rdx
	.cfi_def_cfa_offset 192
.L73:
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	je	.L75
	call	__stack_chk_fail@PLT
.L75:
	addq	$184, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2080:
	.size	_Z40__device_stub__Z12computeNormsPfmS_S_iiiPfmS_S_iii, .-_Z40__device_stub__Z12computeNormsPfmS_S_iiiPfmS_S_iii
	.globl	_Z12computeNormsPfmS_S_iii
	.type	_Z12computeNormsPfmS_S_iii, @function
_Z12computeNormsPfmS_S_iii:
.LFB2081:
	.cfi_startproc
	endbr64
	jmp	_Z40__device_stub__Z12computeNormsPfmS_S_iiiPfmS_S_iii
	.cfi_endproc
.LFE2081:
	.size	_Z12computeNormsPfmS_S_iii, .-_Z12computeNormsPfmS_S_iii
	.section	.rodata.str1.1
.LC14:
	.string	"ERROR"
	.text
	.globl	_Z14matrixNorm_GPUv
	.type	_Z14matrixNorm_GPUv, @function
_Z14matrixNorm_GPUv:
.LFB2047:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	xorl	%esi, %esi
	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
	subq	$1176, %rsp
	.cfi_def_cfa_offset 1232
	movq	%fs:40, %rax
	movq	%rax, 1160(%rsp)
	xorl	%eax, %eax
	leaq	128(%rsp), %rdi
	call	cudaGetDeviceProperties_v2@PLT
	movl	$8, %ecx
	leaq	72(%rsp), %rsi
	movl	752(%rsp), %eax
	leaq	48(%rsp), %rdi
	cltd
	idivl	%ecx
	movslq	N(%rip), %rdx
	cvtsi2ssl	%edx, %xmm0
	divss	.LC13(%rip), %xmm0
	movl	%eax, %ebx
	call	_Z8ceil_h_df
	cvtsi2ssl	%ebx, %xmm1
	cvtsi2ssl	%eax, %xmm0
	movl	%eax, %ecx
	movss	%xmm1, 20(%rsp)
	divss	%xmm1, %xmm0
	call	_Z8ceil_h_df
	leal	-1(%rax), %r12d
	leaq	0(,%rdx,4), %rax
	movl	%r12d, %ebp
	movq	%rax, %rdx
	movq	%rax, (%rsp)
	imull	%ebx, %ebp
	subl	%ebp, %ecx
	imull	$10, %ebp, %ebp
	movl	%ecx, 40(%rsp)
	movq	%rax, %rcx
	call	cudaMallocPitch@PLT
	movslq	N(%rip), %rsi
	leaq	56(%rsp), %rdi
	salq	$2, %rsi
	call	cudaMalloc@PLT
	movslq	N(%rip), %rsi
	leaq	64(%rsp), %rdi
	salq	$2, %rsi
	call	cudaMalloc@PLT
	movl	N(%rip), %r13d
	movl	%r13d, %edi
	imull	%r13d, %edi
	movslq	%edi, %rdi
	salq	$2, %rdi
	call	malloc@PLT
	movslq	%r13d, %r9
	xorl	%edx, %edx
	leaq	A(%rip), %r10
	movq	%rax, 8(%rsp)
	leaq	0(,%r9,4), %r8
	movq	%rax, %r14
	movq	%rax, %rcx
.L79:
	cmpl	%edx, %r13d
	jle	.L113
	movslq	%edx, %rdi
	xorl	%eax, %eax
	imulq	$8000, %rdi, %rdi
.L80:
	movslq	%eax, %rsi
	addq	%rdi, %rsi
	movss	(%r10,%rsi,4), %xmm0
	movss	%xmm0, (%rcx,%rax,4)
	incq	%rax
	cmpl	%eax, %r13d
	jg	.L80
	incl	%edx
	addq	%r8, %rcx
	jmp	.L79
.L113:
	pushq	%rdi
	.cfi_def_cfa_offset 1240
	pushq	$1
	.cfi_def_cfa_offset 1248
	movq	16(%rsp), %rcx
	movq	24(%rsp), %rdx
	movq	88(%rsp), %rsi
	movq	64(%rsp), %rdi
	call	cudaMemcpy2D@PLT
	popq	%r8
	.cfi_def_cfa_offset 1240
	popq	%r9
	.cfi_def_cfa_offset 1232
	testl	%eax, %eax
	je	.L81
	leaq	.LC14(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L81:
	movl	N(%rip), %eax
	movslq	%ebx, %r15
	movl	%ebx, 116(%rsp)
	xorl	%r9d, %r9d
	movl	%r12d, 84(%rsp)
	movl	$1, %esi
	movl	%eax, 80(%rsp)
	imulq	$40, %r15, %rax
	movq	80(%rsp), %rdi
	movl	$1, 88(%rsp)
	movq	%rax, 24(%rsp)
	movq	24(%rsp), %r8
	movabsq	$4294967297, %rax
	movq	%rax, 120(%rsp)
	movl	124(%rsp), %ecx
	movq	116(%rsp), %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L82
	movq	72(%rsp), %rdx
	movq	48(%rsp), %rdi
	pushq	%rbp
	.cfi_def_cfa_offset 1240
	movl	$1, %r9d
	pushq	$10
	.cfi_def_cfa_offset 1248
	movl	N(%rip), %r8d
	movq	%rdx, %rcx
	movq	%rdi, %rsi
	call	_Z39__device_stub__Z11computeSumsPfS_mmiiiiPfS_mmiiii
	popq	%rcx
	.cfi_def_cfa_offset 1240
	popq	%rsi
	.cfi_def_cfa_offset 1232
.L82:
	movslq	40(%rsp), %rax
	movl	N(%rip), %edx
	xorl	%r9d, %r9d
	imulq	$40, %rax, %rax
	movl	%edx, 104(%rsp)
	movq	%rax, 32(%rsp)
	movl	40(%rsp), %eax
	movq	32(%rsp), %r8
	movl	%eax, 116(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 120(%rsp)
	movl	124(%rsp), %ecx
	movq	%rax, 108(%rsp)
	movq	116(%rsp), %rdx
	movq	104(%rsp), %rdi
	movl	112(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L83
	movq	72(%rsp), %rdx
	movq	48(%rsp), %rdi
	pushq	%rbp
	.cfi_def_cfa_offset 1240
	xorl	%r9d, %r9d
	pushq	$10
	.cfi_def_cfa_offset 1248
	movl	N(%rip), %r8d
	movq	%rdx, %rcx
	movq	%rdi, %rsi
	call	_Z39__device_stub__Z11computeSumsPfS_mmiiiiPfS_mmiiii
	popq	%rax
	.cfi_def_cfa_offset 1240
	popq	%rdx
	.cfi_def_cfa_offset 1232
.L83:
	movl	N(%rip), %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	cvtsi2ssl	%edx, %xmm0
	divss	20(%rsp), %xmm0
	call	_Z8ceil_h_df
	movl	%ebx, 116(%rsp)
	leal	-1(%rax), %r13d
	movl	%r13d, %eax
	movl	%r13d, 104(%rsp)
	imull	%ebx, %eax
	subl	%eax, %edx
	movl	%eax, 16(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 120(%rsp)
	movl	124(%rsp), %ecx
	movq	%rax, 108(%rsp)
	movq	104(%rsp), %rdi
	movl	%edx, 44(%rsp)
	movl	112(%rsp), %esi
	movq	116(%rsp), %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L84
	imull	$10, %ebx, %r8d
	pushq	%r11
	.cfi_def_cfa_offset 1240
	movl	N(%rip), %edx
	xorl	%ecx, %ecx
	pushq	64(%rsp)
	.cfi_def_cfa_offset 1248
	movl	32(%rsp), %r9d
	movq	88(%rsp), %rsi
	movq	64(%rsp), %rdi
	call	_Z43__device_stub__Z16computeFinalSumsPfmiiiiS_PfmiiiiS_
	popq	%r12
	.cfi_def_cfa_offset 1240
	popq	%r15
	.cfi_def_cfa_offset 1232
.L84:
	movl	44(%rsp), %eax
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %esi
	movabsq	$4294967297, %rdi
	movl	$1, 112(%rsp)
	movq	%rdi, 120(%rsp)
	movl	124(%rsp), %ecx
	movl	%eax, 116(%rsp)
	movq	116(%rsp), %rdx
	movq	%rdi, 104(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L85
	pushq	%r8
	.cfi_def_cfa_offset 1240
	imull	$10, %ebx, %r8d
	movl	N(%rip), %edx
	movl	$1, %ecx
	pushq	64(%rsp)
	.cfi_def_cfa_offset 1248
	movl	32(%rsp), %r9d
	movq	88(%rsp), %rsi
	movq	64(%rsp), %rdi
	call	_Z43__device_stub__Z16computeFinalSumsPfmiiiiS_PfmiiiiS_
	popq	%r9
	.cfi_def_cfa_offset 1240
	popq	%r10
	.cfi_def_cfa_offset 1232
.L85:
	movslq	N(%rip), %r9
	cvtsi2ssl	%r9d, %xmm0
	divss	20(%rsp), %xmm0
	leaq	0(,%r9,4), %r8
	movl	%r9d, %r15d
	call	_Z8ceil_h_df
	movl	%r9d, 92(%rsp)
	decl	%eax
	movl	$1, 100(%rsp)
	movl	%eax, 96(%rsp)
	movl	%eax, %r12d
	pushq	%rcx
	.cfi_def_cfa_offset 1240
	imull	%ebx, %r12d
	pushq	$1
	.cfi_def_cfa_offset 1248
	movq	16(%rsp), %rcx
	movq	24(%rsp), %rdx
	movq	88(%rsp), %rsi
	movq	64(%rsp), %rdi
	subl	%r12d, %r15d
	call	cudaMemcpy2D@PLT
	popq	%rsi
	.cfi_def_cfa_offset 1240
	popq	%rdi
	.cfi_def_cfa_offset 1232
	testl	%eax, %eax
	je	.L86
	leaq	.LC14(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L86:
	movl	%ebx, 116(%rsp)
	movq	92(%rsp), %rdi
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movabsq	$4294967297, %rax
	movl	100(%rsp), %esi
	movq	%rax, 120(%rsp)
	movl	124(%rsp), %ecx
	movq	116(%rsp), %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L87
	movq	56(%rsp), %rdx
	movq	72(%rsp), %rsi
	movl	%r12d, %r9d
	xorl	%r8d, %r8d
	movl	N(%rip), %ecx
	movq	48(%rsp), %rdi
	call	_Z48__device_stub__Z22computeVarianceSquaresPfmS_iiiPfmS_iii
.L87:
	movl	N(%rip), %edx
	movl	%r15d, 116(%rsp)
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movabsq	$4294967297, %rax
	movq	%rax, 120(%rsp)
	movl	124(%rsp), %ecx
	movl	%edx, 104(%rsp)
	movq	116(%rsp), %rdx
	movq	%rax, 108(%rsp)
	movq	104(%rsp), %rdi
	movl	112(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L88
	movl	N(%rip), %ecx
	movq	56(%rsp), %rdx
	movl	%r12d, %r9d
	movl	$1, %r8d
	movq	72(%rsp), %rsi
	movq	48(%rsp), %rdi
	call	_Z48__device_stub__Z22computeVarianceSquaresPfmS_iiiPfmS_iii
.L88:
	movq	24(%rsp), %r8
	movl	88(%rsp), %esi
	movl	%ebx, 116(%rsp)
	xorl	%r9d, %r9d
	movabsq	$4294967297, %rax
	movq	80(%rsp), %rdi
	movq	%rax, 120(%rsp)
	movl	124(%rsp), %ecx
	movq	116(%rsp), %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L89
	movq	72(%rsp), %rdx
	movq	48(%rsp), %rdi
	pushq	%rbp
	.cfi_def_cfa_offset 1240
	movl	$1, %r9d
	pushq	$10
	.cfi_def_cfa_offset 1248
	movl	N(%rip), %r8d
	movq	%rdx, %rcx
	movq	%rdi, %rsi
	call	_Z39__device_stub__Z11computeSumsPfS_mmiiiiPfS_mmiiii
	popq	%rax
	.cfi_def_cfa_offset 1240
	popq	%rdx
	.cfi_def_cfa_offset 1232
.L89:
	movl	40(%rsp), %eax
	movl	N(%rip), %edx
	xorl	%r9d, %r9d
	movq	32(%rsp), %r8
	movl	%eax, 116(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 120(%rsp)
	movl	124(%rsp), %ecx
	movl	%edx, 104(%rsp)
	movq	116(%rsp), %rdx
	movq	%rax, 108(%rsp)
	movq	104(%rsp), %rdi
	movl	112(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L90
	movq	72(%rsp), %rdx
	movq	48(%rsp), %rdi
	pushq	%rbp
	.cfi_def_cfa_offset 1240
	xorl	%r9d, %r9d
	pushq	$10
	.cfi_def_cfa_offset 1248
	movl	N(%rip), %r8d
	movq	%rdx, %rcx
	movq	%rdi, %rsi
	call	_Z39__device_stub__Z11computeSumsPfS_mmiiiiPfS_mmiiii
	popq	%r11
	.cfi_def_cfa_offset 1240
	popq	%rbp
	.cfi_def_cfa_offset 1232
.L90:
	movabsq	$4294967297, %rax
	movl	%ebx, 116(%rsp)
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	%rax, 120(%rsp)
	movl	124(%rsp), %ecx
	movq	%rax, 108(%rsp)
	movq	116(%rsp), %rdx
	movl	%r13d, 104(%rsp)
	movl	112(%rsp), %esi
	movq	104(%rsp), %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L91
	pushq	%r8
	.cfi_def_cfa_offset 1240
	imull	$10, %ebx, %r8d
	movl	N(%rip), %edx
	xorl	%ecx, %ecx
	pushq	72(%rsp)
	.cfi_def_cfa_offset 1248
	movl	32(%rsp), %r9d
	movq	88(%rsp), %rsi
	movq	64(%rsp), %rdi
	call	_Z43__device_stub__Z16computeFinalSumsPfmiiiiS_PfmiiiiS_
	popq	%r9
	.cfi_def_cfa_offset 1240
	popq	%r10
	.cfi_def_cfa_offset 1232
.L91:
	movl	44(%rsp), %eax
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %esi
	movabsq	$4294967297, %rdi
	movl	$1, 112(%rsp)
	movq	%rdi, 120(%rsp)
	movl	124(%rsp), %ecx
	movl	%eax, 116(%rsp)
	movq	116(%rsp), %rdx
	movq	%rdi, 104(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L92
	imull	$10, %ebx, %r8d
	pushq	%rcx
	.cfi_def_cfa_offset 1240
	movl	N(%rip), %edx
	movl	$1, %ecx
	pushq	72(%rsp)
	.cfi_def_cfa_offset 1248
	movl	32(%rsp), %r9d
	movq	88(%rsp), %rsi
	movq	64(%rsp), %rdi
	call	_Z43__device_stub__Z16computeFinalSumsPfmiiiiS_PfmiiiiS_
	popq	%rsi
	.cfi_def_cfa_offset 1240
	popq	%rdi
	.cfi_def_cfa_offset 1232
.L92:
	movslq	N(%rip), %r9
	pushq	%rbp
	.cfi_def_cfa_offset 1240
	pushq	$1
	.cfi_def_cfa_offset 1248
	movq	16(%rsp), %rcx
	movq	24(%rsp), %rdx
	movq	88(%rsp), %rsi
	leaq	0(,%r9,4), %r8
	movq	64(%rsp), %rdi
	call	cudaMemcpy2D@PLT
	popq	%r13
	.cfi_def_cfa_offset 1240
	popq	%rdx
	.cfi_def_cfa_offset 1232
	testl	%eax, %eax
	je	.L93
	leaq	.LC14(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L93:
	movl	%ebx, 116(%rsp)
	movq	92(%rsp), %rdi
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movabsq	$4294967297, %rax
	movl	100(%rsp), %esi
	movq	%rax, 120(%rsp)
	movl	124(%rsp), %ecx
	movq	116(%rsp), %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L94
	pushq	%r9
	.cfi_def_cfa_offset 1240
	movl	N(%rip), %r8d
	xorl	%r9d, %r9d
	pushq	%r12
	.cfi_def_cfa_offset 1248
	movq	80(%rsp), %rcx
	movq	72(%rsp), %rdx
	movq	88(%rsp), %rsi
	movq	64(%rsp), %rdi
	call	_Z40__device_stub__Z12computeNormsPfmS_S_iiiPfmS_S_iii
	popq	%r10
	.cfi_def_cfa_offset 1240
	popq	%r11
	.cfi_def_cfa_offset 1232
.L94:
	movl	N(%rip), %edx
	movl	%r15d, 116(%rsp)
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movabsq	$4294967297, %rax
	movq	%rax, 120(%rsp)
	movl	124(%rsp), %ecx
	movl	%edx, 104(%rsp)
	movq	116(%rsp), %rdx
	movq	%rax, 108(%rsp)
	movq	104(%rsp), %rdi
	movl	112(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L95
	pushq	%rsi
	.cfi_def_cfa_offset 1240
	movl	N(%rip), %r8d
	movl	$1, %r9d
	pushq	%r12
	.cfi_def_cfa_offset 1248
	movq	80(%rsp), %rcx
	movq	72(%rsp), %rdx
	movq	88(%rsp), %rsi
	movq	64(%rsp), %rdi
	call	_Z40__device_stub__Z12computeNormsPfmS_S_iiiPfmS_S_iii
	popq	%rdi
	.cfi_def_cfa_offset 1240
	popq	%r8
	.cfi_def_cfa_offset 1232
.L95:
	movslq	N(%rip), %r9
	pushq	%rax
	.cfi_def_cfa_offset 1240
	pushq	$2
	.cfi_def_cfa_offset 1248
	movq	16(%rsp), %rsi
	movq	24(%rsp), %rdi
	movq	88(%rsp), %rcx
	leaq	0(,%r9,4), %r8
	movq	64(%rsp), %rdx
	call	cudaMemcpy2D@PLT
	popq	%rdx
	.cfi_def_cfa_offset 1240
	popq	%rcx
	.cfi_def_cfa_offset 1232
	movq	48(%rsp), %rdi
	call	cudaFree@PLT
	movq	56(%rsp), %rdi
	call	cudaFree@PLT
	movq	64(%rsp), %rdi
	call	cudaFree@PLT
	movslq	N(%rip), %rsi
	xorl	%edx, %edx
	leaq	B(%rip), %r9
	movq	%rsi, %rcx
	salq	$2, %rsi
.L96:
	cmpl	%edx, %ecx
	jle	.L114
	movslq	%edx, %r8
	xorl	%eax, %eax
	imulq	$8000, %r8, %r8
.L97:
	movss	(%r14,%rax,4), %xmm0
	movslq	%eax, %rdi
	incq	%rax
	addq	%r8, %rdi
	movss	%xmm0, (%r9,%rdi,4)
	cmpl	%eax, %ecx
	jg	.L97
	incl	%edx
	addq	%rsi, %r14
	jmp	.L96
.L114:
	movq	1160(%rsp), %rax
	subq	%fs:40, %rax
	je	.L98
	call	__stack_chk_fail@PLT
.L98:
	addq	$1176, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2047:
	.size	_Z14matrixNorm_GPUv, .-_Z14matrixNorm_GPUv
	.section	.rodata.str1.1
.LC15:
	.string	"\nStarting clock.\n"
.LC16:
	.string	"Stopped clock.\n"
.LC18:
	.string	"\nElapsed time = %g ms.\n"
.LC20:
	.string	"(CPU times are accurate to the nearest %g ms)\n"
.LC22:
	.string	"My total CPU time for parent = %g ms.\n"
.LC23:
	.string	"My system CPU time for parent = %g ms.\n"
.LC24:
	.string	"My total CPU time for child processes = %g ms.\n"
.LC25:
	.string	"--------------------------------------------\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2045:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	movq	%rsi, %rbp
	movl	$1, %esi
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	movl	%edi, %ebx
	movabsq	$4294967297, %rdi
	movq	%rdi, %rdx
	subq	$112, %rsp
	.cfi_def_cfa_offset 144
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	movq	%rdi, 72(%rsp)
	movl	$1, 80(%rsp)
	movq	%rdi, 40(%rsp)
	movl	$1, 48(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L116
	call	_Z33__device_stub__Z14blank_functionvv
.L116:
	movq	%rbp, %rsi
	movl	%ebx, %edi
	movq	%rsp, %rbx
	call	_Z10parametersiPPc
	call	_Z17initialize_inputsv
	call	_Z12print_inputsv
	leaq	.LC15(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	%rbx, %rsi
	leaq	8(%rsp), %rdi
	call	gettimeofday@PLT
	leaq	40(%rsp), %rdi
	call	times@PLT
	call	_Z14matrixNorm_GPUv
	movq	%rbx, %rsi
	leaq	24(%rsp), %rdi
	call	gettimeofday@PLT
	leaq	72(%rsp), %rdi
	call	times@PLT
	leaq	.LC16(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	32(%rsp), %rbx
	imulq	$1000000, 8(%rsp), %r12
	imulq	$1000000, 24(%rsp), %rbp
	subq	16(%rsp), %rbp
	call	_Z7print_Bv
	leaq	(%rbx,%rbp), %rax
	subq	%r12, %rax
	js	.L117
	cvtsi2ssq	%rax, %xmm0
	jmp	.L118
.L117:
	movq	%rax, %rdx
	andl	$1, %eax
	shrq	%rdx
	orq	%rax, %rdx
	cvtsi2ssq	%rdx, %xmm0
	addss	%xmm0, %xmm0
.L118:
	leaq	.LC18(%rip), %rsi
	movl	$2, %edi
	movb	$1, %al
	divss	.LC17(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	call	__printf_chk@PLT
	movl	$2, %edi
	movb	$1, %al
	movsd	.LC19(%rip), %xmm0
	leaq	.LC20(%rip), %rsi
	call	__printf_chk@PLT
	movq	80(%rsp), %rax
	addq	72(%rsp), %rax
	leaq	.LC22(%rip), %rsi
	subq	48(%rsp), %rax
	subq	40(%rsp), %rax
	movl	$2, %edi
	cvtsi2ssq	%rax, %xmm0
	divss	.LC21(%rip), %xmm0
	movb	$1, %al
	mulss	.LC17(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	call	__printf_chk@PLT
	movq	80(%rsp), %rax
	subq	48(%rsp), %rax
	cvtsi2ssq	%rax, %xmm0
	divss	.LC21(%rip), %xmm0
	leaq	.LC23(%rip), %rsi
	movl	$2, %edi
	mulss	.LC17(%rip), %xmm0
	movb	$1, %al
	cvtss2sd	%xmm0, %xmm0
	call	__printf_chk@PLT
	movq	96(%rsp), %rax
	addq	88(%rsp), %rax
	subq	64(%rsp), %rax
	subq	56(%rsp), %rax
	movl	$2, %edi
	cvtsi2ssq	%rax, %xmm0
	divss	.LC21(%rip), %xmm0
	leaq	.LC24(%rip), %rsi
	movb	$1, %al
	mulss	.LC17(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	call	__printf_chk@PLT
	movl	$2, %edi
	xorl	%eax, %eax
	leaq	.LC25(%rip), %rsi
	call	__printf_chk@PLT
	xorl	%edi, %edi
	call	exit@PLT
	.cfi_endproc
.LFE2045:
	.size	main, .-main
	.section	.rodata.str1.1
.LC26:
	.string	"_Z12computeNormsPfmS_S_iii"
.LC27:
	.string	"_Z22computeVarianceSquaresPfmS_iii"
.LC28:
	.string	"_Z16computeFinalSumsPfmiiiiS_"
.LC29:
	.string	"_Z11computeSumsPfS_mmiiii"
.LC30:
	.string	"_Z14blank_functionv"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2083:
	.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	.LC26(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z12computeNormsPfmS_S_iii(%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	.LC27(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z22computeVarianceSquaresPfmS_iii(%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
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC28(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z16computeFinalSumsPfmiiiiS_(%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
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC29(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z11computeSumsPfS_mmiiii(%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
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC30(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z14blank_functionv(%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
.LFE2083:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.globl	B
	.bss
	.align 32
	.type	B, @object
	.size	B, 256000000
B:
	.zero	256000000
	.globl	A
	.align 32
	.type	A, @object
	.size	A, 256000000
A:
	.zero	256000000
	.globl	N
	.align 4
	.type	N, @object
	.size	N, 4
N:
	.zero	4
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC5:
	.long	939524096
	.align 4
.LC13:
	.long	1092616192
	.align 4
.LC17:
	.long	1148846080
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC19:
	.long	-755914244
	.long	1062232653
	.section	.rodata.cst4
	.align 4
.LC21:
	.long	1232348160
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:

/* Matrix normalization.
 * Compile with "gcc matrixNorm.c"
 */

/* ****** ADD YOUR CODE AT THE END OF THIS FILE. ******
 * You need not submit the provided code.
 */


#include <hip/hip_runtime.h>
#include <stdio.h>
#include <stdlib.h>
#include <unistd.h>
#include <math.h>
#include <sys/types.h>
#include <sys/times.h>
#include <sys/time.h>
#include <time.h>

/* Program Parameters */
#define MAXN 8000  /* Max value of N */
int N;  /* Matrix size */

/* Matrices */
volatile float A[MAXN][MAXN], B[MAXN][MAXN];

/* junk */
#define randm() 4|2[uid]&3

/* Prototype */
void matrixNorm_GPU();

__global__ void blank_function();

/* returns a seed for srand based on the time */
unsigned int time_seed() {
  struct timeval t;
  struct timezone tzdummy;

  gettimeofday(&t, &tzdummy);
  return (unsigned int)(t.tv_usec);
}

/* Set the program parameters from the command-line arguments */
void parameters(int argc, char **argv) {
  int seed = 0;  /* Random seed */
  char uid[32]; /*User name */

  /* Read command-line arguments */
  srand(time_seed());  /* Randomize */

  if (argc == 3) {
    seed = atoi(argv[2]);
    srand(seed);
    printf("Random seed = %i\n", seed);
  }
  if (argc >= 2) {
    N = atoi(argv[1]);
    if (N < 1 || N > MAXN) {
      printf("N = %i is out of range.\n", N);
      exit(0);
    }
  }
  else {
    printf("Usage: %s <matrix_dimension> [random seed]\n",
           argv[0]);
    exit(0);
  }

  /* Print parameters */
  printf("\nMatrix dimension N = %i.\n", N);
}

/* Initialize A and B*/
void initialize_inputs() {
  int row, col;

  printf("\nInitializing...\n");
  for (col = 0; col < N; col++) {
    for (row = 0; row < N; row++) {
      A[row][col] = (float)rand() / 32768.0;
      B[row][col] = 0.0;
    }
  }

}

/* Print input matrices */
void print_inputs() {
  int row, col;

  if (N < 10) {
    printf("\nA =\n\t");
    for (row = 0; row < N; row++) {
      for (col = 0; col < N; col++) {
	    printf("%5.2f%s", A[row][col], (col < N-1) ? ", " : ";\n\t");
      }
    }
  }
}

void print_B() {
    int row, col;

    if (N < 10) {
        printf("\nB =\n\t");
        for (row = 0; row < N; row++) {
            for (col = 0; col < N; col++) {
                printf("%1.10f%s", B[row][col], (col < N-1) ? ", " : ";\n\t");
            }
        }
    }
}

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

  blank_function<<<1, 1>>>(); /* Blank Function (defined below) having no content which is simply used to initialize CUDA environment as it is implicitly initialized on first call and this delay is not caused by parallelized algorithm.
                                 For reference, follow this link - https://devtalk.nvidia.com/default/topic/392429/first-cudamalloc-takes-long-time-/ */
  /* Timing variables */
  struct timeval etstart, etstop;  /* Elapsed times using gettimeofday() */
  struct timezone tzdummy;
  clock_t etstart2, etstop2;  /* Elapsed times using times() */
  unsigned long long usecstart, usecstop;
  struct tms cputstart, cputstop;  /* CPU times for my processes */

  /* Process program parameters */
  parameters(argc, argv);

  /* Initialize A and B */
  initialize_inputs();

  /* Print input matrices */
  print_inputs();

  /* Start Clock */
  printf("\nStarting clock.\n");
  gettimeofday(&etstart, &tzdummy);
  etstart2 = times(&cputstart);

  /* Gaussian Elimination */
  matrixNorm_GPU();

  /* Stop Clock */
  gettimeofday(&etstop, &tzdummy);
  etstop2 = times(&cputstop);
  printf("Stopped clock.\n");
  usecstart = (unsigned long long)etstart.tv_sec * 1000000 + etstart.tv_usec;
  usecstop = (unsigned long long)etstop.tv_sec * 1000000 + etstop.tv_usec;

  /* Display output */
  print_B();

  /* Display timing results */
  printf("\nElapsed time = %g ms.\n",
	 (float)(usecstop - usecstart)/(float)1000);

  printf("(CPU times are accurate to the nearest %g ms)\n",
	 1.0/(float)CLOCKS_PER_SEC * 1000.0);
  printf("My total CPU time for parent = %g ms.\n",
	 (float)( (cputstop.tms_utime + cputstop.tms_stime) -
		  (cputstart.tms_utime + cputstart.tms_stime) ) /
	 (float)CLOCKS_PER_SEC * 1000);
  printf("My system CPU time for parent = %g ms.\n",
	 (float)(cputstop.tms_stime - cputstart.tms_stime) /
	 (float)CLOCKS_PER_SEC * 1000);
  printf("My total CPU time for child processes = %g ms.\n",
	 (float)( (cputstop.tms_cutime + cputstop.tms_cstime) -
		  (cputstart.tms_cutime + cputstart.tms_cstime) ) /
	 (float)CLOCKS_PER_SEC * 1000);
      /* Contrary to the man pages, this appears not to include the parent */
  printf("--------------------------------------------\n");

  exit(0);
}

/* ------------------ Above Was Provided --------------------- */

/****** You will replace this routine with your own parallel version *******/
/* Provided global variables are MAXN, N, A[][] and B[][],
 * defined in the beginning of this code.  B[][] is initialized to zeros.
 */

/* Blank function which is called at first as the first call to any CUDA kernel or function implicitly initializes the CUDA environment.
For reference, follow this link - https://devtalk.nvidia.com/default/topic/392429/first-cudamalloc-takes-long-time-/ */
__global__ void blank_function()
{

}

/* Computes the sums of each group of numbers whose size is denoted by fragSize. */
__global__ void computeSums(float *d_A, float *d_B, size_t pitch_A, size_t pitch_B, int n, int fullBlock, int fragSize, int lastBlockStartRow)
{
    /* Initializing shared memory which will hold the numbers each thread block will be summing. */
    extern __shared__ float chunk[];

    /* Variables to keep track of which number the given is operating on in a 2D perspective. */
    int tx, ty_base, i;
    tx = blockIdx.x;

    /* Calculating the y index of the first element of the numbers being handled by the thread block. */
    if (fullBlock == 0)
    {
        ty_base = threadIdx.x * fragSize + lastBlockStartRow;
    }
    else
    {
        ty_base = (threadIdx.x + blockIdx.y * blockDim.x) * fragSize;
    }

    float *bElem, *aElem;

    /* Getting the starting index of the number the given thread begins operating from on its fragment. This is relative to the chunk[] array. */
    int localStartIndex = threadIdx.x * fragSize;

    /* Copying elements which the block will be summing from global to shared memory of the thread block. */
    for (i = 0; i < fragSize; i++)
    {
        aElem = (float*)((char*)d_A + (pitch_A * (ty_base + i)));
        chunk[localStartIndex + i] = aElem[tx];
    }

    /* Summing up the elements in the block. */
    if (ty_base < (n - 1))
    {
        int inc = 1;
        while(inc < n)
        {
            if ((localStartIndex % (inc * fragSize)) == 0)
            {
                for (i = localStartIndex + inc; (i < localStartIndex + fragSize * inc) && (i < (blockDim.x * fragSize)) && (i < n); i += inc)
                {
                    chunk[localStartIndex] += chunk[i];
                }
                inc *= fragSize;
            }
            else
                break;
            __syncthreads();
        }
    }

    /* Copying result to array d_B on the GPU. */
    for (i = 0; i < fragSize; i++)
    {
        bElem = (float*)((char*)d_B + (pitch_B * (ty_base + i)));
        bElem[tx] = chunk[localStartIndex + i];
    }
}

/* Computes Average Values of the columns in the input 2D array flattened as d_B. */
__global__ void computeFinalSums(float *d_B, size_t pitch_B, int n, int ifLastBlock, int prevBlockSize, int lastBlockStartCol, float *d_Avgs)
{
    int tx, i;
    float *bElem, *bElemBase;

    /* Getting address of starting element of the column which the thread is handling. */
    bElemBase = (float *)((char *)d_B);

    /* Getting x position of the starting element of the column on which the thread begins work. */
    if (ifLastBlock == 0)
    {
        tx = blockIdx.x * blockDim.x + threadIdx.x;
    }
    else
    {
        tx = lastBlockStartCol + threadIdx.x;
    }

    /* Summing up to get final sums for each column and subsequently calculating average. */
    for (i = prevBlockSize; i < n; i += prevBlockSize)
    {
        bElem = (float *)((char *)d_B + (pitch_B * i));
        bElemBase[tx] += bElem[tx];
    }
    bElemBase[tx] /= n;
    d_Avgs[tx] = bElemBase[tx];
}

/* Computes the Variances of each column of input flattened 2D array d_A. */
__global__ void computeVarianceSquares(float *d_A, size_t pitch_A, float *d_Avgs, int n, int ifLastBlock, int lastBlockStartRow)
{
    int tx, ty;
    tx = blockIdx.x;

    /* Getting y position of the starting element of the column on which the thread begins work. */
    if (ifLastBlock == 1)
    {
        ty = threadIdx.x + lastBlockStartRow;
    }
    else
    {
        ty = threadIdx.x + blockIdx.y * blockDim.x;
    }

    /* Getting address of starting element of the column which the thread is handling. */
    float *aElem = (float *)((char *)d_A + (pitch_A * ty));

    /* Calculating and storing the variances in global memory. */
    aElem[tx] = pow((aElem[tx] - d_Avgs[tx]), 2);
}

/* Finally computes the matrix norms for each element in the input flattened 2D matrix d_A. */
__global__ void computeNorms(float *d_A, size_t pitch_A, float *d_Avgs, float *d_Vars, int n, int ifLastBlock, int lastBlockStartRow)
{
    int tx, ty;
    tx = blockIdx.x;

    /* Getting y position of the element on which the thread begins work. */
    if (ifLastBlock == 1)
    {
        ty = threadIdx.x + lastBlockStartRow;
    }
    else
    {
        ty = threadIdx.x + blockIdx.y * blockDim.x;
    }

    float *aElem = (float *)((char *)d_A + (pitch_A * ty));

    /* Calculating and finally storing the norm value of each element. */
    if (d_Vars[tx] != 0)
        aElem[tx] = (aElem[tx] - d_Avgs[tx]) / sqrt(d_Vars[tx]);
    else
        aElem[tx] = 0;
}

/* Self Implemented ceiling function. */
__host__ __device__ int ceil_h_d(float f)
{
    int tmp = (int) f;
	if (f > tmp)
		tmp++;
	return tmp;
}

/* Function to calculate the normalized values of a matrix, implemented on GPU. */
void matrixNorm_GPU()
{
    /* Getting Device Properties of GPU to scale performance according to GPU specs. */
    hipDeviceProp_t prop;
    hipGetDeviceProperties(&prop, 0);

    int numThreadsPerMP = prop.maxThreadsPerMultiProcessor; /* Getting max. number of concurrent threads each Streaming Multiprocessor in the GPU can run. This is a multiple of the warp size which is 32.*/
    int fragSize = 10; /* The number of elements which are summed by a thread. */
    int BLOCKS_PER_MP = 8; /* Number of threadBlocks each Streaming Multiprocessor supports. */

    int i, j;
    int fullblockSize = numThreadsPerMP / BLOCKS_PER_MP; /* Computing number of threads per thread Block shall run. It should be a multiple of the warp size for maximum efficiency and memory coalescing. */

    int numElemsCol = ceil_h_d((float) N / (float) fragSize); /* Getting number of summing blocks per column */

    int blocksReqdPerCol = ceil_h_d((float) numElemsCol / (float) fullblockSize); /* Getting number of thread blocks required to sum each column. */
    int lastBlockSize = numElemsCol - (blocksReqdPerCol - 1) * fullblockSize; /* Since thread block sizes are multiples of warp size, therefore the last blocksize if not a multiple of warp size, is handled by this thread block. */

    int lastBlockStartRow = (blocksReqdPerCol - 1) * fullblockSize * fragSize; /* Getting the starting row index of the chunk of array being handled by the final thread block */

    float *d_A, *d_Avgs, *d_Vars; /* Declaring device pointers */

    size_t dev_pitch_A;
    size_t host_pitch = N * sizeof(float);

    /* Allocating memory for 2D arrays on GPU */
    hipMallocPitch(&d_A, &dev_pitch_A, N * sizeof(float), N * sizeof(float));

    /* Allocating memory for 1D arrays on GPU */
    hipMalloc((void **)&d_Avgs, N * sizeof(float));
    hipMalloc((void **)&d_Vars, N * sizeof(float));

    /* Allocating memory for flattened 1D array on Host */
    float *A_flat = (float *)malloc(N * N * sizeof(float));

    /* Flattening 2D array on host to a 1D array for transfer to GPU */
    for (i = 0; i < N; i++)
    {
        for (j = 0; j < N; j++)
        {
            A_flat[j + i * N] = A[i][j];
        }
    }

    /* Copying flattened 2D array A from host to GPU */
    if (hipMemcpy2D(d_A, dev_pitch_A, A_flat, host_pitch, N * sizeof(float), N, hipMemcpyHostToDevice)!= hipSuccess)
        printf("ERROR");

    dim3 numFullBlocks(N, (blocksReqdPerCol - 1));

    /* Parallel computation of sums of each sum fragment in each column. */
    computeSums<<<numFullBlocks, fullblockSize, (fullblockSize * sizeof(float) * fragSize)>>>(d_A, d_A, dev_pitch_A, dev_pitch_A, N, 1, fragSize, lastBlockStartRow);
    computeSums<<<N, lastBlockSize, (lastBlockSize * sizeof(float) * fragSize)>>>(d_A, d_A, dev_pitch_A, dev_pitch_A, N, 0, fragSize, lastBlockStartRow);


    int numFinalSumBlocksReqd = ceil_h_d((float) N / (float) fullblockSize);
    int lastFinalSumBlockSize = N - (numFinalSumBlocksReqd - 1) * fullblockSize;
    int lastBlockStartCol = (numFinalSumBlocksReqd - 1) * fullblockSize;

    /* Parallel computation of averages of each column which is stored in array d_Avgs. */
    computeFinalSums<<<(numFinalSumBlocksReqd - 1), fullblockSize>>>(d_A, dev_pitch_A, N, 0, (fullblockSize * fragSize), lastBlockStartCol, d_Avgs);
    computeFinalSums<<<1, lastFinalSumBlockSize>>>(d_A, dev_pitch_A, N, 1, (fullblockSize * fragSize), lastBlockStartCol, d_Avgs);

    int numVarBlocksPerCol = ceil_h_d((float) N / (float) fullblockSize);
    int lastVarBlockSize = N - (numVarBlocksPerCol - 1) * fullblockSize;
    int lastVarBlockStartRow = (numVarBlocksPerCol - 1) * fullblockSize;

    dim3 numFullVarBlocks(N, (numVarBlocksPerCol - 1));

    /* Refreshing GPU's copy of array A with original copy from host as it has changed. */
    if (hipMemcpy2D(d_A, dev_pitch_A, A_flat, host_pitch, N * sizeof(float), N, hipMemcpyHostToDevice)!= hipSuccess)
        printf("ERROR");

    /* Parallel computation of variances of each column which is  stored in array d_Vars */
    computeVarianceSquares<<<numFullVarBlocks, fullblockSize>>>(d_A, dev_pitch_A, d_Avgs, N, 0, lastVarBlockStartRow);
    computeVarianceSquares<<<N, lastVarBlockSize>>>(d_A, dev_pitch_A, d_Avgs, N, 1, lastVarBlockStartRow);

    /* Parallel computation of sums of variances in each column. */
    computeSums<<<numFullBlocks, fullblockSize, (fullblockSize * sizeof(float) * fragSize)>>>(d_A, d_A, dev_pitch_A, dev_pitch_A, N, 1, fragSize, lastBlockStartRow);
    computeSums<<<N, lastBlockSize, (lastBlockSize * sizeof(float) * fragSize)>>>(d_A, d_A, dev_pitch_A, dev_pitch_A, N, 0, fragSize, lastBlockStartRow);

    /* Parallel computation of averages of variances in each column */
    computeFinalSums<<<(numFinalSumBlocksReqd - 1), fullblockSize>>>(d_A, dev_pitch_A, N, 0, (fullblockSize * fragSize), lastBlockStartCol, d_Vars);
    computeFinalSums<<<1, lastFinalSumBlockSize>>>(d_A, dev_pitch_A, N, 1, (fullblockSize * fragSize), lastBlockStartCol, d_Vars);

    /* Refreshing GPU's copy of array A with original copy from host as it has changed. */
    if (hipMemcpy2D(d_A, dev_pitch_A, A_flat, host_pitch, N * sizeof(float), N, hipMemcpyHostToDevice)!= hipSuccess)
        printf("ERROR");

    /* Calculating Final Norm values for each column. */
    computeNorms<<<numFullVarBlocks, fullblockSize>>>(d_A, dev_pitch_A, d_Avgs, d_Vars, N, 0, lastVarBlockStartRow);
    computeNorms<<<N, lastVarBlockSize>>>(d_A, dev_pitch_A, d_Avgs, d_Vars, N, 1, lastVarBlockStartRow);

    /* Copying result back to host. */
    hipMemcpy2D(A_flat, host_pitch, d_A, dev_pitch_A, N * sizeof(float), N, hipMemcpyDeviceToHost);

    /* Freeing Memory on GPU */
    hipFree(d_A);
    hipFree(d_Avgs);
    hipFree(d_Vars);

    /* Unflattening returned array from 1D to 2D. */
    for (i = 0; i < N; i++) ///Unflattening array returned from GPU
    {
        for (j = 0; j < N; j++)
        {
            B[i][j] = A_flat[j + i * N];
        }
    }
}

	.text
	.file	"matrixNorm_GPU.hip"
	.globl	_Z9time_seedv                   # -- Begin function _Z9time_seedv
	.type	_Z9time_seedv,@function
_Z9time_seedv:                          # @_Z9time_seedv
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -16
	leaq	16(%rsp), %rbx
	leaq	8(%rsp), %rsi
	movq	%rbx, %rdi
	callq	gettimeofday
	movl	8(%rbx), %eax
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z9time_seedv, .Lfunc_end0-_Z9time_seedv
	.cfi_endproc
                                        # -- End function
	.globl	_Z10parametersiPPc              # -- Begin function _Z10parametersiPPc
	.type	_Z10parametersiPPc,@function
_Z10parametersiPPc:                     # @_Z10parametersiPPc
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	subq	$32, %rsp
	.cfi_def_cfa_offset 64
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %rbp, -16
	movq	%rsi, %rbx
	movl	%edi, %ebp
	leaq	16(%rsp), %r14
	leaq	8(%rsp), %rsi
	movq	%r14, %rdi
	callq	gettimeofday
	movl	8(%r14), %edi
	callq	srand
	cmpl	$3, %ebp
	jne	.LBB1_2
# %bb.1:                                # %.thread
	movq	16(%rbx), %rdi
	callq	atoi
	movl	%eax, %ebp
	movl	%eax, %edi
	callq	srand
	movl	$.L.str, %edi
	movl	%ebp, %esi
	xorl	%eax, %eax
	callq	printf
	jmp	.LBB1_3
.LBB1_2:
	cmpl	$2, %ebp
	jl	.LBB1_7
.LBB1_3:
	movq	8(%rbx), %rdi
	callq	atoi
	movl	%eax, N(%rip)
	movl	$-8001, %ecx                    # imm = 0xE0BF
	movl	%eax, %edx
	addl	%ecx, %edx
	cmpl	%ecx, %edx
	jbe	.LBB1_5
# %bb.4:
	movl	$.L.str.3, %edi
	movl	%eax, %esi
	xorl	%eax, %eax
	addq	$32, %rsp
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	jmp	printf                          # TAILCALL
.LBB1_5:
	.cfi_def_cfa_offset 64
	movl	$.L.str.1, %edi
	movl	%eax, %esi
	xorl	%eax, %eax
	callq	printf
	jmp	.LBB1_6
.LBB1_7:
	movq	(%rbx), %rsi
	movl	$.L.str.2, %edi
	xorl	%eax, %eax
	callq	printf
.LBB1_6:
	xorl	%edi, %edi
	callq	exit
.Lfunc_end1:
	.size	_Z10parametersiPPc, .Lfunc_end1-_Z10parametersiPPc
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function _Z17initialize_inputsv
.LCPI2_0:
	.long	0x38000000                      # float 3.05175781E-5
	.text
	.globl	_Z17initialize_inputsv
	.type	_Z17initialize_inputsv,@function
_Z17initialize_inputsv:                 # @_Z17initialize_inputsv
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r12
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	pushq	%rax
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -40
	.cfi_offset %r12, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movl	$.Lstr, %edi
	callq	puts@PLT
	movl	N(%rip), %eax
	testl	%eax, %eax
	jle	.LBB2_6
# %bb.1:                                # %.preheader.preheader
	xorl	%ebx, %ebx
	xorl	%r14d, %r14d
.LBB2_2:                                # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB2_4 Depth 2
	testl	%eax, %eax
	jle	.LBB2_7
# %bb.3:                                # %.lr.ph.preheader
                                        #   in Loop: Header=BB2_2 Depth=1
	movq	%rbx, %r15
	xorl	%r12d, %r12d
.LBB2_4:                                # %.lr.ph
                                        #   Parent Loop BB2_2 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	callq	rand
	movss	.LCPI2_0(%rip), %xmm1           # xmm1 = mem[0],zero,zero,zero
	xorps	%xmm0, %xmm0
	cvtsi2ss	%eax, %xmm0
	mulss	%xmm1, %xmm0
	movss	%xmm0, A(%r15)
	movl	$0, B(%r15)
	incq	%r12
	movl	N(%rip), %eax
	movslq	%eax, %rcx
	addq	$32000, %r15                    # imm = 0x7D00
	cmpq	%rcx, %r12
	jl	.LBB2_4
	jmp	.LBB2_5
.LBB2_7:                                # %.preheader.._crit_edge_crit_edge
                                        #   in Loop: Header=BB2_2 Depth=1
	movslq	%eax, %rcx
.LBB2_5:                                # %._crit_edge
                                        #   in Loop: Header=BB2_2 Depth=1
	incq	%r14
	addq	$4, %rbx
	cmpq	%rcx, %r14
	jl	.LBB2_2
.LBB2_6:                                # %._crit_edge11
	addq	$8, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end2:
	.size	_Z17initialize_inputsv, .Lfunc_end2-_Z17initialize_inputsv
	.cfi_endproc
                                        # -- End function
	.globl	_Z12print_inputsv               # -- Begin function _Z12print_inputsv
	.type	_Z12print_inputsv,@function
_Z12print_inputsv:                      # @_Z12print_inputsv
	.cfi_startproc
# %bb.0:
	cmpl	$9, N(%rip)
	jg	.LBB3_8
# %bb.1:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r12
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	pushq	%rax
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -40
	.cfi_offset %r12, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movl	$.L.str.5, %edi
	xorl	%eax, %eax
	callq	printf
	movl	N(%rip), %eax
	testl	%eax, %eax
	jle	.LBB3_7
# %bb.2:                                # %.preheader.preheader
	movl	$A, %ebx
	xorl	%r14d, %r14d
	movl	$.L.str.7, %r15d
.LBB3_3:                                # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB3_5 Depth 2
	testl	%eax, %eax
	jle	.LBB3_9
# %bb.4:                                # %.lr.ph.preheader
                                        #   in Loop: Header=BB3_3 Depth=1
	xorl	%r12d, %r12d
.LBB3_5:                                # %.lr.ph
                                        #   Parent Loop BB3_3 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	xorps	%xmm0, %xmm0
	cvtss2sd	(%rbx,%r12,4), %xmm0
	decl	%eax
	cltq
	cmpq	%rax, %r12
	movl	$.L.str.8, %esi
	cmovlq	%r15, %rsi
	movl	$.L.str.6, %edi
	movb	$1, %al
	callq	printf
	incq	%r12
	movl	N(%rip), %eax
	movslq	%eax, %rcx
	cmpq	%rcx, %r12
	jl	.LBB3_5
	jmp	.LBB3_6
.LBB3_9:                                # %.preheader.._crit_edge_crit_edge
                                        #   in Loop: Header=BB3_3 Depth=1
	movslq	%eax, %rcx
.LBB3_6:                                # %._crit_edge
                                        #   in Loop: Header=BB3_3 Depth=1
	incq	%r14
	addq	$32000, %rbx                    # imm = 0x7D00
	cmpq	%rcx, %r14
	jl	.LBB3_3
.LBB3_7:
	addq	$8, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	.cfi_restore %rbx
	.cfi_restore %r12
	.cfi_restore %r14
	.cfi_restore %r15
.LBB3_8:                                # %.loopexit
	retq
.Lfunc_end3:
	.size	_Z12print_inputsv, .Lfunc_end3-_Z12print_inputsv
	.cfi_endproc
                                        # -- End function
	.globl	_Z7print_Bv                     # -- Begin function _Z7print_Bv
	.type	_Z7print_Bv,@function
_Z7print_Bv:                            # @_Z7print_Bv
	.cfi_startproc
# %bb.0:
	cmpl	$9, N(%rip)
	jg	.LBB4_8
# %bb.1:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r12
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	pushq	%rax
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -40
	.cfi_offset %r12, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movl	$.L.str.9, %edi
	xorl	%eax, %eax
	callq	printf
	movl	N(%rip), %eax
	testl	%eax, %eax
	jle	.LBB4_7
# %bb.2:                                # %.preheader.preheader
	movl	$B, %ebx
	xorl	%r14d, %r14d
	movl	$.L.str.7, %r15d
.LBB4_3:                                # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB4_5 Depth 2
	testl	%eax, %eax
	jle	.LBB4_9
# %bb.4:                                # %.lr.ph.preheader
                                        #   in Loop: Header=BB4_3 Depth=1
	xorl	%r12d, %r12d
.LBB4_5:                                # %.lr.ph
                                        #   Parent Loop BB4_3 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	xorps	%xmm0, %xmm0
	cvtss2sd	(%rbx,%r12,4), %xmm0
	decl	%eax
	cltq
	cmpq	%rax, %r12
	movl	$.L.str.8, %esi
	cmovlq	%r15, %rsi
	movl	$.L.str.10, %edi
	movb	$1, %al
	callq	printf
	incq	%r12
	movl	N(%rip), %eax
	movslq	%eax, %rcx
	cmpq	%rcx, %r12
	jl	.LBB4_5
	jmp	.LBB4_6
.LBB4_9:                                # %.preheader.._crit_edge_crit_edge
                                        #   in Loop: Header=BB4_3 Depth=1
	movslq	%eax, %rcx
.LBB4_6:                                # %._crit_edge
                                        #   in Loop: Header=BB4_3 Depth=1
	incq	%r14
	addq	$32000, %rbx                    # imm = 0x7D00
	cmpq	%rcx, %r14
	jl	.LBB4_3
.LBB4_7:
	addq	$8, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	.cfi_restore %rbx
	.cfi_restore %r12
	.cfi_restore %r14
	.cfi_restore %r15
.LBB4_8:                                # %.loopexit
	retq
.Lfunc_end4:
	.size	_Z7print_Bv, .Lfunc_end4-_Z7print_Bv
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function main
.LCPI5_0:
	.long	0x447a0000                      # float 1000
.LCPI5_2:
	.long	0x49742400                      # float 1.0E+6
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0
.LCPI5_1:
	.quad	0x3f50624dd2f1a9fc              # double 0.001
	.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	%rbx
	.cfi_def_cfa_offset 40
	subq	$104, %rsp
	.cfi_def_cfa_offset 144
	.cfi_offset %rbx, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%rsi, %rbx
	movl	%edi, %ebp
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB5_2
# %bb.1:
	callq	_Z29__device_stub__blank_functionv
.LBB5_2:
	movl	%ebp, %edi
	movq	%rbx, %rsi
	callq	_Z10parametersiPPc
	callq	_Z17initialize_inputsv
	callq	_Z12print_inputsv
	movl	$.Lstr.1, %edi
	callq	puts@PLT
	leaq	88(%rsp), %rbx
	leaq	64(%rsp), %r14
	movq	%rbx, %rdi
	movq	%r14, %rsi
	callq	gettimeofday
	leaq	32(%rsp), %rdi
	callq	times
	callq	_Z14matrixNorm_GPUv
	leaq	72(%rsp), %r15
	movq	%r15, %rdi
	movq	%r14, %rsi
	callq	gettimeofday
	movq	%rsp, %rdi
	callq	times
	movl	$.Lstr.2, %edi
	callq	puts@PLT
	movq	(%r15), %r14
	movq	8(%r15), %r15
	subq	(%rbx), %r14
	subq	8(%rbx), %r15
	callq	_Z7print_Bv
	imulq	$1000000, %r14, %rax            # imm = 0xF4240
	addq	%r15, %rax
	js	.LBB5_3
# %bb.4:
	cvtsi2ss	%rax, %xmm0
	jmp	.LBB5_5
.LBB5_3:
	movq	%rax, %rcx
	shrq	%rcx
	andl	$1, %eax
	orq	%rcx, %rax
	cvtsi2ss	%rax, %xmm0
	addss	%xmm0, %xmm0
.LBB5_5:
	divss	.LCPI5_0(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	movl	$.L.str.13, %edi
	movb	$1, %al
	callq	printf
	movsd	.LCPI5_1(%rip), %xmm0           # xmm0 = mem[0],zero
	movl	$.L.str.14, %edi
	movb	$1, %al
	callq	printf
	movq	8(%rsp), %rbx
	movq	40(%rsp), %r14
	movq	(%rsp), %rax
	addq	%rbx, %rax
	movq	32(%rsp), %rcx
	addq	%r14, %rcx
	subq	%rcx, %rax
	xorps	%xmm0, %xmm0
	cvtsi2ss	%rax, %xmm0
	divss	.LCPI5_2(%rip), %xmm0
	mulss	.LCPI5_0(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	movl	$.L.str.15, %edi
	movb	$1, %al
	callq	printf
	subq	%r14, %rbx
	xorps	%xmm0, %xmm0
	cvtsi2ss	%rbx, %xmm0
	divss	.LCPI5_2(%rip), %xmm0
	mulss	.LCPI5_0(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	movl	$.L.str.16, %edi
	movb	$1, %al
	callq	printf
	movq	24(%rsp), %rax
	movq	48(%rsp), %rcx
	addq	16(%rsp), %rax
	addq	56(%rsp), %rcx
	subq	%rcx, %rax
	xorps	%xmm0, %xmm0
	cvtsi2ss	%rax, %xmm0
	divss	.LCPI5_2(%rip), %xmm0
	mulss	.LCPI5_0(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	movl	$.L.str.17, %edi
	movb	$1, %al
	callq	printf
	movl	$.Lstr.3, %edi
	callq	puts@PLT
	xorl	%edi, %edi
	callq	exit
.Lfunc_end5:
	.size	main, .Lfunc_end5-main
	.cfi_endproc
                                        # -- End function
	.globl	_Z29__device_stub__blank_functionv # -- Begin function _Z29__device_stub__blank_functionv
	.type	_Z29__device_stub__blank_functionv,@function
_Z29__device_stub__blank_functionv:     # @_Z29__device_stub__blank_functionv
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r12
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	subq	$56, %rsp
	.cfi_def_cfa_offset 96
	.cfi_offset %rbx, -40
	.cfi_offset %r12, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rbx
	leaq	24(%rsp), %r14
	leaq	16(%rsp), %r15
	leaq	8(%rsp), %r12
	movq	%rbx, %rdi
	movq	%r14, %rsi
	movq	%r15, %rdx
	movq	%r12, %rcx
	callq	__hipPopCallConfiguration
	movq	(%rbx), %rsi
	movl	8(%rbx), %edx
	movq	(%r14), %rcx
	movl	8(%r14), %r8d
	movq	%rsp, %r9
	movl	$_Z14blank_functionv, %edi
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	pushq	(%r15)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$72, %rsp
	.cfi_adjust_cfa_offset -72
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end6:
	.size	_Z29__device_stub__blank_functionv, .Lfunc_end6-_Z29__device_stub__blank_functionv
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function _Z14matrixNorm_GPUv
.LCPI7_0:
	.long	0x41200000                      # float 10
	.text
	.globl	_Z14matrixNorm_GPUv
	.type	_Z14matrixNorm_GPUv,@function
_Z14matrixNorm_GPUv:                    # @_Z14matrixNorm_GPUv
	.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	$1608, %rsp                     # imm = 0x648
	.cfi_def_cfa_offset 1664
	.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	136(%rsp), %rbx
	xorl	%r15d, %r15d
	movq	%rbx, %rdi
	xorl	%esi, %esi
	callq	hipGetDevicePropertiesR0600
	movl	624(%rbx), %eax
	leal	7(%rax), %r12d
	testl	%eax, %eax
	cmovnsl	%eax, %r12d
	movslq	N(%rip), %rdx
	cvtsi2ss	%edx, %xmm0
	divss	.LCPI7_0(%rip), %xmm0
	cvttss2si	%xmm0, %eax
	sarl	$3, %r12d
	cvttps2dq	%xmm0, %xmm1
	cvtdq2ps	%xmm1, %xmm1
	xorl	%ebp, %ebp
	ucomiss	%xmm1, %xmm0
	seta	%bpl
	addl	%eax, %ebp
	xorps	%xmm0, %xmm0
	cvtsi2ss	%ebp, %xmm0
	xorps	%xmm1, %xmm1
	cvtsi2ss	%r12d, %xmm1
	movss	%xmm1, 36(%rsp)                 # 4-byte Spill
	divss	%xmm1, %xmm0
	cvttss2si	%xmm0, %eax
	cvttps2dq	%xmm0, %xmm1
	cvtdq2ps	%xmm1, %xmm1
	xorl	%ecx, %ecx
	ucomiss	%xmm1, %xmm0
	seta	%cl
	leal	(%rax,%rcx), %r13d
	decl	%r13d
	movq	%r13, 88(%rsp)                  # 8-byte Spill
                                        # kill: def $r13d killed $r13d killed $r13 def $r13
	imull	%r12d, %r13d
	leal	(,%r13,2), %eax
	movq	%rax, 56(%rsp)                  # 8-byte Spill
	shlq	$2, %rdx
	leaq	24(%rsp), %rdi
	leaq	40(%rsp), %rsi
	movq	%rdx, 80(%rsp)                  # 8-byte Spill
	movq	%rdx, %rcx
	callq	hipMallocPitch
	movslq	N(%rip), %rsi
	shlq	$2, %rsi
	leaq	48(%rsp), %rdi
	callq	hipMalloc
	movslq	N(%rip), %rsi
	shlq	$2, %rsi
	leaq	64(%rsp), %rdi
	callq	hipMalloc
	movslq	N(%rip), %rbx
	movl	%ebx, %edi
	imull	%edi, %edi
	shlq	$2, %rdi
	callq	malloc
	movq	%rax, %r14
	testq	%rbx, %rbx
	jle	.LBB7_6
# %bb.1:                                # %.preheader233.lr.ph
	movl	%ebx, %eax
	movl	$A, %ecx
	xorl	%edx, %edx
.LBB7_2:                                # %.lr.ph
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB7_3 Depth 2
	movl	%r15d, %esi
	leaq	(%r14,%rsi,4), %rsi
	xorl	%edi, %edi
.LBB7_3:                                #   Parent Loop BB7_2 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movss	(%rcx,%rdi,4), %xmm0            # xmm0 = mem[0],zero,zero,zero
	movss	%xmm0, (%rsi,%rdi,4)
	incq	%rdi
	cmpq	%rdi, %rax
	jne	.LBB7_3
# %bb.4:                                # %._crit_edge
                                        #   in Loop: Header=BB7_2 Depth=1
	incq	%rdx
	addl	%ebx, %r15d
	addq	$32000, %rcx                    # imm = 0x7D00
	cmpq	%rax, %rdx
	jne	.LBB7_2
# %bb.5:
	movq	%rax, %rbx
.LBB7_6:                                # %._crit_edge236
	movq	24(%rsp), %rdi
	movq	40(%rsp), %rsi
	leaq	(,%rbx,4), %r8
	movl	$1, (%rsp)
	movq	%r14, %rdx
	movq	80(%rsp), %rcx                  # 8-byte Reload
	movq	%rbx, %r9
	callq	hipMemcpy2D
	testl	%eax, %eax
	je	.LBB7_8
# %bb.7:
	movl	$.L.str.19, %edi
	xorl	%eax, %eax
	callq	printf
.LBB7_8:
	subl	%r13d, %ebp
	movq	56(%rsp), %rax                  # 8-byte Reload
	leal	(%rax,%rax,4), %r13d
	movabsq	$4294967296, %rbx               # imm = 0x100000000
	movl	N(%rip), %eax
	movq	88(%rsp), %rdi                  # 8-byte Reload
	shlq	$32, %rdi
	orq	%rax, %rdi
	movslq	%r12d, %rax
	shlq	$3, %rax
	leaq	(%rax,%rax,4), %r8
	movl	%r12d, %r15d
	btsq	$32, %r15
	movq	%rdi, 88(%rsp)                  # 8-byte Spill
	movl	$1, %esi
	movq	%r15, %rdx
	movl	$1, %ecx
	movq	%r8, 128(%rsp)                  # 8-byte Spill
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB7_10
# %bb.9:
	movq	24(%rsp), %rdi
	movq	40(%rsp), %rcx
	movl	N(%rip), %r8d
	movl	%r13d, 8(%rsp)
	movl	$10, (%rsp)
	movq	%rdi, %rsi
	movq	%rcx, %rdx
	movl	$1, %r9d
	callq	_Z26__device_stub__computeSumsPfS_mmiiii
.LBB7_10:
	movl	N(%rip), %edi
	movslq	%ebp, %rax
	leaq	(,%rax,8), %rcx
	leaq	(%rcx,%rcx,4), %r8
	orq	%rbx, %rdi
	movl	%eax, %edx
	orq	%rbx, %rdx
	movl	$1, %esi
	movq	%rdx, 112(%rsp)                 # 8-byte Spill
	movl	$1, %ecx
	movq	%r8, 120(%rsp)                  # 8-byte Spill
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	movq	%rbx, %rbp
	jne	.LBB7_12
# %bb.11:
	movq	24(%rsp), %rdi
	movq	40(%rsp), %rcx
	movl	N(%rip), %r8d
	movl	%r13d, 8(%rsp)
	movl	$10, (%rsp)
	movq	%rdi, %rsi
	movq	%rcx, %rdx
	xorl	%r9d, %r9d
	callq	_Z26__device_stub__computeSumsPfS_mmiiii
.LBB7_12:
	movl	%r13d, 76(%rsp)                 # 4-byte Spill
	movl	N(%rip), %ebx
	xorps	%xmm0, %xmm0
	cvtsi2ss	%ebx, %xmm0
	divss	36(%rsp), %xmm0                 # 4-byte Folded Reload
	cvttss2si	%xmm0, %eax
	cvttps2dq	%xmm0, %xmm1
	cvtdq2ps	%xmm1, %xmm1
	xorl	%ecx, %ecx
	ucomiss	%xmm1, %xmm0
	seta	%cl
	leal	(%rax,%rcx), %r13d
	decl	%r13d
	movl	%r13d, %eax
	imull	%r12d, %eax
	movl	%eax, 56(%rsp)                  # 4-byte Spill
	subl	%eax, %ebx
	orq	%rbp, %r13
	movq	%r13, %rdi
	movl	$1, %esi
	movq	%r15, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB7_14
# %bb.13:
	movq	24(%rsp), %rdi
	movq	40(%rsp), %rsi
	movl	N(%rip), %edx
	leal	(%r12,%r12), %eax
	leal	(%rax,%rax,4), %r8d
	movq	48(%rsp), %rax
	movq	%rax, (%rsp)
	xorl	%ecx, %ecx
	movl	56(%rsp), %r9d                  # 4-byte Reload
	callq	_Z31__device_stub__computeFinalSumsPfmiiiiS_
.LBB7_14:
	movl	%ebx, %edx
	orq	%rbp, %rdx
	leaq	1(%rbp), %rdi
	movq	%rdi, 96(%rsp)                  # 8-byte Spill
	movl	$1, %esi
	movq	%rdx, 104(%rsp)                 # 8-byte Spill
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB7_16
# %bb.15:
	movq	24(%rsp), %rdi
	movq	40(%rsp), %rsi
	movl	N(%rip), %edx
	leal	(%r12,%r12), %eax
	leal	(%rax,%rax,4), %r8d
	movq	48(%rsp), %rax
	movq	%rax, (%rsp)
	movl	$1, %ecx
	movl	56(%rsp), %r9d                  # 4-byte Reload
	callq	_Z31__device_stub__computeFinalSumsPfmiiiiS_
.LBB7_16:
	movslq	N(%rip), %rbx
	xorps	%xmm0, %xmm0
	cvtsi2ss	%ebx, %xmm0
	divss	36(%rsp), %xmm0                 # 4-byte Folded Reload
	cvttss2si	%xmm0, %eax
	cvttps2dq	%xmm0, %xmm1
	cvtdq2ps	%xmm1, %xmm1
	xorl	%ecx, %ecx
	ucomiss	%xmm1, %xmm0
	seta	%cl
	leal	(%rax,%rcx), %ebp
	decl	%ebp
	movl	%ebx, %eax
	movl	%ebp, %ecx
	imull	%r12d, %ecx
	movl	%ecx, 36(%rsp)                  # 4-byte Spill
	shlq	$32, %rbp
	orq	%rax, %rbp
	movq	24(%rsp), %rdi
	movq	40(%rsp), %rsi
	leaq	(,%rbx,4), %r8
	movl	$1, (%rsp)
	movq	%r14, %rdx
	movq	80(%rsp), %rcx                  # 8-byte Reload
	movq	%rbx, %r9
	callq	hipMemcpy2D
	testl	%eax, %eax
	je	.LBB7_18
# %bb.17:
	movl	$.L.str.19, %edi
	xorl	%eax, %eax
	callq	printf
.LBB7_18:
	subl	36(%rsp), %ebx                  # 4-byte Folded Reload
	movq	%rbp, %rdi
	movl	$1, %esi
	movq	%r15, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB7_20
# %bb.19:
	movq	24(%rsp), %rdi
	movq	40(%rsp), %rsi
	movq	48(%rsp), %rdx
	movl	N(%rip), %ecx
	xorl	%r8d, %r8d
	movl	36(%rsp), %r9d                  # 4-byte Reload
	callq	_Z37__device_stub__computeVarianceSquaresPfmS_iii
.LBB7_20:
	movl	N(%rip), %edi
	movabsq	$4294967296, %rax               # imm = 0x100000000
	orq	%rax, %rdi
	movl	%ebx, %ebx
	orq	%rax, %rbx
	movl	$1, %esi
	movq	%rbx, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB7_22
# %bb.21:
	movq	24(%rsp), %rdi
	movq	40(%rsp), %rsi
	movq	48(%rsp), %rdx
	movl	N(%rip), %ecx
	movl	$1, %r8d
	movl	36(%rsp), %r9d                  # 4-byte Reload
	callq	_Z37__device_stub__computeVarianceSquaresPfmS_iii
.LBB7_22:
	movq	88(%rsp), %rdi                  # 8-byte Reload
	movl	$1, %esi
	movq	%r15, %rdx
	movl	$1, %ecx
	movq	128(%rsp), %r8                  # 8-byte Reload
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB7_24
# %bb.23:
	movq	24(%rsp), %rdi
	movq	40(%rsp), %rcx
	movl	N(%rip), %r8d
	movl	76(%rsp), %eax                  # 4-byte Reload
	movl	%eax, 8(%rsp)
	movl	$10, (%rsp)
	movq	%rdi, %rsi
	movq	%rcx, %rdx
	movl	$1, %r9d
	callq	_Z26__device_stub__computeSumsPfS_mmiiii
.LBB7_24:
	movl	N(%rip), %edi
	movabsq	$4294967296, %rax               # imm = 0x100000000
	orq	%rax, %rdi
	movl	$1, %esi
	movq	112(%rsp), %rdx                 # 8-byte Reload
	movl	$1, %ecx
	movq	120(%rsp), %r8                  # 8-byte Reload
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB7_26
# %bb.25:
	movq	24(%rsp), %rdi
	movq	40(%rsp), %rcx
	movl	N(%rip), %r8d
	movl	76(%rsp), %eax                  # 4-byte Reload
	movl	%eax, 8(%rsp)
	movl	$10, (%rsp)
	movq	%rdi, %rsi
	movq	%rcx, %rdx
	xorl	%r9d, %r9d
	callq	_Z26__device_stub__computeSumsPfS_mmiiii
.LBB7_26:
	movq	%r13, %rdi
	movl	$1, %esi
	movq	%r15, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB7_28
# %bb.27:
	movq	24(%rsp), %rdi
	movq	40(%rsp), %rsi
	movl	N(%rip), %edx
	leal	(%r12,%r12), %eax
	leal	(%rax,%rax,4), %r8d
	movq	64(%rsp), %rax
	movq	%rax, (%rsp)
	xorl	%ecx, %ecx
	movl	56(%rsp), %r9d                  # 4-byte Reload
	callq	_Z31__device_stub__computeFinalSumsPfmiiiiS_
.LBB7_28:
	movq	96(%rsp), %rdi                  # 8-byte Reload
	movl	$1, %esi
	movq	104(%rsp), %rdx                 # 8-byte Reload
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB7_30
# %bb.29:
	movq	24(%rsp), %rdi
	movq	40(%rsp), %rsi
	movl	N(%rip), %edx
	addl	%r12d, %r12d
	leal	(%r12,%r12,4), %r8d
	movq	64(%rsp), %rax
	movq	%rax, (%rsp)
	movl	$1, %ecx
	movl	56(%rsp), %r9d                  # 4-byte Reload
	callq	_Z31__device_stub__computeFinalSumsPfmiiiiS_
.LBB7_30:
	movq	24(%rsp), %rdi
	movq	40(%rsp), %rsi
	movslq	N(%rip), %r9
	leaq	(,%r9,4), %r8
	movl	$1, (%rsp)
	movq	%r14, %rdx
	movq	80(%rsp), %rcx                  # 8-byte Reload
	callq	hipMemcpy2D
	testl	%eax, %eax
	je	.LBB7_32
# %bb.31:
	movl	$.L.str.19, %edi
	xorl	%eax, %eax
	callq	printf
.LBB7_32:
	movq	%rbp, %rdi
	movl	$1, %esi
	movq	%r15, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB7_34
# %bb.33:
	movq	24(%rsp), %rdi
	movq	40(%rsp), %rsi
	movq	48(%rsp), %rdx
	movq	64(%rsp), %rcx
	movl	N(%rip), %r8d
	movl	36(%rsp), %eax                  # 4-byte Reload
	movl	%eax, (%rsp)
	xorl	%r9d, %r9d
	callq	_Z27__device_stub__computeNormsPfmS_S_iii
.LBB7_34:
	movl	N(%rip), %edi
	movabsq	$4294967296, %rax               # imm = 0x100000000
	orq	%rax, %rdi
	movl	$1, %esi
	movq	%rbx, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB7_36
# %bb.35:
	movq	24(%rsp), %rdi
	movq	40(%rsp), %rsi
	movq	48(%rsp), %rdx
	movq	64(%rsp), %rcx
	movl	N(%rip), %r8d
	movl	36(%rsp), %eax                  # 4-byte Reload
	movl	%eax, (%rsp)
	movl	$1, %r9d
	callq	_Z27__device_stub__computeNormsPfmS_S_iii
.LBB7_36:
	movq	24(%rsp), %rdx
	movq	40(%rsp), %rcx
	movslq	N(%rip), %r9
	leaq	(,%r9,4), %r8
	movl	$2, (%rsp)
	movq	%r14, %rdi
	movq	80(%rsp), %rsi                  # 8-byte Reload
	callq	hipMemcpy2D
	movq	24(%rsp), %rdi
	callq	hipFree
	movq	48(%rsp), %rdi
	callq	hipFree
	movq	64(%rsp), %rdi
	callq	hipFree
	movl	N(%rip), %eax
	testl	%eax, %eax
	jle	.LBB7_41
# %bb.37:                               # %.preheader.lr.ph
	xorl	%ecx, %ecx
	movl	$B, %edx
	xorl	%esi, %esi
.LBB7_38:                               # %.lr.ph238
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB7_39 Depth 2
	movl	%ecx, %edi
	leaq	(%r14,%rdi,4), %rdi
	xorl	%r8d, %r8d
.LBB7_39:                               #   Parent Loop BB7_38 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movss	(%rdi,%r8,4), %xmm0             # xmm0 = mem[0],zero,zero,zero
	movss	%xmm0, (%rdx,%r8,4)
	incq	%r8
	cmpq	%r8, %rax
	jne	.LBB7_39
# %bb.40:                               # %._crit_edge239
                                        #   in Loop: Header=BB7_38 Depth=1
	incq	%rsi
	addl	%eax, %ecx
	addq	$32000, %rdx                    # imm = 0x7D00
	cmpq	%rax, %rsi
	jne	.LBB7_38
.LBB7_41:                               # %._crit_edge241
	addq	$1608, %rsp                     # imm = 0x648
	.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_end7:
	.size	_Z14matrixNorm_GPUv, .Lfunc_end7-_Z14matrixNorm_GPUv
	.cfi_endproc
                                        # -- End function
	.globl	_Z26__device_stub__computeSumsPfS_mmiiii # -- Begin function _Z26__device_stub__computeSumsPfS_mmiiii
	.type	_Z26__device_stub__computeSumsPfS_mmiiii,@function
_Z26__device_stub__computeSumsPfS_mmiiii: # @_Z26__device_stub__computeSumsPfS_mmiiii
	.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	$160, %rsp
	.cfi_def_cfa_offset 208
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	56(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	48(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	40(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	12(%rsp), %rcx
	movl	%r8d, (%rcx)
	leaq	8(%rsp), %r8
	movl	%r9d, (%r8)
	leaq	96(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r8, 40(%rbx)
	leaq	208(%rsp), %rax
	movq	%rax, 48(%rbx)
	leaq	216(%rsp), %rax
	movq	%rax, 56(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%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	$_Z11computeSumsPfS_mmiiii, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$176, %rsp
	.cfi_adjust_cfa_offset -176
	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_end8:
	.size	_Z26__device_stub__computeSumsPfS_mmiiii, .Lfunc_end8-_Z26__device_stub__computeSumsPfS_mmiiii
	.cfi_endproc
                                        # -- End function
	.globl	_Z31__device_stub__computeFinalSumsPfmiiiiS_ # -- Begin function _Z31__device_stub__computeFinalSumsPfmiiiiS_
	.type	_Z31__device_stub__computeFinalSumsPfmiiiiS_,@function
_Z31__device_stub__computeFinalSumsPfmiiiiS_: # @_Z31__device_stub__computeFinalSumsPfmiiiiS_
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$144, %rsp
	.cfi_def_cfa_offset 192
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	12(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	8(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	4(%rsp), %rcx
	movl	%r8d, (%rcx)
	movq	%rsp, %r8
	movl	%r9d, (%r8)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r8, 40(%rbx)
	leaq	192(%rsp), %rax
	movq	%rax, 48(%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	$_Z16computeFinalSumsPfmiiiiS_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$160, %rsp
	.cfi_adjust_cfa_offset -160
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end9:
	.size	_Z31__device_stub__computeFinalSumsPfmiiiiS_, .Lfunc_end9-_Z31__device_stub__computeFinalSumsPfmiiiiS_
	.cfi_endproc
                                        # -- End function
	.globl	_Z37__device_stub__computeVarianceSquaresPfmS_iii # -- Begin function _Z37__device_stub__computeVarianceSquaresPfmS_iii
	.type	_Z37__device_stub__computeVarianceSquaresPfmS_iii,@function
_Z37__device_stub__computeVarianceSquaresPfmS_iii: # @_Z37__device_stub__computeVarianceSquaresPfmS_iii
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$144, %rsp
	.cfi_def_cfa_offset 192
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	56(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	48(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	40(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	20(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	16(%rsp), %rcx
	movl	%r8d, (%rcx)
	leaq	12(%rsp), %r8
	movl	%r9d, (%r8)
	leaq	96(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r8, 40(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%rsp), %r15
	leaq	32(%rsp), %r12
	leaq	24(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z22computeVarianceSquaresPfmS_iii, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$160, %rsp
	.cfi_adjust_cfa_offset -160
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end10:
	.size	_Z37__device_stub__computeVarianceSquaresPfmS_iii, .Lfunc_end10-_Z37__device_stub__computeVarianceSquaresPfmS_iii
	.cfi_endproc
                                        # -- End function
	.globl	_Z27__device_stub__computeNormsPfmS_S_iii # -- Begin function _Z27__device_stub__computeNormsPfmS_S_iii
	.type	_Z27__device_stub__computeNormsPfmS_S_iii,@function
_Z27__device_stub__computeNormsPfmS_S_iii: # @_Z27__device_stub__computeNormsPfmS_S_iii
	.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	$160, %rsp
	.cfi_def_cfa_offset 208
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	56(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	48(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	40(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	12(%rsp), %rcx
	movl	%r8d, (%rcx)
	leaq	8(%rsp), %r8
	movl	%r9d, (%r8)
	leaq	96(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r8, 40(%rbx)
	leaq	208(%rsp), %rax
	movq	%rax, 48(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%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	$_Z12computeNormsPfmS_S_iii, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$176, %rsp
	.cfi_adjust_cfa_offset -176
	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_end11:
	.size	_Z27__device_stub__computeNormsPfmS_S_iii, .Lfunc_end11-_Z27__device_stub__computeNormsPfmS_S_iii
	.cfi_endproc
                                        # -- End function
	.globl	_Z8ceil_h_df                    # -- Begin function _Z8ceil_h_df
	.type	_Z8ceil_h_df,@function
_Z8ceil_h_df:                           # @_Z8ceil_h_df
	.cfi_startproc
# %bb.0:
	cvttss2si	%xmm0, %ecx
	cvttps2dq	%xmm0, %xmm1
	cvtdq2ps	%xmm1, %xmm1
	xorl	%eax, %eax
	ucomiss	%xmm1, %xmm0
	seta	%al
	addl	%ecx, %eax
	retq
.Lfunc_end12:
	.size	_Z8ceil_h_df, .Lfunc_end12-_Z8ceil_h_df
	.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	.LBB13_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB13_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z14blank_functionv, %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	$_Z11computeSumsPfS_mmiiii, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z16computeFinalSumsPfmiiiiS_, %esi
	movl	$.L__unnamed_3, %edx
	movl	$.L__unnamed_3, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z22computeVarianceSquaresPfmS_iii, %esi
	movl	$.L__unnamed_4, %edx
	movl	$.L__unnamed_4, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z12computeNormsPfmS_S_iii, %esi
	movl	$.L__unnamed_5, %edx
	movl	$.L__unnamed_5, %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_end13:
	.size	__hip_module_ctor, .Lfunc_end13-__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	.LBB14_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
.LBB14_2:
	retq
.Lfunc_end14:
	.size	__hip_module_dtor, .Lfunc_end14-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	N,@object                       # @N
	.bss
	.globl	N
	.p2align	2, 0x0
N:
	.long	0                               # 0x0
	.size	N, 4

	.type	A,@object                       # @A
	.globl	A
	.p2align	4, 0x0
A:
	.zero	256000000
	.size	A, 256000000

	.type	B,@object                       # @B
	.globl	B
	.p2align	4, 0x0
B:
	.zero	256000000
	.size	B, 256000000

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"Random seed = %i\n"
	.size	.L.str, 18

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"N = %i is out of range.\n"
	.size	.L.str.1, 25

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"Usage: %s <matrix_dimension> [random seed]\n"
	.size	.L.str.2, 44

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"\nMatrix dimension N = %i.\n"
	.size	.L.str.3, 27

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"\nA =\n\t"
	.size	.L.str.5, 7

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"%5.2f%s"
	.size	.L.str.6, 8

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

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	";\n\t"
	.size	.L.str.8, 4

	.type	.L.str.9,@object                # @.str.9
.L.str.9:
	.asciz	"\nB =\n\t"
	.size	.L.str.9, 7

	.type	.L.str.10,@object               # @.str.10
.L.str.10:
	.asciz	"%1.10f%s"
	.size	.L.str.10, 9

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

	.type	.L.str.13,@object               # @.str.13
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.13:
	.asciz	"\nElapsed time = %g ms.\n"
	.size	.L.str.13, 24

	.type	.L.str.14,@object               # @.str.14
.L.str.14:
	.asciz	"(CPU times are accurate to the nearest %g ms)\n"
	.size	.L.str.14, 47

	.type	.L.str.15,@object               # @.str.15
.L.str.15:
	.asciz	"My total CPU time for parent = %g ms.\n"
	.size	.L.str.15, 39

	.type	.L.str.16,@object               # @.str.16
.L.str.16:
	.asciz	"My system CPU time for parent = %g ms.\n"
	.size	.L.str.16, 40

	.type	.L.str.17,@object               # @.str.17
.L.str.17:
	.asciz	"My total CPU time for child processes = %g ms.\n"
	.size	.L.str.17, 48

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

	.type	_Z16computeFinalSumsPfmiiiiS_,@object # @_Z16computeFinalSumsPfmiiiiS_
	.globl	_Z16computeFinalSumsPfmiiiiS_
	.p2align	3, 0x0
_Z16computeFinalSumsPfmiiiiS_:
	.quad	_Z31__device_stub__computeFinalSumsPfmiiiiS_
	.size	_Z16computeFinalSumsPfmiiiiS_, 8

	.type	_Z22computeVarianceSquaresPfmS_iii,@object # @_Z22computeVarianceSquaresPfmS_iii
	.globl	_Z22computeVarianceSquaresPfmS_iii
	.p2align	3, 0x0
_Z22computeVarianceSquaresPfmS_iii:
	.quad	_Z37__device_stub__computeVarianceSquaresPfmS_iii
	.size	_Z22computeVarianceSquaresPfmS_iii, 8

	.type	_Z12computeNormsPfmS_S_iii,@object # @_Z12computeNormsPfmS_S_iii
	.globl	_Z12computeNormsPfmS_S_iii
	.p2align	3, 0x0
_Z12computeNormsPfmS_S_iii:
	.quad	_Z27__device_stub__computeNormsPfmS_S_iii
	.size	_Z12computeNormsPfmS_S_iii, 8

	.type	.L.str.19,@object               # @.str.19
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.19:
	.asciz	"ERROR"
	.size	.L.str.19, 6

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z14blank_functionv"
	.size	.L__unnamed_1, 20

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z11computeSumsPfS_mmiiii"
	.size	.L__unnamed_2, 26

	.type	.L__unnamed_3,@object           # @2
.L__unnamed_3:
	.asciz	"_Z16computeFinalSumsPfmiiiiS_"
	.size	.L__unnamed_3, 30

	.type	.L__unnamed_4,@object           # @3
.L__unnamed_4:
	.asciz	"_Z22computeVarianceSquaresPfmS_iii"
	.size	.L__unnamed_4, 35

	.type	.L__unnamed_5,@object           # @4
.L__unnamed_5:
	.asciz	"_Z12computeNormsPfmS_S_iii"
	.size	.L__unnamed_5, 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

	.type	.Lstr,@object                   # @str
	.section	.rodata.str1.1,"aMS",@progbits,1
.Lstr:
	.asciz	"\nInitializing..."
	.size	.Lstr, 17

	.type	.Lstr.1,@object                 # @str.1
.Lstr.1:
	.asciz	"\nStarting clock."
	.size	.Lstr.1, 17

	.type	.Lstr.2,@object                 # @str.2
.Lstr.2:
	.asciz	"Stopped clock."
	.size	.Lstr.2, 15

	.type	.Lstr.3,@object                 # @str.3
.Lstr.3:
	.asciz	"--------------------------------------------"
	.size	.Lstr.3, 45

	.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 _Z29__device_stub__blank_functionv
	.addrsig_sym _Z26__device_stub__computeSumsPfS_mmiiii
	.addrsig_sym _Z31__device_stub__computeFinalSumsPfmiiiiS_
	.addrsig_sym _Z37__device_stub__computeVarianceSquaresPfmS_iii
	.addrsig_sym _Z27__device_stub__computeNormsPfmS_S_iii
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym A
	.addrsig_sym B
	.addrsig_sym _Z14blank_functionv
	.addrsig_sym _Z11computeSumsPfS_mmiiii
	.addrsig_sym _Z16computeFinalSumsPfmiiiiS_
	.addrsig_sym _Z22computeVarianceSquaresPfmS_iii
	.addrsig_sym _Z12computeNormsPfmS_S_iii
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z14blank_functionv     ; -- Begin function _Z14blank_functionv
	.globl	_Z14blank_functionv
	.p2align	8
	.type	_Z14blank_functionv,@function
_Z14blank_functionv:                    ; @_Z14blank_functionv
; %bb.0:
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z14blank_functionv
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 0
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 0
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 1
		.amdhsa_next_free_sgpr 1
		.amdhsa_reserve_vcc 0
		.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	_Z14blank_functionv, .Lfunc_end0-_Z14blank_functionv
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 4
; NumSgprs: 0
; NumVgprs: 0
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 0
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 1
; NumVGPRsForWavesPerEU: 1
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.protected	_Z11computeSumsPfS_mmiiii ; -- Begin function _Z11computeSumsPfS_mmiiii
	.globl	_Z11computeSumsPfS_mmiiii
	.p2align	8
	.type	_Z11computeSumsPfS_mmiiii,@function
_Z11computeSumsPfS_mmiiii:              ; @_Z11computeSumsPfS_mmiiii
; %bb.0:
	s_load_b128 s[16:19], s[0:1], 0x20
	s_mov_b32 s2, 0
	s_waitcnt lgkmcnt(0)
	v_mul_lo_u32 v1, v0, s18
	s_cmp_lg_u32 s17, 0
	s_cbranch_scc0 .LBB1_2
; %bb.1:
	s_load_b32 s3, s[0:1], 0x3c
	v_mul_lo_u32 v4, v0, s18
	s_waitcnt lgkmcnt(0)
	s_and_b32 s3, s3, 0xffff
	s_delay_alu instid0(VALU_DEP_2) | instid1(SALU_CYCLE_1)
	v_mad_u64_u32 v[2:3], null, s15, s3, v[0:1]
	s_delay_alu instid0(VALU_DEP_1)
	v_mul_lo_u32 v2, v2, s18
	s_branch .LBB1_3
.LBB1_2:
	s_mov_b32 s2, -1
                                        ; implicit-def: $vgpr2
                                        ; implicit-def: $vgpr4
.LBB1_3:                                ; %Flow121
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s2
	s_cbranch_vccnz .LBB1_5
; %bb.4:
	s_delay_alu instid0(VALU_DEP_1)
	v_add_nc_u32_e32 v2, s19, v1
	v_mov_b32_e32 v4, v1
.LBB1_5:
	s_load_b256 s[4:11], s[0:1], 0x0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v5, 31, v2
	v_lshl_add_u32 v3, v4, 2, 0
	s_cmp_gt_i32 s18, 0
	s_cselect_b32 s12, -1, 0
	s_cmp_lt_i32 s18, 1
	s_cbranch_scc1 .LBB1_8
; %bb.6:                                ; %.lr.ph
	s_ashr_i32 s15, s14, 31
	s_waitcnt lgkmcnt(0)
	v_mul_lo_u32 v7, s9, v2
	s_lshl_b64 s[2:3], s[14:15], 2
	v_mul_lo_u32 v8, s8, v5
	s_add_u32 s2, s4, s2
	s_addc_u32 s3, s5, s3
	v_lshl_add_u32 v6, v4, 2, 0
	v_mad_u64_u32 v[0:1], null, s8, v2, s[2:3]
	s_mov_b32 s2, s18
	s_delay_alu instid0(VALU_DEP_1)
	v_add3_u32 v1, v7, v1, v8
.LBB1_7:                                ; =>This Inner Loop Header: Depth=1
	global_load_b32 v7, v[0:1], off
	v_add_co_u32 v0, vcc_lo, v0, s8
	v_add_co_ci_u32_e32 v1, vcc_lo, s9, v1, vcc_lo
	s_add_i32 s2, s2, -1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_lg_u32 s2, 0
	s_waitcnt vmcnt(0)
	ds_store_b32 v6, v7
	v_add_nc_u32_e32 v6, 4, v6
	s_cbranch_scc1 .LBB1_7
.LBB1_8:                                ; %._crit_edge
	s_add_i32 s2, s16, -1
	s_cmp_gt_i32 s16, 1
	v_cmp_gt_i32_e32 vcc_lo, s2, v2
	s_cselect_b32 s2, -1, 0
	s_waitcnt lgkmcnt(0)
	s_mov_b32 s8, 1
	s_and_b32 s2, vcc_lo, s2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_saveexec_b32 s4, s2
	s_cbranch_execz .LBB1_19
; %bb.9:                                ; %.lr.ph62
	v_ashrrev_i32_e32 v0, 31, v4
	v_lshl_add_u32 v1, v4, 2, 0
	s_add_u32 s2, s0, 48
	s_addc_u32 s3, s1, 0
	s_mov_b32 s1, 0
	v_add_nc_u32_e32 v6, v4, v0
                                        ; implicit-def: $sgpr5
	s_delay_alu instid0(VALU_DEP_1)
	v_xor_b32_e32 v6, v6, v0
.LBB1_10:                               ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB1_13 Depth 2
	s_mul_i32 s9, s8, s18
	s_or_b32 s5, s5, exec_lo
	s_ashr_i32 s0, s9, 31
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_add_i32 s13, s9, s0
	s_xor_b32 s0, s13, s0
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cvt_f32_u32_e32 v7, s0
	s_sub_i32 s13, 0, s0
	v_rcp_iflag_f32_e32 v7, v7
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v7, 0x4f7ffffe, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_u32_f32_e32 v7, v7
	v_mul_lo_u32 v8, s13, v7
	s_mov_b32 s13, exec_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v8, v7, v8
	v_add_nc_u32_e32 v7, v7, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v7, v6, v7
	v_mul_lo_u32 v7, v7, s0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v7, v6, v7
	v_subrev_nc_u32_e32 v8, s0, v7
	v_cmp_le_u32_e32 vcc_lo, s0, v7
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v7, v7, v8, vcc_lo
	v_subrev_nc_u32_e32 v8, s0, v7
	v_cmp_le_u32_e32 vcc_lo, s0, v7
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v7, v7, v8, vcc_lo
	v_xor_b32_e32 v7, v7, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v7, v7, v0
	v_cmpx_eq_u32_e32 0, v7
	s_cbranch_execz .LBB1_18
; %bb.11:                               ; %.preheader
                                        ;   in Loop: Header=BB1_10 Depth=1
	s_cmp_ge_i32 s8, s9
	s_cbranch_scc1 .LBB1_17
; %bb.12:                               ; %.lr.ph59
                                        ;   in Loop: Header=BB1_10 Depth=1
	s_load_b32 s0, s[2:3], 0xc
	v_add_nc_u32_e32 v7, s8, v4
	v_add_nc_u32_e32 v8, s9, v4
	s_lshl_b32 s15, s8, 2
	s_mov_b32 s17, 0
                                        ; implicit-def: $sgpr20
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(SALU_CYCLE_1)
	v_lshl_add_u32 v9, v7, 2, 0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s19, s0, 0xffff
	s_mul_i32 s19, s19, s18
.LBB1_13:                               ;   Parent Loop BB1_10 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_cmp_gt_u32_e32 vcc_lo, s19, v7
	v_cmp_gt_i32_e64 s0, s16, v7
	s_or_b32 s20, s20, exec_lo
	s_and_b32 s21, s0, vcc_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_saveexec_b32 s0, s21
	s_cbranch_execz .LBB1_15
; %bb.14:                               ;   in Loop: Header=BB1_13 Depth=2
	ds_load_b32 v10, v9
	ds_load_b32 v11, v1
	v_add_nc_u32_e32 v7, s8, v7
	s_and_not1_b32 s20, s20, exec_lo
	v_add_nc_u32_e32 v9, s15, v9
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_cmp_ge_i32_e32 vcc_lo, v7, v8
	s_and_b32 s21, vcc_lo, exec_lo
	s_or_b32 s20, s20, s21
	s_waitcnt lgkmcnt(0)
	v_add_f32_e32 v10, v10, v11
	ds_store_b32 v1, v10
.LBB1_15:                               ; %Flow113
                                        ;   in Loop: Header=BB1_13 Depth=2
	s_or_b32 exec_lo, exec_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s0, exec_lo, s20
	s_or_b32 s17, s0, s17
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s17
	s_cbranch_execnz .LBB1_13
; %bb.16:                               ; %Flow114
                                        ;   in Loop: Header=BB1_10 Depth=1
	s_or_b32 exec_lo, exec_lo, s17
.LBB1_17:                               ; %.critedge
                                        ;   in Loop: Header=BB1_10 Depth=1
	s_cmp_ge_i32 s9, s16
	s_waitcnt lgkmcnt(0)
	s_cselect_b32 s0, -1, 0
	s_and_not1_b32 s5, s5, exec_lo
	s_and_b32 s0, s0, exec_lo
	s_barrier
	s_or_b32 s5, s5, s0
	buffer_gl0_inv
.LBB1_18:                               ; %Flow116
                                        ;   in Loop: Header=BB1_10 Depth=1
	s_or_b32 exec_lo, exec_lo, s13
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(SALU_CYCLE_1)
	s_and_b32 s0, exec_lo, s5
	s_mov_b32 s8, s9
	s_or_b32 s1, s0, s1
	s_and_not1_b32 exec_lo, exec_lo, s1
	s_cbranch_execnz .LBB1_10
.LBB1_19:                               ; %Flow118
	s_or_b32 exec_lo, exec_lo, s4
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s12
	s_cbranch_vccnz .LBB1_22
; %bb.20:                               ; %.lr.ph66
	s_ashr_i32 s15, s14, 31
	v_mul_lo_u32 v4, s11, v2
	s_lshl_b64 s[0:1], s[14:15], 2
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	s_add_u32 s0, s6, s0
	s_addc_u32 s1, s7, s1
	v_mad_u64_u32 v[0:1], null, s10, v2, s[0:1]
	v_mul_lo_u32 v2, s10, v5
	s_delay_alu instid0(VALU_DEP_1)
	v_add3_u32 v1, v4, v1, v2
.LBB1_21:                               ; =>This Inner Loop Header: Depth=1
	ds_load_b32 v2, v3
	v_add_nc_u32_e32 v3, 4, v3
	s_add_i32 s18, s18, -1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_lg_u32 s18, 0
	s_waitcnt lgkmcnt(0)
	global_store_b32 v[0:1], v2, off
	v_add_co_u32 v0, vcc_lo, v0, s10
	v_add_co_ci_u32_e32 v1, vcc_lo, s11, v1, vcc_lo
	s_cbranch_scc1 .LBB1_21
.LBB1_22:                               ; %._crit_edge67
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z11computeSumsPfS_mmiiii
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 304
		.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 12
		.amdhsa_next_free_sgpr 22
		.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	_Z11computeSumsPfS_mmiiii, .Lfunc_end1-_Z11computeSumsPfS_mmiiii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 908
; NumSgprs: 24
; NumVgprs: 12
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 24
; NumVGPRsForWavesPerEU: 12
; 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
	.protected	_Z16computeFinalSumsPfmiiiiS_ ; -- Begin function _Z16computeFinalSumsPfmiiiiS_
	.globl	_Z16computeFinalSumsPfmiiiiS_
	.p2align	8
	.type	_Z16computeFinalSumsPfmiiiiS_,@function
_Z16computeFinalSumsPfmiiiiS_:          ; @_Z16computeFinalSumsPfmiiiiS_
; %bb.0:
	s_load_b128 s[4:7], s[0:1], 0x10
	s_mov_b32 s2, 0
	s_waitcnt lgkmcnt(0)
	s_cmp_lg_u32 s5, 0
	s_cbranch_scc0 .LBB2_2
; %bb.1:
	v_mov_b32_e32 v1, v0
	s_branch .LBB2_3
.LBB2_2:
	s_mov_b32 s2, -1
                                        ; implicit-def: $vgpr1
.LBB2_3:                                ; %Flow49
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s2
	s_cbranch_vccnz .LBB2_5
; %bb.4:
	s_load_b32 s2, s[0:1], 0x34
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_mul_i32 s2, s15, s2
	v_mov_b32_e32 v1, s2
	s_branch .LBB2_6
.LBB2_5:
	v_mov_b32_e32 v0, s7
.LBB2_6:
	s_load_b128 s[8:11], s[0:1], 0x0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v0, v1, v0
	s_cmp_ge_i32 s6, s4
	v_ashrrev_i32_e32 v1, 31, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[2:3], 2, v[0:1]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v0, vcc_lo, s8, v2
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v1, vcc_lo, s9, v3, vcc_lo
	global_load_b32 v6, v[0:1], off
	s_cbranch_scc1 .LBB2_9
; %bb.7:                                ; %.lr.ph.preheader
	s_ashr_i32 s2, s6, 31
	s_mul_hi_u32 s3, s10, s6
	s_mul_i32 s2, s10, s2
	s_mul_i32 s5, s11, s6
	s_add_i32 s3, s3, s2
	s_mul_i32 s2, s10, s6
	s_add_i32 s3, s3, s5
	s_add_u32 s5, s8, s2
	s_addc_u32 s7, s9, s3
	v_add_co_u32 v4, vcc_lo, s5, v2
	v_add_co_ci_u32_e32 v5, vcc_lo, s7, v3, vcc_lo
	s_mov_b32 s5, s6
.LBB2_8:                                ; %.lr.ph
                                        ; =>This Inner Loop Header: Depth=1
	global_load_b32 v7, v[4:5], off
	v_add_co_u32 v4, vcc_lo, v4, s2
	v_add_co_ci_u32_e32 v5, vcc_lo, s3, v5, vcc_lo
	s_add_i32 s5, s5, s6
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_lt_i32 s5, s4
	s_waitcnt vmcnt(0)
	v_add_f32_e32 v6, v7, v6
	global_store_b32 v[0:1], v6, off
	s_cbranch_scc1 .LBB2_8
.LBB2_9:                                ; %._crit_edge
	v_cvt_f32_i32_e32 v4, s4
	s_load_b64 s[0:1], s[0:1], 0x20
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_div_scale_f32 v5, null, v4, v4, v6
	v_div_scale_f32 v9, vcc_lo, v6, v4, v6
	v_rcp_f32_e32 v7, v5
	s_waitcnt_depctr 0xfff
	v_fma_f32 v8, -v5, v7, 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v7, v8, v7
	v_mul_f32_e32 v8, v9, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v10, -v5, v8, v9
	v_fmac_f32_e32 v8, v10, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v5, -v5, v8, v9
	v_div_fmas_f32 v5, v5, v7, v8
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v2, vcc_lo, s0, v2
	v_add_co_ci_u32_e32 v3, vcc_lo, s1, v3, vcc_lo
	s_delay_alu instid0(VALU_DEP_3)
	v_div_fixup_f32 v4, v5, v4, v6
	global_store_b32 v[0:1], v4, off
	global_store_b32 v[2:3], v4, off
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z16computeFinalSumsPfmiiiiS_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 296
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 11
		.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_end2:
	.size	_Z16computeFinalSumsPfmiiiiS_, .Lfunc_end2-_Z16computeFinalSumsPfmiiiiS_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 416
; NumSgprs: 18
; NumVgprs: 11
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 11
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.protected	_Z22computeVarianceSquaresPfmS_iii ; -- Begin function _Z22computeVarianceSquaresPfmS_iii
	.globl	_Z22computeVarianceSquaresPfmS_iii
	.p2align	8
	.type	_Z22computeVarianceSquaresPfmS_iii,@function
_Z22computeVarianceSquaresPfmS_iii:     ; @_Z22computeVarianceSquaresPfmS_iii
; %bb.0:
	s_load_b64 s[2:3], s[0:1], 0x1c
	s_waitcnt lgkmcnt(0)
	s_cmp_lg_u32 s2, 1
	s_cbranch_scc0 .LBB3_2
; %bb.1:
	s_load_b32 s2, s[0:1], 0x34
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1)
	v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
	s_mov_b32 s2, 0
	s_branch .LBB3_3
.LBB3_2:
	s_mov_b32 s2, -1
                                        ; implicit-def: $vgpr1
.LBB3_3:                                ; %Flow
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s2
	s_cbranch_vccnz .LBB3_5
; %bb.4:
	v_add_nc_u32_e32 v1, s3, v0
.LBB3_5:
	s_load_b128 s[4:7], s[0:1], 0x0
	s_ashr_i32 s15, s14, 31
	s_delay_alu instid0(VALU_DEP_1)
	v_ashrrev_i32_e32 v0, 31, v1
	s_lshl_b64 s[2:3], s[14:15], 2
	s_load_b64 s[0:1], s[0:1], 0x10
	s_waitcnt lgkmcnt(0)
	s_add_u32 s4, s2, s4
	s_addc_u32 s5, s3, s5
	v_mul_lo_u32 v4, v1, s7
	v_mul_lo_u32 v0, v0, s6
	v_mad_u64_u32 v[2:3], null, v1, s6, s[4:5]
	s_add_u32 s0, s0, s2
	s_addc_u32 s1, s1, s3
	s_load_b32 s0, s[0:1], 0x0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_4) | instid1(VALU_DEP_1)
	v_add3_u32 v3, v0, v3, v4
	global_load_b32 v0, v[2:3], off
	s_waitcnt vmcnt(0) lgkmcnt(0)
	v_subrev_f32_e32 v0, s0, v0
	s_mov_b32 s0, 0x3e76c4e1
	v_frexp_mant_f32_e64 v1, |v0|
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cmp_gt_f32_e32 vcc_lo, 0x3f2aaaab, v1
	v_cndmask_b32_e64 v4, 0, 1, vcc_lo
	v_ldexp_f32 v1, v1, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_f32_e32 v4, 1.0, v1
	v_add_f32_e32 v6, -1.0, v1
	v_add_f32_e32 v8, -1.0, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_sub_f32_e32 v1, v1, v8
	v_rcp_f32_e32 v5, v4
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v7, v6, v5
	v_mul_f32_e32 v9, v4, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v4, v7, v4, -v9
	v_fmac_f32_e32 v4, v7, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v1, v9, v4
	v_sub_f32_e32 v8, v6, v1
	v_sub_f32_e32 v9, v1, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_sub_f32_e32 v4, v9, v4
	v_sub_f32_e32 v6, v6, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_f32_e32 v1, v6, v1
	v_add_f32_e32 v1, v4, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v1, v8, v1
	v_mul_f32_e32 v1, v5, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v4, v7, v1
	v_sub_f32_e32 v5, v4, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_mul_f32 v6, v4, v4 :: v_dual_sub_f32 v1, v1, v5
	v_fma_f32 v5, v4, v4, -v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v7, v1, v1
	v_fmac_f32_e32 v5, v4, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v7, v6, v5
	v_fmaak_f32 v8, s0, v7, 0x3e91f4c4
	v_sub_f32_e32 v6, v7, v6
	v_cmp_neq_f32_e64 s0, 0x7f800000, |v0|
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_fmaak_f32 v8, v7, v8, 0x3ecccdef :: v_dual_sub_f32 v5, v5, v6
	v_mul_f32_e32 v9, v7, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v6, v7, v8, -v9
	v_fmac_f32_e32 v6, v5, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v8, v9, v6
	v_dual_sub_f32 v9, v8, v9 :: v_dual_add_f32 v10, 0x3f2aaaaa, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_dual_sub_f32 v6, v6, v9 :: v_dual_add_f32 v9, 0xbf2aaaaa, v10
	v_mul_f32_e32 v11, v4, v7
	v_sub_f32_e32 v8, v8, v9
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v12, v7, v4, -v11
	v_fmac_f32_e32 v12, v7, v1
	v_ldexp_f32 v1, v1, 1
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v12, v5, v4
	v_ldexp_f32 v4, v4, 1
	v_dual_add_f32 v7, v11, v12 :: v_dual_add_f32 v6, 0x31739010, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v6, v6, v8
	v_add_f32_e32 v5, v10, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_mul_f32_e32 v9, v7, v5
	v_sub_f32_e32 v8, v10, v5
	v_sub_f32_e32 v10, v7, v11
	v_add_f32_e32 v6, v6, v8
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f32 v8, v7, v5, -v9
	v_sub_f32_e32 v10, v12, v10
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v8, v7, v6
	v_frexp_exp_i32_f32_e32 v6, v0
	v_fmac_f32_e32 v8, v10, v5
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_subrev_co_ci_u32_e32 v5, vcc_lo, 0, v6, vcc_lo
	v_add_f32_e32 v6, v9, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cvt_f32_i32_e32 v5, v5
	v_add_f32_e32 v7, v4, v6
	v_sub_f32_e32 v9, v6, v9
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_f32_e32 v10, 0x3f317218, v5
	v_sub_f32_e32 v4, v7, v4
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_sub_f32_e32 v8, v8, v9
	v_fma_f32 v9, 0x3f317218, v5, -v10
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_sub_f32_e32 v4, v6, v4
	v_add_f32_e32 v1, v1, v8
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmamk_f32 v5, v5, 0xb102e308, v9
	v_dual_add_f32 v1, v1, v4 :: v_dual_add_f32 v4, v10, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v6, v7, v1
	v_dual_add_f32 v8, v4, v6 :: v_dual_sub_f32 v7, v6, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_f32_e32 v9, v8, v4
	v_dual_sub_f32 v10, v4, v10 :: v_dual_sub_f32 v1, v1, v7
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_f32_e32 v11, v8, v9
	v_dual_sub_f32 v5, v5, v10 :: v_dual_sub_f32 v6, v6, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_sub_f32 v4, v4, v11 :: v_dual_add_f32 v7, v5, v1
	v_add_f32_e32 v4, v6, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_f32_e32 v6, v7, v5
	v_add_f32_e32 v4, v7, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_sub_f32_e32 v7, v7, v6
	v_sub_f32_e32 v1, v1, v6
	v_sub_f32_e32 v5, v5, v7
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f32_e32 v9, v8, v4
	v_add_f32_e32 v1, v1, v5
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_f32_e32 v6, v9, v8
	v_sub_f32_e32 v4, v4, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v1, v1, v4
	v_add_f32_e32 v4, v9, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_sub_f32_e32 v5, v4, v9
	v_add_f32_e32 v6, v4, v4
	v_mul_f32_e32 v7, 0, v4
	v_sub_f32_e32 v1, v1, v5
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_fma_f32 v4, v4, 2.0, -v6
	v_cmp_class_f32_e64 vcc_lo, v6, 0x204
	v_fmac_f32_e32 v7, 2.0, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v1, v4, v7
	v_add_f32_e32 v4, v6, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v5, v4, v6, vcc_lo
	v_cmp_eq_f32_e32 vcc_lo, 0x42b17218, v5
	v_cndmask_b32_e64 v7, 0, 0x37000000, vcc_lo
	v_cmp_neq_f32_e64 vcc_lo, 0x7f800000, |v5|
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_f32_e32 v8, v5, v7
	v_mul_f32_e32 v9, 0x3fb8aa3b, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_fma_f32 v10, 0x3fb8aa3b, v8, -v9
	v_rndne_f32_e32 v11, v9
	v_dual_fmamk_f32 v10, v8, 0x32a5705f, v10 :: v_dual_sub_f32 v9, v9, v11
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_add_f32_e32 v9, v9, v10
	v_sub_f32_e32 v4, v4, v6
	v_cvt_i32_f32_e32 v6, v11
	v_exp_f32_e32 v9, v9
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_f32_e32 v1, v1, v4
	v_cndmask_b32_e32 v1, 0, v1, vcc_lo
	v_cmp_ngt_f32_e32 vcc_lo, 0xc2ce8ed0, v8
	s_waitcnt_depctr 0xfff
	v_ldexp_f32 v4, v9, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_dual_add_f32 v1, v7, v1 :: v_dual_cndmask_b32 v4, 0, v4
	v_cmp_nlt_f32_e32 vcc_lo, 0x42b17218, v8
	v_cndmask_b32_e32 v4, 0x7f800000, v4, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_fma_f32 v1, v4, v1, v4
	v_cmp_eq_f32_e32 vcc_lo, 0x7f800000, v4
	v_cndmask_b32_e32 v1, v1, v4, vcc_lo
	v_cmp_neq_f32_e32 vcc_lo, 0, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e64 v1, 0x7f800000, |v1|, s0
	v_cndmask_b32_e32 v0, 0, v1, vcc_lo
	global_store_b32 v[2:3], v0, off
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z22computeVarianceSquaresPfmS_iii
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 296
		.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 13
		.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_end3:
	.size	_Z22computeVarianceSquaresPfmS_iii, .Lfunc_end3-_Z22computeVarianceSquaresPfmS_iii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 1136
; NumSgprs: 18
; NumVgprs: 13
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 13
; 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
	.protected	_Z12computeNormsPfmS_S_iii ; -- Begin function _Z12computeNormsPfmS_S_iii
	.globl	_Z12computeNormsPfmS_S_iii
	.p2align	8
	.type	_Z12computeNormsPfmS_S_iii,@function
_Z12computeNormsPfmS_S_iii:             ; @_Z12computeNormsPfmS_S_iii
; %bb.0:
	s_load_b64 s[2:3], s[0:1], 0x24
	s_waitcnt lgkmcnt(0)
	s_cmp_lg_u32 s2, 1
	s_cbranch_scc0 .LBB4_2
; %bb.1:
	s_load_b32 s2, s[0:1], 0x3c
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1)
	v_mad_u64_u32 v[2:3], null, s15, s2, v[0:1]
	s_mov_b32 s2, 0
	s_branch .LBB4_3
.LBB4_2:
	s_mov_b32 s2, -1
                                        ; implicit-def: $vgpr2
.LBB4_3:                                ; %Flow
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s2
	s_cbranch_vccnz .LBB4_5
; %bb.4:
	v_add_nc_u32_e32 v2, s3, v0
.LBB4_5:
	s_load_b256 s[0:7], s[0:1], 0x0
	s_ashr_i32 s15, s14, 31
	s_delay_alu instid0(VALU_DEP_1)
	v_ashrrev_i32_e32 v3, 31, v2
	s_lshl_b64 s[8:9], s[14:15], 2
	s_waitcnt lgkmcnt(0)
	s_add_u32 s6, s6, s8
	s_addc_u32 s7, s7, s9
	s_add_u32 s0, s8, s0
	s_load_b32 s6, s[6:7], 0x0
	s_addc_u32 s1, s9, s1
	v_mul_lo_u32 v3, v3, s2
	v_mad_u64_u32 v[0:1], null, v2, s2, s[0:1]
	v_mul_lo_u32 v2, v2, s3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_add3_u32 v1, v3, v1, v2
	s_waitcnt lgkmcnt(0)
	v_cmp_eq_f32_e64 s0, s6, 0
	s_and_b32 vcc_lo, exec_lo, s0
	s_cbranch_vccnz .LBB4_7
; %bb.6:
	global_load_b32 v2, v[0:1], off
	v_mul_f32_e64 v3, 0x4f800000, s6
	v_cmp_gt_f32_e64 vcc_lo, 0xf800000, s6
	s_add_u32 s2, s4, s8
	s_addc_u32 s3, s5, s9
	s_load_b32 s1, s[2:3], 0x0
	v_cndmask_b32_e32 v3, s6, v3, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_2)
	v_sqrt_f32_e32 v4, v3
	s_waitcnt_depctr 0xfff
	v_add_nc_u32_e32 v6, 1, v4
	v_add_nc_u32_e32 v5, -1, v4
	v_fma_f32 v8, -v6, v4, v3
	s_waitcnt vmcnt(0) lgkmcnt(0)
	v_subrev_f32_e32 v2, s1, v2
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v7, -v5, v4, v3
	v_cmp_ge_f32_e64 s0, 0, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cndmask_b32_e64 v4, v4, v5, s0
	v_cmp_lt_f32_e64 s0, 0, v8
	v_cndmask_b32_e64 v4, v4, v6, s0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v5, 0x37800000, v4
	v_cndmask_b32_e32 v4, v4, v5, vcc_lo
	v_cmp_class_f32_e64 vcc_lo, v3, 0x260
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v3, v4, v3, vcc_lo
	v_div_scale_f32 v4, null, v3, v3, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f32_e32 v5, v4
	s_waitcnt_depctr 0xfff
	v_fma_f32 v6, -v4, v5, 1.0
	v_fmac_f32_e32 v5, v6, v5
	v_div_scale_f32 v6, vcc_lo, v2, v3, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v7, v6, v5
	v_fma_f32 v8, -v4, v7, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v7, v8, v5
	v_fma_f32 v4, -v4, v7, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fmas_f32 v4, v4, v5, v7
	v_div_fixup_f32 v2, v4, v3, v2
	s_branch .LBB4_8
.LBB4_7:
	v_mov_b32_e32 v2, 0
.LBB4_8:
	global_store_b32 v[0:1], v2, off
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z12computeNormsPfmS_S_iii
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 304
		.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 9
		.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_end4:
	.size	_Z12computeNormsPfmS_S_iii, .Lfunc_end4-_Z12computeNormsPfmS_S_iii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 480
; NumSgprs: 18
; NumVgprs: 9
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 9
; 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:           []
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 4
    .kernarg_segment_size: 0
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z14blank_functionv
    .private_segment_fixed_size: 0
    .sgpr_count:     0
    .sgpr_spill_count: 0
    .symbol:         _Z14blank_functionv.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     0
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .offset:         16
        .size:           8
        .value_kind:     by_value
      - .offset:         24
        .size:           8
        .value_kind:     by_value
      - .offset:         32
        .size:           4
        .value_kind:     by_value
      - .offset:         36
        .size:           4
        .value_kind:     by_value
      - .offset:         40
        .size:           4
        .value_kind:     by_value
      - .offset:         44
        .size:           4
        .value_kind:     by_value
      - .offset:         48
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         52
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         56
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         60
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         62
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         64
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         66
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         68
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         70
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         96
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         104
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         112
        .size:           2
        .value_kind:     hidden_grid_dims
      - .offset:         168
        .size:           4
        .value_kind:     hidden_dynamic_lds_size
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 304
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z11computeSumsPfS_mmiiii
    .private_segment_fixed_size: 0
    .sgpr_count:     24
    .sgpr_spill_count: 0
    .symbol:         _Z11computeSumsPfS_mmiiii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     12
    .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:           8
        .value_kind:     by_value
      - .offset:         16
        .size:           4
        .value_kind:     by_value
      - .offset:         20
        .size:           4
        .value_kind:     by_value
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         28
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         32
        .size:           8
        .value_kind:     global_buffer
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         44
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         48
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         52
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         54
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         56
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         58
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         60
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         62
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         96
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         104
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 296
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z16computeFinalSumsPfmiiiiS_
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z16computeFinalSumsPfmiiiiS_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     11
    .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:           8
        .value_kind:     by_value
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         28
        .size:           4
        .value_kind:     by_value
      - .offset:         32
        .size:           4
        .value_kind:     by_value
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         44
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         48
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         52
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         54
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         56
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         58
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         60
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         62
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         96
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         104
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 296
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z22computeVarianceSquaresPfmS_iii
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z22computeVarianceSquaresPfmS_iii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     13
    .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:           8
        .value_kind:     by_value
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .offset:         32
        .size:           4
        .value_kind:     by_value
      - .offset:         36
        .size:           4
        .value_kind:     by_value
      - .offset:         40
        .size:           4
        .value_kind:     by_value
      - .offset:         48
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         52
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         56
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         60
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         62
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         64
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         66
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         68
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         70
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         96
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         104
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         112
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 304
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z12computeNormsPfmS_S_iii
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z12computeNormsPfmS_S_iii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     9
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
