/* 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();
void cudaErrorCheck(cudaError_t err, const char *s);
__global__ void matrixCuda(float *d_A,int N);
__global__ void BlockMean(float *d_A,float *d_Sum, int N);
__global__ void BlockDev(float *d_A, float *d_Dev, float *d_mu,int N);
__global__ void Normalize(float *d_A, float *d_mu, float *d_sigma, int N);

/* 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;
		}
	}
	/*
	   for (col = 0; col < N; col++) {
	   for (row = 0; row < N; row++) {
	   A[row][col] = col + row;
	   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) {
	/* 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();

	/* 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.
 */


/* Matrix Normalization for the Cuda platform
   Overview
   1.Copy matrix A to Device
   2.CUDA: Split matrix into 2D grid of blocks, each block calculates a partial sum 
   for each column in their section via BlockMean.
   3.Sequentially add all of these partial sums to form the total mean for
   each column.
   4.CUDA: Each block calculate a partial sum of the mean difference for each
   column in their section via BlockDev.
   5.Sequentially add all of these partial mean differences to form the total
   standard deviation for each column.
   6.CUDA: Each block normalize their portion of the matrix using the means
   and standard deviations calculated.
   7.Copy matrix A to Host's B
   */

//Number of threads per block
#define BlockSize 32
void matrixNorm() 
{
	printf("Executing on GPU\n");

	//Set grid size to divide among number of threads
	int GridSize = ceil((float)N/BlockSize); 

	//Create CUDA grid and block size for matrix
	dim3 grid(GridSize,GridSize);
	dim3 block(BlockSize,BlockSize);

	//Device matrices and vectors for calculation 
	float *d_A; //the matrix 
	float *d_Sum; //a partial sum holder
	float *d_Dev; //a partial deviation holder
	float *d_mu; //a vector of means for each column
	float *d_sigma; //a vector of standard deviations for each column

	//Host copies of variables to initalize the device's copies to 0's
	float *h_Sum = (float *)malloc(GridSize*N*sizeof(float));
	float *h_Dev = (float *)malloc(GridSize*N*sizeof(float));
	float *h_mu = (float *)malloc(N*sizeof(float));
	float *h_sigma = (float *)malloc(N*sizeof(float));
	size_t size = N*N*sizeof(float);

	// Allocate Matrices on Devices 
	cudaErrorCheck(cudaMalloc((void **)&d_A,size), "cudaMalloc A");
	cudaErrorCheck(cudaMalloc(&d_Sum,GridSize*N*sizeof(float)), "cudaMalloc d_Sum" );
	cudaErrorCheck(cudaMalloc(&d_Dev,GridSize*N*sizeof(float)), "cudaMalloc d_Dev ");
	cudaErrorCheck(cudaMalloc((void **)&d_mu,N*sizeof(float)), "cudaMalloc d_mu");
	cudaErrorCheck(cudaMalloc((void **)&d_sigma,N*sizeof(float)), "cudaMalloc d_sigma");

	// Copy over matrix to device 
	cudaErrorCheck(cudaMemcpy(d_A,(const void
					*)A[0],size,cudaMemcpyHostToDevice), "cudaMemcpy A");

	//Initalize h_sum and h_dev to 0
	int row,col;
	for(row = 0; row < GridSize; row++)
	{
		h_mu[row] = 0.0;
		h_sigma[row] = 0.0;
		for(col = 0; col < N; col++)
		{
			h_Sum[row*N + col] = 0.0;
			h_Dev[row*N + col] = 0.0;
		}
	}

	//Initalize the device sum and std dev arrays to 0s.
	cudaErrorCheck(cudaMemcpy((void *)d_Sum,(const void
				*)h_Sum,GridSize*N*sizeof(float),cudaMemcpyHostToDevice),
			"cudaMemcpy to d_Sum");
	cudaErrorCheck(cudaMemcpy((void *)d_Dev,(const void
				*)h_Dev,GridSize*N*sizeof(float),cudaMemcpyHostToDevice),"cudaMemcpy to d_Dev" );
	cudaErrorCheck(cudaMemcpy((void *)d_mu,(const void
				*)h_mu,N*sizeof(float),cudaMemcpyHostToDevice), "cudaMemcpy to d_mu" );
	cudaErrorCheck(cudaMemcpy((void *)d_sigma,(const void
				*)h_sigma,N*sizeof(float),cudaMemcpyHostToDevice), "cudaMemcpy to d_sigma");


	//Calcuate a sub mean for each block
	BlockMean<<<grid,block>>>(d_A,d_Sum,N);
	cudaDeviceSynchronize();

	//Calculate total mean for each column sequentially 
	cudaErrorCheck(cudaMemcpy((void *)h_Sum,(const void
					*)d_Sum,GridSize*N*sizeof(float),cudaMemcpyDeviceToHost),
			"cudaMemcpy to h_Sum");

	for(row = 0; row < GridSize; row++)
	{
		for(col = 0; col < N; col++)
		{
			h_mu[col] += h_Sum[row*N + col];
		}
	} 
	for(col = 0; col < N; col++)
		h_mu[col] /= N; 

	//Copy over host calculated mu vector to host
	cudaErrorCheck(cudaMemcpy((void *)d_mu,(const void
				*)h_mu,N*sizeof(float),cudaMemcpyHostToDevice), "cudaMemcpy to d_mu");

	//Calculate a sub standard deviation for each block
	BlockDev<<<grid,block>>>(d_A,d_Dev,d_mu,N);
	cudaDeviceSynchronize();

	//Calculate total standard deviation from each block sequentially
	cudaErrorCheck(cudaMemcpy((void *)h_Dev,(const void
				*)d_Dev,GridSize*N*sizeof(float),cudaMemcpyDeviceToHost),
			"cudaMemcpy to h_dev");

	for(row = 0; row < GridSize; row++)
	{
		for(col = 0; col < N; col++)
		{
			h_sigma[col] += h_Dev[row*N + col];
		}
	}
	for(col = 0; col < N; col++)
	{
		h_sigma[col] /= N;
		h_sigma[col] = sqrt(h_sigma[col]);
	}

	//Copy sigma vector to device
	cudaErrorCheck(cudaMemcpy((void *)d_sigma,(const void
				*)h_sigma,N*sizeof(float),cudaMemcpyHostToDevice), "cudaMemcpy to d_sigma" );

	//Normalize with means and standard deviations by splitting into blocks
	Normalize<<<grid,block>>>(d_A,d_mu,d_sigma,N);
	cudaDeviceSynchronize();

	//Copy Normalized array back to B
	cudaErrorCheck(cudaMemcpy((void *)B[0],(const void
					*)d_A,size,cudaMemcpyDeviceToHost), "cudaMemcpy to B");

	//Free all host and device pointers
	cudaFree(d_A);
	cudaFree(d_Sum);
	cudaFree(d_Dev);
	cudaFree(d_mu);
	cudaFree(d_sigma);
	free(h_Sum);
	free(h_Dev);
	free(h_mu);
	free(h_sigma);
}

/* Calculates a partial sum for a section of the matrix for each block.
   This is done by allocating a shared sub matrix, calculated a
   sum for each column using the algorithm from class, and returning
   the corresponding sub sum from each block.
   d_Sum holds all of these partial sums for every block.
   */
__global__ void BlockMean(float *d_A,float *d_Sum, int N)
{
	//Shared sub matrix
	__shared__ float sum[BlockSize*BlockSize];

	//Block and thread indices
	unsigned int x = blockDim.x * blockIdx.x + threadIdx.x;
	unsigned int y = blockDim.y * blockIdx.y + threadIdx.y;
	unsigned int tx = threadIdx.x;
	unsigned int ty = threadIdx.y;

	//Size of entire matrix
	int size = N*N*sizeof(float);

	//Indices for flattened input matrices
	// i + j == [i][j]
	unsigned int i = blockIdx.x*BlockSize*N + tx*N;
	unsigned int j = blockIdx.y*BlockSize + ty;

	//Row index into sum
	unsigned int sx = tx*BlockSize;

	//Ensure block and thread within bounds of matrix
	if(x >= N || y >= N)
		return; 

	//Transfer section of d_A into sum
	if(i + j < size) 
		sum[sx + ty] = d_A[i + j];  
	else 
		sum[sx + ty] = 0.0;

	//Apply partial sum algorithm from class
	for(unsigned int stride = blockDim.x >> 1; stride > 0; stride >>= 1)
	{
		__syncthreads();
		if(tx < stride)
			sum[sx + ty] += sum[sx + ty + stride];
	}
	
	//Transfer shared sub sum matrix to global memory
	if(tx == 0)
	{
		d_Sum[blockIdx.x*N + ty] = sum[ty];
	}
}

/* Calculate a partial sum of the square of the difference between the mean 
   for each block, in the same way as BlockMean, except by squaring a 
   difference of a calculated mean for the column in d_mu.
   */
__global__ void BlockDev(float *d_A, float *d_Dev, float *d_mu, int N)
{
	//Shared mu vector
	__shared__ float mu[BlockSize];
	//shared partial sum sub matrix
	__shared__ float sum[BlockSize*BlockSize];
	//Size of entire matrix
	int size = N*N*sizeof(float);

	//Block and thread indices
	unsigned int x = blockDim.x * blockIdx.x + threadIdx.x;
	unsigned int y = blockDim.y * blockIdx.y + threadIdx.y;
	unsigned int tx = threadIdx.x;
	unsigned int ty = threadIdx.y;

	//Indices for flattened input matrices
	// i + j == [i][j]
	unsigned int i = blockIdx.x*BlockSize*N + tx*N;
	unsigned int j = blockIdx.y*BlockSize + ty;

	//Row index into sum
	unsigned int sx = tx*BlockSize;

	//Ensure block and thread within bounds of matrix
	if(x >= N || y >= N)
		return; 

	//Transfer sub mu vector into shared memory
	if(tx == 0)
	{
		mu[ty] = d_mu[j];
	}
	
	//Ensure mu vector populated
	__syncthreads();

	//Transfer section of d_A into sum and square the mean difference
	if((i + j < size) && (j < N))
		sum[sx + ty] = powf(d_A[i + j] - mu[ty],2.0);  
	else 
		sum[sx + ty] = 0.0;

	
	//Apply partial sum algorithm shown in class
	for(unsigned int stride = blockDim.x >> 1; stride > 0; stride >>= 1)
	{
		__syncthreads();
		if(tx < stride)
			sum[sx + ty] += sum[sx + ty + stride];
	}

	//Tranfer shared sub matrix to global memory
	if(tx == 0)
	{
		d_Dev[blockIdx.x*N + ty] = sum[ty];
	}
}
/* Normalizing function, which each thread among all blocks corresponds
   to a single element in the matrix. Each applies the normalizing function
   to its element*/
__global__ void Normalize(float *d_A, float *d_mu, float *d_sigma, int N)
{
	// Thread position variables
	unsigned int tx = blockDim.x * blockIdx.x + threadIdx.x;
	unsigned int ty = blockDim.y * blockIdx.y + threadIdx.y;

	//Ensure within bounds of matrix
	if(tx >= N || ty >= N)
		return;

	//If sigma 0, set 0, else calculate normalized value
	if(d_sigma[ty] == 0)
		d_A[tx + ty*N] = 0;
	else
		d_A[tx + ty*N] = (d_A[tx + ty*N] - d_mu[ty])/(d_sigma[ty]);
}

/*Simple printing error function */
void cudaErrorCheck(cudaError_t err, const char *s)
{
	if(err != cudaSuccess)
	{
		printf("%s error: %s\n",s,cudaGetErrorString(err));
		exit(0);
	}
}


	.file	"tmpxft_00250dc2_00000000-6_matrixNorm.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
	.section	.rodata.str1.1
.LC13:
	.string	"%s error: %s\n"
	.text
	.globl	_Z14cudaErrorCheck9cudaErrorPKc
	.type	_Z14cudaErrorCheck9cudaErrorPKc, @function
_Z14cudaErrorCheck9cudaErrorPKc:
.LFB2047:
	.cfi_startproc
	endbr64
	testl	%edi, %edi
	je	.L49
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movq	%rsi, %rbx
	call	cudaGetErrorString@PLT
	movl	$2, %edi
	movq	%rbx, %rdx
	leaq	.LC13(%rip), %rsi
	movq	%rax, %rcx
	xorl	%eax, %eax
	call	__printf_chk@PLT
	xorl	%edi, %edi
	call	exit@PLT
.L49:
	.cfi_def_cfa_offset 8
	.cfi_restore 3
	ret
	.cfi_endproc
.LFE2047:
	.size	_Z14cudaErrorCheck9cudaErrorPKc, .-_Z14cudaErrorCheck9cudaErrorPKc
	.globl	_Z31__device_stub__Z9BlockMeanPfS_iPfS_i
	.type	_Z31__device_stub__Z9BlockMeanPfS_iPfS_i, @function
_Z31__device_stub__Z9BlockMeanPfS_iPfS_i:
.LFB2072:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 24(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L54
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z9BlockMeanPfS_i(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 160
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L54:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L56
	call	__stack_chk_fail@PLT
.L56:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2072:
	.size	_Z31__device_stub__Z9BlockMeanPfS_iPfS_i, .-_Z31__device_stub__Z9BlockMeanPfS_iPfS_i
	.globl	_Z9BlockMeanPfS_i
	.type	_Z9BlockMeanPfS_i, @function
_Z9BlockMeanPfS_i:
.LFB2073:
	.cfi_startproc
	endbr64
	jmp	_Z31__device_stub__Z9BlockMeanPfS_iPfS_i
	.cfi_endproc
.LFE2073:
	.size	_Z9BlockMeanPfS_i, .-_Z9BlockMeanPfS_i
	.globl	_Z32__device_stub__Z8BlockDevPfS_S_iPfS_S_i
	.type	_Z32__device_stub__Z8BlockDevPfS_S_iPfS_S_i, @function
_Z32__device_stub__Z8BlockDevPfS_S_iPfS_S_i:
.LFB2074:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movl	%ecx, 4(%rsp)
	leaq	48(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 104(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L59
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z8BlockDevPfS_S_i(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 176
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	120(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L59:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L61
	call	__stack_chk_fail@PLT
.L61:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2074:
	.size	_Z32__device_stub__Z8BlockDevPfS_S_iPfS_S_i, .-_Z32__device_stub__Z8BlockDevPfS_S_iPfS_S_i
	.globl	_Z8BlockDevPfS_S_i
	.type	_Z8BlockDevPfS_S_i, @function
_Z8BlockDevPfS_S_i:
.LFB2075:
	.cfi_startproc
	endbr64
	jmp	_Z32__device_stub__Z8BlockDevPfS_S_iPfS_S_i
	.cfi_endproc
.LFE2075:
	.size	_Z8BlockDevPfS_S_i, .-_Z8BlockDevPfS_S_i
	.globl	_Z33__device_stub__Z9NormalizePfS_S_iPfS_S_i
	.type	_Z33__device_stub__Z9NormalizePfS_S_iPfS_S_i, @function
_Z33__device_stub__Z9NormalizePfS_S_iPfS_S_i:
.LFB2076:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movl	%ecx, 4(%rsp)
	leaq	48(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 104(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L64
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z9NormalizePfS_S_i(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 176
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	120(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L64:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L66
	call	__stack_chk_fail@PLT
.L66:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2076:
	.size	_Z33__device_stub__Z9NormalizePfS_S_iPfS_S_i, .-_Z33__device_stub__Z9NormalizePfS_S_iPfS_S_i
	.globl	_Z9NormalizePfS_S_i
	.type	_Z9NormalizePfS_S_i, @function
_Z9NormalizePfS_S_i:
.LFB2077:
	.cfi_startproc
	endbr64
	jmp	_Z33__device_stub__Z9NormalizePfS_S_iPfS_S_i
	.cfi_endproc
.LFE2077:
	.size	_Z9NormalizePfS_S_i, .-_Z9NormalizePfS_S_i
	.section	.rodata.str1.1
.LC14:
	.string	"Executing on GPU\n"
.LC16:
	.string	"cudaMalloc A"
.LC17:
	.string	"cudaMalloc d_Sum"
.LC18:
	.string	"cudaMalloc d_Dev "
.LC19:
	.string	"cudaMalloc d_mu"
.LC20:
	.string	"cudaMalloc d_sigma"
.LC21:
	.string	"cudaMemcpy A"
.LC22:
	.string	"cudaMemcpy to d_Sum"
.LC23:
	.string	"cudaMemcpy to d_Dev"
.LC24:
	.string	"cudaMemcpy to d_mu"
.LC25:
	.string	"cudaMemcpy to d_sigma"
.LC26:
	.string	"cudaMemcpy to h_Sum"
.LC27:
	.string	"cudaMemcpy to h_dev"
.LC28:
	.string	"cudaMemcpy to B"
	.text
	.globl	_Z10matrixNormv
	.type	_Z10matrixNormv, @function
_Z10matrixNormv:
.LFB2046:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	leaq	.LC14(%rip), %rsi
	movl	$2, %edi
	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	$104, %rsp
	.cfi_def_cfa_offset 160
	movq	%fs:40, %rax
	movq	%rax, 88(%rsp)
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	N(%rip), %r15d
	cvtsi2ssl	%r15d, %xmm0
	mulss	.LC15(%rip), %xmm0
	movl	%r15d, %ebp
	call	ceilf@PLT
	movl	$1, 72(%rsp)
	movabsq	$137438953504, %rax
	cvttss2sil	%xmm0, %ebx
	movq	%rax, 76(%rsp)
	movl	$1, 84(%rsp)
	imull	%ebx, %ebp
	movl	%ebx, 64(%rsp)
	movl	%ebx, 68(%rsp)
	movslq	%ebp, %rbp
	salq	$2, %rbp
	movq	%rbp, %rdi
	call	malloc@PLT
	movq	%rbp, %rdi
	movslq	%r15d, %rbp
	salq	$2, %rbp
	imull	%r15d, %r15d
	movq	%rax, %r14
	call	malloc@PLT
	movq	%rbp, %rdi
	movq	%rax, %r13
	call	malloc@PLT
	movq	%rbp, %rdi
	movslq	%r15d, %r15
	movq	%rax, %r12
	call	malloc@PLT
	leaq	24(%rsp), %rdi
	movq	%rax, %rbp
	leaq	0(,%r15,4), %rax
	movq	%rax, %rsi
	movq	%rax, 8(%rsp)
	call	cudaMalloc@PLT
	leaq	.LC16(%rip), %rsi
	movl	%eax, %edi
	call	_Z14cudaErrorCheck9cudaErrorPKc
	movl	N(%rip), %esi
	leaq	32(%rsp), %rdi
	imull	%ebx, %esi
	movslq	%esi, %rsi
	salq	$2, %rsi
	call	cudaMalloc@PLT
	leaq	.LC17(%rip), %rsi
	movl	%eax, %edi
	call	_Z14cudaErrorCheck9cudaErrorPKc
	movl	N(%rip), %esi
	leaq	40(%rsp), %rdi
	imull	%ebx, %esi
	movslq	%esi, %rsi
	salq	$2, %rsi
	call	cudaMalloc@PLT
	leaq	.LC18(%rip), %rsi
	movl	%eax, %edi
	call	_Z14cudaErrorCheck9cudaErrorPKc
	movslq	N(%rip), %rsi
	leaq	48(%rsp), %rdi
	salq	$2, %rsi
	call	cudaMalloc@PLT
	leaq	.LC19(%rip), %rsi
	movl	%eax, %edi
	call	_Z14cudaErrorCheck9cudaErrorPKc
	movslq	N(%rip), %rsi
	leaq	56(%rsp), %rdi
	salq	$2, %rsi
	call	cudaMalloc@PLT
	leaq	.LC20(%rip), %rsi
	movl	%eax, %edi
	call	_Z14cudaErrorCheck9cudaErrorPKc
	movq	8(%rsp), %rdx
	movq	24(%rsp), %rdi
	movl	$1, %ecx
	leaq	A(%rip), %rsi
	call	cudaMemcpy@PLT
	leaq	.LC21(%rip), %rsi
	movl	%eax, %edi
	call	_Z14cudaErrorCheck9cudaErrorPKc
	movl	N(%rip), %edx
	xorl	%esi, %esi
	xorl	%ecx, %ecx
.L70:
	cmpl	%ecx, %ebx
	jle	.L92
	movl	$0x00000000, (%r12,%rcx,4)
	movslq	%esi, %rax
	xorl	%edi, %edi
	movl	$0x00000000, 0(%rbp,%rcx,4)
	salq	$2, %rax
.L71:
	cmpl	%edi, %edx
	jle	.L93
	movl	$0x00000000, (%r14,%rax)
	incl	%edi
	movl	$0x00000000, 0(%r13,%rax)
	addq	$4, %rax
	jmp	.L71
.L93:
	incq	%rcx
	addl	%edx, %esi
	jmp	.L70
.L92:
	imull	%ebx, %edx
	movq	32(%rsp), %rdi
	movl	$1, %ecx
	movq	%r14, %rsi
	leaq	.LC24(%rip), %r15
	movslq	%edx, %rdx
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	leaq	.LC22(%rip), %rsi
	movl	%eax, %edi
	call	_Z14cudaErrorCheck9cudaErrorPKc
	movl	N(%rip), %edx
	movq	40(%rsp), %rdi
	movq	%r13, %rsi
	movl	$1, %ecx
	imull	%ebx, %edx
	movslq	%edx, %rdx
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	leaq	.LC23(%rip), %rsi
	movl	%eax, %edi
	call	_Z14cudaErrorCheck9cudaErrorPKc
	movslq	N(%rip), %rdx
	movq	48(%rsp), %rdi
	movq	%r12, %rsi
	movl	$1, %ecx
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	movq	%r15, %rsi
	movl	%eax, %edi
	call	_Z14cudaErrorCheck9cudaErrorPKc
	movslq	N(%rip), %rdx
	movq	56(%rsp), %rdi
	movq	%rbp, %rsi
	movl	$1, %ecx
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	leaq	.LC25(%rip), %rsi
	movl	%eax, %edi
	call	_Z14cudaErrorCheck9cudaErrorPKc
	movl	84(%rsp), %ecx
	movl	72(%rsp), %esi
	xorl	%r9d, %r9d
	movq	76(%rsp), %rdx
	movq	64(%rsp), %rdi
	xorl	%r8d, %r8d
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L74
	movl	N(%rip), %edx
	movq	32(%rsp), %rsi
	movq	24(%rsp), %rdi
	call	_Z31__device_stub__Z9BlockMeanPfS_iPfS_i
.L74:
	call	cudaDeviceSynchronize@PLT
	movl	N(%rip), %edx
	movq	32(%rsp), %rsi
	movq	%r14, %rdi
	movl	$2, %ecx
	imull	%ebx, %edx
	movslq	%edx, %rdx
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	leaq	.LC26(%rip), %rsi
	movl	%eax, %edi
	call	_Z14cudaErrorCheck9cudaErrorPKc
	movslq	N(%rip), %rdx
	xorl	%ecx, %ecx
	xorl	%esi, %esi
.L75:
	cmpl	%ebx, %esi
	jge	.L89
	movslq	%ecx, %rax
	leaq	(%r14,%rax,4), %rdi
	xorl	%eax, %eax
.L78:
	cmpl	%eax, %edx
	jle	.L94
	movss	(%r12,%rax,4), %xmm0
	addss	(%rdi,%rax,4), %xmm0
	movss	%xmm0, (%r12,%rax,4)
	incq	%rax
	jmp	.L78
.L94:
	incl	%esi
	addl	%edx, %ecx
	jmp	.L75
.L89:
	xorl	%eax, %eax
.L77:
	cmpl	%eax, %edx
	jle	.L95
	cvtsi2ssl	%edx, %xmm1
	movss	(%r12,%rax,4), %xmm0
	divss	%xmm1, %xmm0
	movss	%xmm0, (%r12,%rax,4)
	incq	%rax
	jmp	.L77
.L95:
	movq	48(%rsp), %rdi
	salq	$2, %rdx
	movl	$1, %ecx
	movq	%r12, %rsi
	call	cudaMemcpy@PLT
	movq	%r15, %rsi
	movl	%eax, %edi
	call	_Z14cudaErrorCheck9cudaErrorPKc
	movl	84(%rsp), %ecx
	movl	72(%rsp), %esi
	xorl	%r9d, %r9d
	movq	76(%rsp), %rdx
	movq	64(%rsp), %rdi
	xorl	%r8d, %r8d
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L80
	movl	N(%rip), %ecx
	movq	48(%rsp), %rdx
	movq	40(%rsp), %rsi
	movq	24(%rsp), %rdi
	call	_Z32__device_stub__Z8BlockDevPfS_S_iPfS_S_i
.L80:
	call	cudaDeviceSynchronize@PLT
	movl	N(%rip), %edx
	movq	40(%rsp), %rsi
	movq	%r13, %rdi
	movl	$2, %ecx
	imull	%ebx, %edx
	movslq	%edx, %rdx
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	leaq	.LC27(%rip), %rsi
	movl	%eax, %edi
	call	_Z14cudaErrorCheck9cudaErrorPKc
	movl	N(%rip), %r15d
	xorl	%edx, %edx
	xorl	%ecx, %ecx
.L81:
	cmpl	%ebx, %ecx
	jge	.L90
	movslq	%edx, %rax
	leaq	0(%r13,%rax,4), %rsi
	xorl	%eax, %eax
.L84:
	cmpl	%eax, %r15d
	jle	.L96
	movss	0(%rbp,%rax,4), %xmm0
	addss	(%rsi,%rax,4), %xmm0
	movss	%xmm0, 0(%rbp,%rax,4)
	incq	%rax
	jmp	.L84
.L96:
	incl	%ecx
	addl	%r15d, %edx
	jmp	.L81
.L90:
	xorl	%ebx, %ebx
.L83:
	cmpl	%ebx, %r15d
	jle	.L97
	cvtsi2ssl	%r15d, %xmm1
	movss	0(%rbp,%rbx,4), %xmm0
	divss	%xmm1, %xmm0
	call	sqrtf@PLT
	movss	%xmm0, 0(%rbp,%rbx,4)
	incq	%rbx
	jmp	.L83
.L97:
	movq	56(%rsp), %rdi
	movslq	%r15d, %rdx
	movl	$1, %ecx
	movq	%rbp, %rsi
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	leaq	.LC25(%rip), %rsi
	movl	%eax, %edi
	call	_Z14cudaErrorCheck9cudaErrorPKc
	movl	84(%rsp), %ecx
	movl	72(%rsp), %esi
	xorl	%r9d, %r9d
	movq	76(%rsp), %rdx
	movq	64(%rsp), %rdi
	xorl	%r8d, %r8d
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L86
	movl	N(%rip), %ecx
	movq	56(%rsp), %rdx
	movq	48(%rsp), %rsi
	movq	24(%rsp), %rdi
	call	_Z33__device_stub__Z9NormalizePfS_S_iPfS_S_i
.L86:
	call	cudaDeviceSynchronize@PLT
	movq	8(%rsp), %rdx
	movq	24(%rsp), %rsi
	movl	$2, %ecx
	leaq	B(%rip), %rdi
	call	cudaMemcpy@PLT
	leaq	.LC28(%rip), %rsi
	movl	%eax, %edi
	call	_Z14cudaErrorCheck9cudaErrorPKc
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
	movq	32(%rsp), %rdi
	call	cudaFree@PLT
	movq	40(%rsp), %rdi
	call	cudaFree@PLT
	movq	48(%rsp), %rdi
	call	cudaFree@PLT
	movq	56(%rsp), %rdi
	call	cudaFree@PLT
	movq	%r14, %rdi
	call	free@PLT
	movq	%r13, %rdi
	call	free@PLT
	movq	%r12, %rdi
	call	free@PLT
	movq	88(%rsp), %rax
	subq	%fs:40, %rax
	je	.L87
	call	__stack_chk_fail@PLT
.L87:
	addq	$104, %rsp
	.cfi_def_cfa_offset 56
	movq	%rbp, %rdi
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	jmp	free@PLT
	.cfi_endproc
.LFE2046:
	.size	_Z10matrixNormv, .-_Z10matrixNormv
	.section	.rodata.str1.1
.LC29:
	.string	"\nStarting clock.\n"
.LC30:
	.string	"Stopped clock.\n"
.LC32:
	.string	"\nElapsed time = %g ms.\n"
.LC34:
	.string	"(CPU times are accurate to the nearest %g ms)\n"
.LC36:
	.string	"My total CPU time for parent = %g ms.\n"
.LC37:
	.string	"My system CPU time for parent = %g ms.\n"
.LC38:
	.string	"My total CPU time for child processes = %g ms.\n"
.LC39:
	.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
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	subq	$112, %rsp
	.cfi_def_cfa_offset 144
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	movq	%rsp, %rbx
	call	_Z10parametersiPPc
	call	_Z17initialize_inputsv
	call	_Z12print_inputsv
	leaq	.LC29(%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	_Z10matrixNormv
	movq	%rbx, %rsi
	leaq	24(%rsp), %rdi
	call	gettimeofday@PLT
	leaq	72(%rsp), %rdi
	call	times@PLT
	leaq	.LC30(%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	.L99
	cvtsi2ssq	%rax, %xmm0
	jmp	.L100
.L99:
	movq	%rax, %rdx
	andl	$1, %eax
	shrq	%rdx
	orq	%rax, %rdx
	cvtsi2ssq	%rdx, %xmm0
	addss	%xmm0, %xmm0
.L100:
	leaq	.LC32(%rip), %rsi
	movl	$2, %edi
	movb	$1, %al
	divss	.LC31(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	call	__printf_chk@PLT
	movl	$2, %edi
	movb	$1, %al
	movsd	.LC33(%rip), %xmm0
	leaq	.LC34(%rip), %rsi
	call	__printf_chk@PLT
	movq	80(%rsp), %rax
	addq	72(%rsp), %rax
	leaq	.LC36(%rip), %rsi
	subq	48(%rsp), %rax
	subq	40(%rsp), %rax
	movl	$2, %edi
	cvtsi2ssq	%rax, %xmm0
	divss	.LC35(%rip), %xmm0
	movb	$1, %al
	mulss	.LC31(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	call	__printf_chk@PLT
	movq	80(%rsp), %rax
	subq	48(%rsp), %rax
	cvtsi2ssq	%rax, %xmm0
	divss	.LC35(%rip), %xmm0
	leaq	.LC37(%rip), %rsi
	movl	$2, %edi
	mulss	.LC31(%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	.LC35(%rip), %xmm0
	leaq	.LC38(%rip), %rsi
	movb	$1, %al
	mulss	.LC31(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	call	__printf_chk@PLT
	movl	$2, %edi
	xorl	%eax, %eax
	leaq	.LC39(%rip), %rsi
	call	__printf_chk@PLT
	xorl	%edi, %edi
	call	exit@PLT
	.cfi_endproc
.LFE2045:
	.size	main, .-main
	.section	.rodata.str1.1
.LC40:
	.string	"_Z9NormalizePfS_S_i"
.LC41:
	.string	"_Z8BlockDevPfS_S_i"
.LC42:
	.string	"_Z9BlockMeanPfS_i"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2079:
	.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	.LC40(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z9NormalizePfS_S_i(%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	.LC41(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z8BlockDevPfS_S_i(%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	.LC42(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z9BlockMeanPfS_i(%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
.LFE2079:
	.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
.LC15:
	.long	1023410176
	.align 4
.LC31:
	.long	1148846080
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC33:
	.long	-755914244
	.long	1062232653
	.section	.rodata.cst4
	.align 4
.LC35:
	.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();
void cudaErrorCheck(hipError_t err, const char *s);
__global__ void matrixCuda(float *d_A,int N);
__global__ void BlockMean(float *d_A,float *d_Sum, int N);
__global__ void BlockDev(float *d_A, float *d_Dev, float *d_mu,int N);
__global__ void Normalize(float *d_A, float *d_mu, float *d_sigma, int N);

/* 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;
		}
	}
	/*
	   for (col = 0; col < N; col++) {
	   for (row = 0; row < N; row++) {
	   A[row][col] = col + row;
	   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) {
	/* 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();

	/* 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.
 */


/* Matrix Normalization for the Cuda platform
   Overview
   1.Copy matrix A to Device
   2.CUDA: Split matrix into 2D grid of blocks, each block calculates a partial sum 
   for each column in their section via BlockMean.
   3.Sequentially add all of these partial sums to form the total mean for
   each column.
   4.CUDA: Each block calculate a partial sum of the mean difference for each
   column in their section via BlockDev.
   5.Sequentially add all of these partial mean differences to form the total
   standard deviation for each column.
   6.CUDA: Each block normalize their portion of the matrix using the means
   and standard deviations calculated.
   7.Copy matrix A to Host's B
   */

//Number of threads per block
#define BlockSize 32
void matrixNorm() 
{
	printf("Executing on GPU\n");

	//Set grid size to divide among number of threads
	int GridSize = ceil((float)N/BlockSize); 

	//Create CUDA grid and block size for matrix
	dim3 grid(GridSize,GridSize);
	dim3 block(BlockSize,BlockSize);

	//Device matrices and vectors for calculation 
	float *d_A; //the matrix 
	float *d_Sum; //a partial sum holder
	float *d_Dev; //a partial deviation holder
	float *d_mu; //a vector of means for each column
	float *d_sigma; //a vector of standard deviations for each column

	//Host copies of variables to initalize the device's copies to 0's
	float *h_Sum = (float *)malloc(GridSize*N*sizeof(float));
	float *h_Dev = (float *)malloc(GridSize*N*sizeof(float));
	float *h_mu = (float *)malloc(N*sizeof(float));
	float *h_sigma = (float *)malloc(N*sizeof(float));
	size_t size = N*N*sizeof(float);

	// Allocate Matrices on Devices 
	cudaErrorCheck(hipMalloc((void **)&d_A,size), "hipMalloc A");
	cudaErrorCheck(hipMalloc(&d_Sum,GridSize*N*sizeof(float)), "hipMalloc d_Sum" );
	cudaErrorCheck(hipMalloc(&d_Dev,GridSize*N*sizeof(float)), "hipMalloc d_Dev ");
	cudaErrorCheck(hipMalloc((void **)&d_mu,N*sizeof(float)), "hipMalloc d_mu");
	cudaErrorCheck(hipMalloc((void **)&d_sigma,N*sizeof(float)), "hipMalloc d_sigma");

	// Copy over matrix to device 
	cudaErrorCheck(hipMemcpy(d_A,(const void
					*)A[0],size,hipMemcpyHostToDevice), "hipMemcpy A");

	//Initalize h_sum and h_dev to 0
	int row,col;
	for(row = 0; row < GridSize; row++)
	{
		h_mu[row] = 0.0;
		h_sigma[row] = 0.0;
		for(col = 0; col < N; col++)
		{
			h_Sum[row*N + col] = 0.0;
			h_Dev[row*N + col] = 0.0;
		}
	}

	//Initalize the device sum and std dev arrays to 0s.
	cudaErrorCheck(hipMemcpy((void *)d_Sum,(const void
				*)h_Sum,GridSize*N*sizeof(float),hipMemcpyHostToDevice),
			"hipMemcpy to d_Sum");
	cudaErrorCheck(hipMemcpy((void *)d_Dev,(const void
				*)h_Dev,GridSize*N*sizeof(float),hipMemcpyHostToDevice),"hipMemcpy to d_Dev" );
	cudaErrorCheck(hipMemcpy((void *)d_mu,(const void
				*)h_mu,N*sizeof(float),hipMemcpyHostToDevice), "hipMemcpy to d_mu" );
	cudaErrorCheck(hipMemcpy((void *)d_sigma,(const void
				*)h_sigma,N*sizeof(float),hipMemcpyHostToDevice), "hipMemcpy to d_sigma");


	//Calcuate a sub mean for each block
	BlockMean<<<grid,block>>>(d_A,d_Sum,N);
	hipDeviceSynchronize();

	//Calculate total mean for each column sequentially 
	cudaErrorCheck(hipMemcpy((void *)h_Sum,(const void
					*)d_Sum,GridSize*N*sizeof(float),hipMemcpyDeviceToHost),
			"hipMemcpy to h_Sum");

	for(row = 0; row < GridSize; row++)
	{
		for(col = 0; col < N; col++)
		{
			h_mu[col] += h_Sum[row*N + col];
		}
	} 
	for(col = 0; col < N; col++)
		h_mu[col] /= N; 

	//Copy over host calculated mu vector to host
	cudaErrorCheck(hipMemcpy((void *)d_mu,(const void
				*)h_mu,N*sizeof(float),hipMemcpyHostToDevice), "hipMemcpy to d_mu");

	//Calculate a sub standard deviation for each block
	BlockDev<<<grid,block>>>(d_A,d_Dev,d_mu,N);
	hipDeviceSynchronize();

	//Calculate total standard deviation from each block sequentially
	cudaErrorCheck(hipMemcpy((void *)h_Dev,(const void
				*)d_Dev,GridSize*N*sizeof(float),hipMemcpyDeviceToHost),
			"hipMemcpy to h_dev");

	for(row = 0; row < GridSize; row++)
	{
		for(col = 0; col < N; col++)
		{
			h_sigma[col] += h_Dev[row*N + col];
		}
	}
	for(col = 0; col < N; col++)
	{
		h_sigma[col] /= N;
		h_sigma[col] = sqrt(h_sigma[col]);
	}

	//Copy sigma vector to device
	cudaErrorCheck(hipMemcpy((void *)d_sigma,(const void
				*)h_sigma,N*sizeof(float),hipMemcpyHostToDevice), "hipMemcpy to d_sigma" );

	//Normalize with means and standard deviations by splitting into blocks
	Normalize<<<grid,block>>>(d_A,d_mu,d_sigma,N);
	hipDeviceSynchronize();

	//Copy Normalized array back to B
	cudaErrorCheck(hipMemcpy((void *)B[0],(const void
					*)d_A,size,hipMemcpyDeviceToHost), "hipMemcpy to B");

	//Free all host and device pointers
	hipFree(d_A);
	hipFree(d_Sum);
	hipFree(d_Dev);
	hipFree(d_mu);
	hipFree(d_sigma);
	free(h_Sum);
	free(h_Dev);
	free(h_mu);
	free(h_sigma);
}

/* Calculates a partial sum for a section of the matrix for each block.
   This is done by allocating a shared sub matrix, calculated a
   sum for each column using the algorithm from class, and returning
   the corresponding sub sum from each block.
   d_Sum holds all of these partial sums for every block.
   */
__global__ void BlockMean(float *d_A,float *d_Sum, int N)
{
	//Shared sub matrix
	__shared__ float sum[BlockSize*BlockSize];

	//Block and thread indices
	unsigned int x = blockDim.x * blockIdx.x + threadIdx.x;
	unsigned int y = blockDim.y * blockIdx.y + threadIdx.y;
	unsigned int tx = threadIdx.x;
	unsigned int ty = threadIdx.y;

	//Size of entire matrix
	int size = N*N*sizeof(float);

	//Indices for flattened input matrices
	// i + j == [i][j]
	unsigned int i = blockIdx.x*BlockSize*N + tx*N;
	unsigned int j = blockIdx.y*BlockSize + ty;

	//Row index into sum
	unsigned int sx = tx*BlockSize;

	//Ensure block and thread within bounds of matrix
	if(x >= N || y >= N)
		return; 

	//Transfer section of d_A into sum
	if(i + j < size) 
		sum[sx + ty] = d_A[i + j];  
	else 
		sum[sx + ty] = 0.0;

	//Apply partial sum algorithm from class
	for(unsigned int stride = blockDim.x >> 1; stride > 0; stride >>= 1)
	{
		__syncthreads();
		if(tx < stride)
			sum[sx + ty] += sum[sx + ty + stride];
	}
	
	//Transfer shared sub sum matrix to global memory
	if(tx == 0)
	{
		d_Sum[blockIdx.x*N + ty] = sum[ty];
	}
}

/* Calculate a partial sum of the square of the difference between the mean 
   for each block, in the same way as BlockMean, except by squaring a 
   difference of a calculated mean for the column in d_mu.
   */
__global__ void BlockDev(float *d_A, float *d_Dev, float *d_mu, int N)
{
	//Shared mu vector
	__shared__ float mu[BlockSize];
	//shared partial sum sub matrix
	__shared__ float sum[BlockSize*BlockSize];
	//Size of entire matrix
	int size = N*N*sizeof(float);

	//Block and thread indices
	unsigned int x = blockDim.x * blockIdx.x + threadIdx.x;
	unsigned int y = blockDim.y * blockIdx.y + threadIdx.y;
	unsigned int tx = threadIdx.x;
	unsigned int ty = threadIdx.y;

	//Indices for flattened input matrices
	// i + j == [i][j]
	unsigned int i = blockIdx.x*BlockSize*N + tx*N;
	unsigned int j = blockIdx.y*BlockSize + ty;

	//Row index into sum
	unsigned int sx = tx*BlockSize;

	//Ensure block and thread within bounds of matrix
	if(x >= N || y >= N)
		return; 

	//Transfer sub mu vector into shared memory
	if(tx == 0)
	{
		mu[ty] = d_mu[j];
	}
	
	//Ensure mu vector populated
	__syncthreads();

	//Transfer section of d_A into sum and square the mean difference
	if((i + j < size) && (j < N))
		sum[sx + ty] = powf(d_A[i + j] - mu[ty],2.0);  
	else 
		sum[sx + ty] = 0.0;

	
	//Apply partial sum algorithm shown in class
	for(unsigned int stride = blockDim.x >> 1; stride > 0; stride >>= 1)
	{
		__syncthreads();
		if(tx < stride)
			sum[sx + ty] += sum[sx + ty + stride];
	}

	//Tranfer shared sub matrix to global memory
	if(tx == 0)
	{
		d_Dev[blockIdx.x*N + ty] = sum[ty];
	}
}
/* Normalizing function, which each thread among all blocks corresponds
   to a single element in the matrix. Each applies the normalizing function
   to its element*/
__global__ void Normalize(float *d_A, float *d_mu, float *d_sigma, int N)
{
	// Thread position variables
	unsigned int tx = blockDim.x * blockIdx.x + threadIdx.x;
	unsigned int ty = blockDim.y * blockIdx.y + threadIdx.y;

	//Ensure within bounds of matrix
	if(tx >= N || ty >= N)
		return;

	//If sigma 0, set 0, else calculate normalized value
	if(d_sigma[ty] == 0)
		d_A[tx + ty*N] = 0;
	else
		d_A[tx + ty*N] = (d_A[tx + ty*N] - d_mu[ty])/(d_sigma[ty]);
}

/*Simple printing error function */
void cudaErrorCheck(hipError_t err, const char *s)
{
	if(err != hipSuccess)
	{
		printf("%s error: %s\n",s,hipGetErrorString(err));
		exit(0);
	}
}


	.text
	.file	"matrixNorm.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	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	subq	$112, %rsp
	.cfi_def_cfa_offset 144
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	callq	_Z10parametersiPPc
	callq	_Z17initialize_inputsv
	callq	_Z12print_inputsv
	movl	$.Lstr.1, %edi
	callq	puts@PLT
	leaq	96(%rsp), %rbx
	leaq	72(%rsp), %r14
	movq	%rbx, %rdi
	movq	%r14, %rsi
	callq	gettimeofday
	leaq	40(%rsp), %rdi
	callq	times
	callq	_Z10matrixNormv
	leaq	80(%rsp), %r15
	movq	%r15, %rdi
	movq	%r14, %rsi
	callq	gettimeofday
	leaq	8(%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_1
# %bb.2:
	cvtsi2ss	%rax, %xmm0
	jmp	.LBB5_3
.LBB5_1:
	movq	%rax, %rcx
	shrq	%rcx
	andl	$1, %eax
	orq	%rcx, %rax
	cvtsi2ss	%rax, %xmm0
	addss	%xmm0, %xmm0
.LBB5_3:
	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	16(%rsp), %rbx
	movq	48(%rsp), %r14
	movq	8(%rsp), %rax
	addq	%rbx, %rax
	movq	40(%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	32(%rsp), %rax
	movq	56(%rsp), %rcx
	addq	24(%rsp), %rax
	addq	64(%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
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function _Z10matrixNormv
.LCPI6_0:
	.long	0x3d000000                      # float 0.03125
	.text
	.globl	_Z10matrixNormv
	.type	_Z10matrixNormv,@function
_Z10matrixNormv:                        # @_Z10matrixNormv
	.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	$104, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	$.Lstr.4, %edi
	callq	puts@PLT
	movslq	N(%rip), %rbx
	cvtsi2ss	%ebx, %xmm0
	mulss	.LCPI6_0(%rip), %xmm0
	callq	ceilf@PLT
	cvttss2si	%xmm0, %ebp
	movq	%rbp, %rax
	shlq	$32, %rax
	orq	%rbp, %rax
	movq	%rax, 48(%rsp)                  # 8-byte Spill
	movl	%ebx, %eax
	imull	%ebp, %eax
	movslq	%eax, %r14
	shlq	$2, %r14
	movq	%r14, %rdi
	callq	malloc
	movq	%rax, 72(%rsp)                  # 8-byte Spill
	movq	%r14, %rdi
	callq	malloc
	movq	%rax, 64(%rsp)                  # 8-byte Spill
	movq	%rbx, %r12
	shlq	$2, %r12
	movq	%r12, %rdi
	callq	malloc
	movq	%rax, %r15
	movq	%r12, %rdi
	callq	malloc
	movq	%rax, %r12
	imull	%ebx, %ebx
	shlq	$2, %rbx
	leaq	16(%rsp), %r14
	movq	%r14, %rdi
	movq	%rbx, %rsi
	callq	hipMalloc
	movl	$.L.str.20, %esi
	movl	%eax, %edi
	callq	_Z14cudaErrorCheck10hipError_tPKc
	movslq	N(%rip), %rsi
	movslq	%ebp, %r13
	imulq	%r13, %rsi
	shlq	$2, %rsi
	leaq	40(%rsp), %rdi
	callq	hipMalloc
	movl	$.L.str.21, %esi
	movl	%eax, %edi
	callq	_Z14cudaErrorCheck10hipError_tPKc
	movslq	N(%rip), %rsi
	imulq	%r13, %rsi
	shlq	$2, %rsi
	leaq	32(%rsp), %rdi
	callq	hipMalloc
	movl	$.L.str.22, %esi
	movl	%eax, %edi
	callq	_Z14cudaErrorCheck10hipError_tPKc
	movslq	N(%rip), %rsi
	shlq	$2, %rsi
	leaq	8(%rsp), %rdi
	callq	hipMalloc
	movl	$.L.str.23, %esi
	movl	%eax, %edi
	callq	_Z14cudaErrorCheck10hipError_tPKc
	movslq	N(%rip), %rsi
	shlq	$2, %rsi
	leaq	24(%rsp), %rdi
	callq	hipMalloc
	movl	$.L.str.24, %esi
	movl	%eax, %edi
	callq	_Z14cudaErrorCheck10hipError_tPKc
	movq	(%r14), %rdi
	movl	$A, %esi
	movq	%rbx, 80(%rsp)                  # 8-byte Spill
	movq	%rbx, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movl	$.L.str.25, %esi
	movl	%eax, %edi
	callq	_Z14cudaErrorCheck10hipError_tPKc
	movq	%rbp, %rax
	movl	N(%rip), %edx
	testl	%r13d, %r13d
	movq	%rbp, 56(%rsp)                  # 8-byte Spill
	jle	.LBB6_5
# %bb.1:                                # %.lr.ph123
	leaq	(,%rdx,4), %rcx
	movq	%rcx, 88(%rsp)                  # 8-byte Spill
	xorl	%ebx, %ebx
	xorl	%ebp, %ebp
	movq	88(%rsp), %r14                  # 8-byte Reload
.LBB6_2:                                # =>This Inner Loop Header: Depth=1
	xorl	%ecx, %ecx
	movl	%ecx, (%r15,%rbp,4)
	movl	%ecx, (%r12,%rbp,4)
	testl	%edx, %edx
	jle	.LBB6_4
# %bb.3:                                # %.lr.ph
                                        #   in Loop: Header=BB6_2 Depth=1
	movl	%ebx, %eax
	movq	72(%rsp), %rcx                  # 8-byte Reload
	leaq	(%rcx,%rax,4), %rdi
	movq	64(%rsp), %rcx                  # 8-byte Reload
	leaq	(%rcx,%rax,4), %rax
	movq	%rax, 96(%rsp)                  # 8-byte Spill
	xorl	%esi, %esi
	movq	%rdx, %r13
	movq	%r14, %rdx
	callq	memset@PLT
	movq	96(%rsp), %rdi                  # 8-byte Reload
	xorl	%esi, %esi
	movq	%r14, %rdx
	callq	memset@PLT
	movq	%r13, %rdx
	movq	56(%rsp), %rax                  # 8-byte Reload
.LBB6_4:                                # %._crit_edge
                                        #   in Loop: Header=BB6_2 Depth=1
	incq	%rbp
	addl	%edx, %ebx
	cmpq	%rbp, %rax
	jne	.LBB6_2
.LBB6_5:                                # %._crit_edge124
	movabsq	$137438953504, %rbx             # imm = 0x2000000020
	movq	40(%rsp), %rdi
	imull	%eax, %edx
	movslq	%edx, %rdx
	shlq	$2, %rdx
	movq	72(%rsp), %r14                  # 8-byte Reload
	movq	%r14, %rsi
	movl	$1, %ecx
	movq	%rax, %rbp
	callq	hipMemcpy
	movl	$.L.str.26, %esi
	movl	%eax, %edi
	callq	_Z14cudaErrorCheck10hipError_tPKc
	movq	32(%rsp), %rdi
	movslq	N(%rip), %rax
	movslq	%ebp, %rdx
	imulq	%rax, %rdx
	shlq	$2, %rdx
	movq	64(%rsp), %r13                  # 8-byte Reload
	movq	%r13, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movl	$.L.str.27, %esi
	movl	%eax, %edi
	callq	_Z14cudaErrorCheck10hipError_tPKc
	movq	8(%rsp), %rdi
	movslq	N(%rip), %rdx
	shlq	$2, %rdx
	movq	%r15, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movl	$.L.str.28, %esi
	movl	%eax, %edi
	callq	_Z14cudaErrorCheck10hipError_tPKc
	movq	24(%rsp), %rdi
	movslq	N(%rip), %rdx
	shlq	$2, %rdx
	movq	%r12, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movl	$.L.str.29, %esi
	movl	%eax, %edi
	callq	_Z14cudaErrorCheck10hipError_tPKc
	movq	48(%rsp), %rdi                  # 8-byte Reload
	movl	$1, %esi
	movq	%rbx, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB6_7
# %bb.6:
	movq	16(%rsp), %rdi
	movq	40(%rsp), %rsi
	movl	N(%rip), %edx
	callq	_Z24__device_stub__BlockMeanPfS_i
.LBB6_7:
	callq	hipDeviceSynchronize
	movq	40(%rsp), %rsi
	movslq	N(%rip), %rdx
	movslq	%ebp, %rbx
	imulq	%rbx, %rdx
	shlq	$2, %rdx
	movq	%r14, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movl	$.L.str.30, %esi
	movl	%eax, %edi
	callq	_Z14cudaErrorCheck10hipError_tPKc
	movl	N(%rip), %eax
	testl	%ebx, %ebx
	jle	.LBB6_13
# %bb.8:                                # %.preheader118.lr.ph
	xorl	%ecx, %ecx
	xorl	%edx, %edx
.LBB6_9:                                # %.preheader118
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB6_11 Depth 2
	testl	%eax, %eax
	jle	.LBB6_12
# %bb.10:                               # %.lr.ph126
                                        #   in Loop: Header=BB6_9 Depth=1
	movl	%ecx, %esi
	leaq	(%r14,%rsi,4), %rsi
	xorl	%edi, %edi
.LBB6_11:                               #   Parent Loop BB6_9 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movss	(%rsi,%rdi,4), %xmm0            # xmm0 = mem[0],zero,zero,zero
	addss	(%r15,%rdi,4), %xmm0
	movss	%xmm0, (%r15,%rdi,4)
	incq	%rdi
	cmpq	%rdi, %rax
	jne	.LBB6_11
.LBB6_12:                               # %._crit_edge127
                                        #   in Loop: Header=BB6_9 Depth=1
	incq	%rdx
	addl	%eax, %ecx
	cmpq	%rbp, %rdx
	jne	.LBB6_9
.LBB6_13:                               # %.preheader117
	movslq	%eax, %rdx
	testl	%edx, %edx
	jle	.LBB6_16
# %bb.14:                               # %.lr.ph130
	xorps	%xmm0, %xmm0
	cvtsi2ss	%eax, %xmm0
	xorl	%ecx, %ecx
.LBB6_15:                               # =>This Inner Loop Header: Depth=1
	movss	(%r15,%rcx,4), %xmm1            # xmm1 = mem[0],zero,zero,zero
	divss	%xmm0, %xmm1
	movss	%xmm1, (%r15,%rcx,4)
	incq	%rcx
	cmpq	%rcx, %rax
	jne	.LBB6_15
.LBB6_16:                               # %._crit_edge131
	movq	8(%rsp), %rdi
	shlq	$2, %rdx
	movq	%r15, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movl	$.L.str.28, %esi
	movl	%eax, %edi
	callq	_Z14cudaErrorCheck10hipError_tPKc
	movq	48(%rsp), %rdi                  # 8-byte Reload
	movl	$1, %esi
	movabsq	$137438953504, %rdx             # imm = 0x2000000020
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB6_18
# %bb.17:
	movq	16(%rsp), %rdi
	movq	32(%rsp), %rsi
	movq	8(%rsp), %rdx
	movl	N(%rip), %ecx
	callq	_Z23__device_stub__BlockDevPfS_S_i
.LBB6_18:
	callq	hipDeviceSynchronize
	movq	32(%rsp), %rsi
	movslq	N(%rip), %rdx
	movq	56(%rsp), %rbp                  # 8-byte Reload
	movslq	%ebp, %rbx
	imulq	%rbx, %rdx
	shlq	$2, %rdx
	movq	%r13, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movl	$.L.str.31, %esi
	movl	%eax, %edi
	callq	_Z14cudaErrorCheck10hipError_tPKc
	movl	N(%rip), %eax
	testl	%ebx, %ebx
	jle	.LBB6_24
# %bb.19:                               # %.preheader116.lr.ph
	movq	%rbp, %r8
	xorl	%ecx, %ecx
	xorl	%edx, %edx
.LBB6_20:                               # %.preheader116
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB6_22 Depth 2
	testl	%eax, %eax
	jle	.LBB6_23
# %bb.21:                               # %.lr.ph133
                                        #   in Loop: Header=BB6_20 Depth=1
	movl	%ecx, %esi
	leaq	(,%rsi,4), %rsi
	addq	%r13, %rsi
	xorl	%edi, %edi
.LBB6_22:                               #   Parent Loop BB6_20 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movss	(%rsi,%rdi,4), %xmm0            # xmm0 = mem[0],zero,zero,zero
	addss	(%r12,%rdi,4), %xmm0
	movss	%xmm0, (%r12,%rdi,4)
	incq	%rdi
	cmpq	%rdi, %rax
	jne	.LBB6_22
.LBB6_23:                               # %._crit_edge134
                                        #   in Loop: Header=BB6_20 Depth=1
	incq	%rdx
	addl	%eax, %ecx
	cmpq	%r8, %rdx
	jne	.LBB6_20
.LBB6_24:                               # %.preheader
	testl	%eax, %eax
	jle	.LBB6_25
# %bb.26:                               # %.lr.ph137.preheader
	xorl	%ebx, %ebx
	xorps	%xmm2, %xmm2
.LBB6_27:                               # %.lr.ph137
                                        # =>This Inner Loop Header: Depth=1
	xorps	%xmm1, %xmm1
	cvtsi2ss	%eax, %xmm1
	movss	(%r12,%rbx,4), %xmm0            # xmm0 = mem[0],zero,zero,zero
	divss	%xmm1, %xmm0
	ucomiss	%xmm2, %xmm0
	jb	.LBB6_29
# %bb.28:                               #   in Loop: Header=BB6_27 Depth=1
	sqrtss	%xmm0, %xmm0
	jmp	.LBB6_30
.LBB6_29:                               # %call.sqrt
                                        #   in Loop: Header=BB6_27 Depth=1
	callq	sqrtf
	xorps	%xmm2, %xmm2
.LBB6_30:                               # %.lr.ph137.split
                                        #   in Loop: Header=BB6_27 Depth=1
	movss	%xmm0, (%r12,%rbx,4)
	incq	%rbx
	movl	N(%rip), %eax
	movslq	%eax, %rdx
	cmpq	%rdx, %rbx
	jl	.LBB6_27
	jmp	.LBB6_31
.LBB6_25:
	movslq	%eax, %rdx
.LBB6_31:                               # %._crit_edge138
	movq	24(%rsp), %rdi
	shlq	$2, %rdx
	movq	%r12, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movl	$.L.str.29, %esi
	movl	%eax, %edi
	callq	_Z14cudaErrorCheck10hipError_tPKc
	movq	48(%rsp), %rdi                  # 8-byte Reload
	movl	$1, %esi
	movabsq	$137438953504, %rdx             # imm = 0x2000000020
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB6_33
# %bb.32:
	movq	16(%rsp), %rdi
	movq	8(%rsp), %rsi
	movq	24(%rsp), %rdx
	movl	N(%rip), %ecx
	callq	_Z24__device_stub__NormalizePfS_S_i
.LBB6_33:
	callq	hipDeviceSynchronize
	movq	16(%rsp), %rsi
	movl	$B, %edi
	movq	80(%rsp), %rdx                  # 8-byte Reload
	movl	$2, %ecx
	callq	hipMemcpy
	movl	$.L.str.32, %esi
	movl	%eax, %edi
	callq	_Z14cudaErrorCheck10hipError_tPKc
	movq	16(%rsp), %rdi
	callq	hipFree
	movq	40(%rsp), %rdi
	callq	hipFree
	movq	32(%rsp), %rdi
	callq	hipFree
	movq	8(%rsp), %rdi
	callq	hipFree
	movq	24(%rsp), %rdi
	callq	hipFree
	movq	%r14, %rdi
	callq	free
	movq	%r13, %rdi
	callq	free
	movq	%r15, %rdi
	callq	free
	movq	%r12, %rdi
	callq	free
	addq	$104, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end6:
	.size	_Z10matrixNormv, .Lfunc_end6-_Z10matrixNormv
	.cfi_endproc
                                        # -- End function
	.globl	_Z14cudaErrorCheck10hipError_tPKc # -- Begin function _Z14cudaErrorCheck10hipError_tPKc
	.type	_Z14cudaErrorCheck10hipError_tPKc,@function
_Z14cudaErrorCheck10hipError_tPKc:      # @_Z14cudaErrorCheck10hipError_tPKc
	.cfi_startproc
# %bb.0:
	testl	%edi, %edi
	jne	.LBB7_2
# %bb.1:
	retq
.LBB7_2:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset %rbx, -16
	movq	%rsi, %rbx
	callq	hipGetErrorString
	movl	$.L.str.33, %edi
	movq	%rbx, %rsi
	movq	%rax, %rdx
	xorl	%eax, %eax
	callq	printf
	xorl	%edi, %edi
	callq	exit
.Lfunc_end7:
	.size	_Z14cudaErrorCheck10hipError_tPKc, .Lfunc_end7-_Z14cudaErrorCheck10hipError_tPKc
	.cfi_endproc
                                        # -- End function
	.globl	_Z24__device_stub__BlockMeanPfS_i # -- Begin function _Z24__device_stub__BlockMeanPfS_i
	.type	_Z24__device_stub__BlockMeanPfS_i,@function
_Z24__device_stub__BlockMeanPfS_i:      # @_Z24__device_stub__BlockMeanPfS_i
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$112, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	12(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	movq	%rsi, 16(%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	$_Z9BlockMeanPfS_i, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end8:
	.size	_Z24__device_stub__BlockMeanPfS_i, .Lfunc_end8-_Z24__device_stub__BlockMeanPfS_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z23__device_stub__BlockDevPfS_S_i # -- Begin function _Z23__device_stub__BlockDevPfS_S_i
	.type	_Z23__device_stub__BlockDevPfS_S_i,@function
_Z23__device_stub__BlockDevPfS_S_i:     # @_Z23__device_stub__BlockDevPfS_S_i
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$112, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	24(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	4(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z8BlockDevPfS_S_i, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end9:
	.size	_Z23__device_stub__BlockDevPfS_S_i, .Lfunc_end9-_Z23__device_stub__BlockDevPfS_S_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z24__device_stub__NormalizePfS_S_i # -- Begin function _Z24__device_stub__NormalizePfS_S_i
	.type	_Z24__device_stub__NormalizePfS_S_i,@function
_Z24__device_stub__NormalizePfS_S_i:    # @_Z24__device_stub__NormalizePfS_S_i
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$112, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	24(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	4(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z9NormalizePfS_S_i, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end10:
	.size	_Z24__device_stub__NormalizePfS_S_i, .Lfunc_end10-_Z24__device_stub__NormalizePfS_S_i
	.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	.LBB11_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB11_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z9BlockMeanPfS_i, %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	$_Z8BlockDevPfS_S_i, %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	$_Z9NormalizePfS_S_i, %esi
	movl	$.L__unnamed_3, %edx
	movl	$.L__unnamed_3, %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_end11:
	.size	__hip_module_ctor, .Lfunc_end11-__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	.LBB12_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
.LBB12_2:
	retq
.Lfunc_end12:
	.size	__hip_module_dtor, .Lfunc_end12-__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	.L.str.13,@object               # @.str.13
.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	.L.str.20,@object               # @.str.20
.L.str.20:
	.asciz	"hipMalloc A"
	.size	.L.str.20, 12

	.type	.L.str.21,@object               # @.str.21
.L.str.21:
	.asciz	"hipMalloc d_Sum"
	.size	.L.str.21, 16

	.type	.L.str.22,@object               # @.str.22
.L.str.22:
	.asciz	"hipMalloc d_Dev "
	.size	.L.str.22, 17

	.type	.L.str.23,@object               # @.str.23
.L.str.23:
	.asciz	"hipMalloc d_mu"
	.size	.L.str.23, 15

	.type	.L.str.24,@object               # @.str.24
.L.str.24:
	.asciz	"hipMalloc d_sigma"
	.size	.L.str.24, 18

	.type	.L.str.25,@object               # @.str.25
.L.str.25:
	.asciz	"hipMemcpy A"
	.size	.L.str.25, 12

	.type	.L.str.26,@object               # @.str.26
.L.str.26:
	.asciz	"hipMemcpy to d_Sum"
	.size	.L.str.26, 19

	.type	.L.str.27,@object               # @.str.27
.L.str.27:
	.asciz	"hipMemcpy to d_Dev"
	.size	.L.str.27, 19

	.type	.L.str.28,@object               # @.str.28
.L.str.28:
	.asciz	"hipMemcpy to d_mu"
	.size	.L.str.28, 18

	.type	.L.str.29,@object               # @.str.29
.L.str.29:
	.asciz	"hipMemcpy to d_sigma"
	.size	.L.str.29, 21

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

	.type	.L.str.30,@object               # @.str.30
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.30:
	.asciz	"hipMemcpy to h_Sum"
	.size	.L.str.30, 19

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

	.type	.L.str.31,@object               # @.str.31
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.31:
	.asciz	"hipMemcpy to h_dev"
	.size	.L.str.31, 19

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

	.type	.L.str.32,@object               # @.str.32
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.32:
	.asciz	"hipMemcpy to B"
	.size	.L.str.32, 15

	.type	.L.str.33,@object               # @.str.33
.L.str.33:
	.asciz	"%s error: %s\n"
	.size	.L.str.33, 14

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z9BlockMeanPfS_i"
	.size	.L__unnamed_1, 18

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z8BlockDevPfS_S_i"
	.size	.L__unnamed_2, 19

	.type	.L__unnamed_3,@object           # @2
.L__unnamed_3:
	.asciz	"_Z9NormalizePfS_S_i"
	.size	.L__unnamed_3, 20

	.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

	.type	.Lstr.4,@object                 # @str.4
.Lstr.4:
	.asciz	"Executing on GPU"
	.size	.Lstr.4, 17

	.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 _Z24__device_stub__BlockMeanPfS_i
	.addrsig_sym _Z23__device_stub__BlockDevPfS_S_i
	.addrsig_sym _Z24__device_stub__NormalizePfS_S_i
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym A
	.addrsig_sym B
	.addrsig_sym _Z9BlockMeanPfS_i
	.addrsig_sym _Z8BlockDevPfS_S_i
	.addrsig_sym _Z9NormalizePfS_S_i
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z9BlockMeanPfS_i       ; -- Begin function _Z9BlockMeanPfS_i
	.globl	_Z9BlockMeanPfS_i
	.p2align	8
	.type	_Z9BlockMeanPfS_i,@function
_Z9BlockMeanPfS_i:                      ; @_Z9BlockMeanPfS_i
; %bb.0:
	s_clause 0x1
	s_load_b32 s5, s[0:1], 0x24
	s_load_b32 s4, s[0:1], 0x10
	v_and_b32_e32 v1, 0x3ff, v0
	s_add_u32 s2, s0, 24
	s_addc_u32 s3, s1, 0
	s_mov_b32 s6, exec_lo
	s_waitcnt lgkmcnt(0)
	s_and_b32 s5, s5, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[2:3], null, s14, s5, v[1:2]
	v_cmpx_gt_u32_e64 s4, v2
	s_cbranch_execz .LBB0_10
; %bb.1:
	s_load_b32 s2, s[2:3], 0xc
	v_bfe_u32 v0, v0, 10, 10
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s2, s2, 16
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_mad_u64_u32 v[2:3], null, s15, s2, v[0:1]
	s_delay_alu instid0(VALU_DEP_1)
	v_cmp_gt_u32_e32 vcc_lo, s4, v2
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_10
; %bb.2:
	v_lshl_add_u32 v2, s14, 5, v1
	s_load_b128 s[0:3], s[0:1], 0x0
	s_lshl_b32 s6, s15, 5
	v_mov_b32_e32 v3, 0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v2, v2, s4
	v_add3_u32 v2, s6, v0, v2
	s_mul_i32 s6, s4, s4
	s_delay_alu instid0(SALU_CYCLE_1)
	s_lshl_b32 s6, s6, 2
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_cmp_gt_u32_e32 vcc_lo, s6, v2
	s_and_saveexec_b32 s6, vcc_lo
	s_cbranch_execz .LBB0_4
; %bb.3:
	v_mov_b32_e32 v3, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[2:3], 2, v[2:3]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v2, vcc_lo, s0, v2
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v3, vcc_lo, s1, v3, vcc_lo
	global_load_b32 v3, v[2:3], off
.LBB0_4:
	s_or_b32 exec_lo, exec_lo, s6
	v_lshlrev_b32_e32 v2, 5, v1
	s_cmp_lt_u32 s5, 2
	s_delay_alu instid0(VALU_DEP_1)
	v_add_lshl_u32 v2, v2, v0, 2
	s_waitcnt vmcnt(0)
	ds_store_b32 v2, v3
	s_cbranch_scc1 .LBB0_8
.LBB0_5:                                ; =>This Inner Loop Header: Depth=1
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s0, s5, 1
	s_mov_b32 s1, exec_lo
	s_barrier
	buffer_gl0_inv
	v_cmpx_gt_u32_e64 s0, v1
	s_cbranch_execz .LBB0_7
; %bb.6:                                ;   in Loop: Header=BB0_5 Depth=1
	v_lshl_add_u32 v3, s0, 2, v2
	ds_load_b32 v3, v3
	ds_load_b32 v4, v2
	s_waitcnt lgkmcnt(0)
	v_add_f32_e32 v3, v3, v4
	ds_store_b32 v2, v3
.LBB0_7:                                ;   in Loop: Header=BB0_5 Depth=1
	s_or_b32 exec_lo, exec_lo, s1
	s_cmp_lt_u32 s5, 4
	s_mov_b32 s5, s0
	s_cbranch_scc0 .LBB0_5
.LBB0_8:                                ; %._crit_edge
	v_cmp_eq_u32_e32 vcc_lo, 0, v1
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_10
; %bb.9:
	v_lshlrev_b32_e32 v3, 2, v0
	v_mad_u64_u32 v[1:2], null, s14, s4, v[0:1]
	v_mov_b32_e32 v2, 0
	ds_load_b32 v3, v3
	v_lshlrev_b64 v[0:1], 2, v[1:2]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v0, vcc_lo, s2, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s3, v1, vcc_lo
	global_store_b32 v[0:1], v3, off
.LBB0_10:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z9BlockMeanPfS_i
		.amdhsa_group_segment_fixed_size 4096
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 280
		.amdhsa_user_sgpr_count 14
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 1
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 1
		.amdhsa_next_free_vgpr 5
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end0:
	.size	_Z9BlockMeanPfS_i, .Lfunc_end0-_Z9BlockMeanPfS_i
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 460
; NumSgprs: 18
; NumVgprs: 5
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 4096 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 5
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 14
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 1
	.text
	.protected	_Z8BlockDevPfS_S_i      ; -- Begin function _Z8BlockDevPfS_S_i
	.globl	_Z8BlockDevPfS_S_i
	.p2align	8
	.type	_Z8BlockDevPfS_S_i,@function
_Z8BlockDevPfS_S_i:                     ; @_Z8BlockDevPfS_S_i
; %bb.0:
	s_clause 0x1
	s_load_b32 s4, s[0:1], 0x2c
	s_load_b32 s8, s[0:1], 0x18
	v_and_b32_e32 v1, 0x3ff, v0
	s_add_u32 s2, s0, 32
	s_addc_u32 s3, s1, 0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s9, s4, 0xffff
	s_mov_b32 s4, exec_lo
	v_mad_u64_u32 v[2:3], null, s14, s9, v[1:2]
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_gt_u32_e64 s8, v2
	s_cbranch_execz .LBB1_12
; %bb.1:
	s_load_b32 s2, s[2:3], 0xc
	v_bfe_u32 v0, v0, 10, 10
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s2, s2, 16
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_mad_u64_u32 v[2:3], null, s15, s2, v[0:1]
	s_delay_alu instid0(VALU_DEP_1)
	v_cmp_gt_u32_e32 vcc_lo, s8, v2
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB1_12
; %bb.2:
	s_clause 0x1
	s_load_b128 s[4:7], s[0:1], 0x0
	s_load_b64 s[0:1], s[0:1], 0x10
	v_lshl_add_u32 v2, s15, 5, v0
	v_cmp_eq_u32_e32 vcc_lo, 0, v1
	v_lshlrev_b32_e32 v5, 2, v0
	s_and_saveexec_b32 s2, vcc_lo
	s_cbranch_execz .LBB1_4
; %bb.3:
	v_mov_b32_e32 v3, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[3:4], 2, v[2:3]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v3, s0, s0, v3
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e64 v4, s0, s1, v4, s0
	global_load_b32 v3, v[3:4], off
	s_waitcnt vmcnt(0)
	ds_store_b32 v5, v3 offset:4096
.LBB1_4:
	s_or_b32 exec_lo, exec_lo, s2
	v_lshl_add_u32 v6, s14, 5, v1
	s_waitcnt lgkmcnt(0)
	s_mul_i32 s0, s8, s8
	v_cmp_gt_u32_e64 s1, s8, v2
	s_lshl_b32 s0, s0, 2
	v_mad_u64_u32 v[3:4], null, v6, s8, v[2:3]
	v_mov_b32_e32 v4, 0
	s_barrier
	buffer_gl0_inv
	v_cmp_gt_u32_e64 s0, s0, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s0, s1, s0
	s_and_saveexec_b32 s3, s0
	s_cbranch_execz .LBB1_6
; %bb.5:
	v_mov_b32_e32 v4, 0
	s_mov_b32 s1, 0x3e76c4e1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[2:3], 2, v[3:4]
	v_add_co_u32 v2, s0, s4, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_4) | instid1(VALU_DEP_1)
	v_add_co_ci_u32_e64 v3, s0, s5, v3, s0
	global_load_b32 v2, v[2:3], off
	ds_load_b32 v3, v5 offset:4096
	s_waitcnt vmcnt(0) lgkmcnt(0)
	v_sub_f32_e32 v2, v2, v3
	v_frexp_mant_f32_e64 v3, |v2|
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_gt_f32_e64 s0, 0x3f2aaaab, v3
	v_cndmask_b32_e64 v4, 0, 1, s0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ldexp_f32 v3, v3, v4
	v_add_f32_e32 v4, 1.0, v3
	v_add_f32_e32 v7, -1.0, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v9, -1.0, v4
	v_sub_f32_e32 v3, v3, v9
	v_rcp_f32_e32 v6, v4
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v8, v7, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v10, v4, v8
	v_fma_f32 v4, v8, v4, -v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v4, v8, v3
	v_add_f32_e32 v3, v10, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_sub_f32_e32 v9, v7, v3
	v_sub_f32_e32 v10, v3, v10
	v_dual_sub_f32 v7, v7, v9 :: v_dual_sub_f32 v4, v10, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_f32_e32 v3, v7, v3
	v_add_f32_e32 v3, v4, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v3, v9, v3
	v_mul_f32_e32 v3, v6, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v4, v8, v3
	v_sub_f32_e32 v6, v4, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_f32_e32 v3, v3, v6
	v_dual_mul_f32 v7, v4, v4 :: v_dual_add_f32 v8, v3, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v6, v4, v4, -v7
	v_fmac_f32_e32 v6, v4, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v8, v7, v6
	v_sub_f32_e32 v7, v8, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_fmaak_f32 v9, s1, v8, 0x3e91f4c4 :: v_dual_sub_f32 v6, v6, v7
	v_fmaak_f32 v9, v8, v9, 0x3ecccdef
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v10, v8, v9
	v_fma_f32 v7, v8, v9, -v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_fmac_f32 v7, v6, v9 :: v_dual_mul_f32 v12, v4, v8
	v_add_f32_e32 v9, v10, v7
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v13, v8, v4, -v12
	v_dual_sub_f32 v10, v9, v10 :: v_dual_fmac_f32 v13, v8, v3
	v_ldexp_f32 v3, v3, 1
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v13, v6, v4
	v_ldexp_f32 v4, v4, 1
	v_add_f32_e32 v8, v12, v13
	v_add_f32_e32 v11, 0x3f2aaaaa, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_sub_f32 v7, v7, v10 :: v_dual_add_f32 v10, 0xbf2aaaaa, v11
	v_add_f32_e32 v7, 0x31739010, v7
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_f32_e32 v9, v9, v10
	v_add_f32_e32 v7, v7, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v6, v11, v7
	v_mul_f32_e32 v10, v8, v6
	v_sub_f32_e32 v9, v11, v6
	v_sub_f32_e32 v11, v8, v12
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_f32_e32 v7, v7, v9
	v_fma_f32 v9, v8, v6, -v10
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_f32_e32 v11, v13, v11
	v_fmac_f32_e32 v9, v8, v7
	v_frexp_exp_i32_f32_e32 v7, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v9, v11, v6
	v_subrev_co_ci_u32_e64 v6, s0, 0, v7, s0
	v_cmp_eq_f32_e64 s0, 1.0, v2
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f32_e32 v7, v10, v9
	v_cvt_f32_i32_e32 v6, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_f32_e32 v8, v4, v7
	v_sub_f32_e32 v10, v7, v10
	v_dual_mul_f32 v11, 0x3f317218, v6 :: v_dual_sub_f32 v4, v8, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_f32_e32 v9, v9, v10
	v_fma_f32 v10, 0x3f317218, v6, -v11
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_f32_e32 v4, v7, v4
	v_dual_add_f32 v3, v3, v9 :: v_dual_fmamk_f32 v6, v6, 0xb102e308, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v3, v3, v4
	v_dual_add_f32 v4, v11, v6 :: v_dual_add_f32 v7, v8, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_f32_e32 v11, v4, v11
	v_dual_add_f32 v9, v4, v7 :: v_dual_sub_f32 v8, v7, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_f32_e32 v6, v6, v11
	v_sub_f32_e32 v10, v9, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_sub_f32 v3, v3, v8 :: v_dual_sub_f32 v12, v9, v10
	v_dual_sub_f32 v7, v7, v10 :: v_dual_add_f32 v8, v6, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_f32_e32 v4, v4, v12
	v_dual_add_f32 v4, v7, v4 :: v_dual_sub_f32 v7, v8, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f32_e32 v4, v8, v4
	v_sub_f32_e32 v8, v8, v7
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_dual_sub_f32 v3, v3, v7 :: v_dual_add_f32 v10, v9, v4
	v_sub_f32_e32 v6, v6, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_add_f32_e32 v3, v3, v6
	v_cndmask_b32_e64 v6, 2.0, 1.0, s0
	v_sub_f32_e32 v7, v10, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_f32_e32 v4, v4, v7
	v_add_f32_e32 v3, v3, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v4, v10, v3
	v_dual_sub_f32 v7, v4, v10 :: v_dual_mul_f32 v8, v6, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_f32_e32 v3, v3, v7
	v_fma_f32 v4, v6, v4, -v8
	v_cmp_class_f32_e64 s0, v8, 0x204
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v4, v6, v3
	v_add_f32_e32 v3, v8, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e64 v7, v3, v8, s0
	v_cmp_eq_f32_e64 s0, 0x42b17218, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e64 v9, 0, 0x37000000, s0
	v_cmp_neq_f32_e64 s0, 0x7f800000, |v7|
	v_sub_f32_e32 v10, v7, v9
	v_sub_f32_e32 v3, v3, v8
	v_trunc_f32_e32 v7, v6
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_f32_e32 v11, 0x3fb8aa3b, v10
	v_sub_f32_e32 v3, v4, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_fma_f32 v12, 0x3fb8aa3b, v10, -v11
	v_rndne_f32_e32 v13, v11
	v_cndmask_b32_e64 v3, 0, v3, s0
	v_cmp_ngt_f32_e64 s0, 0xc2ce8ed0, v10
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_dual_fmamk_f32 v12, v10, 0x32a5705f, v12 :: v_dual_sub_f32 v11, v11, v13
	v_cvt_i32_f32_e32 v8, v13
	v_add_f32_e32 v3, v9, v3
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v11, v11, v12
	v_exp_f32_e32 v11, v11
	s_waitcnt_depctr 0xfff
	v_ldexp_f32 v4, v11, v8
	v_mul_f32_e32 v8, 0.5, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_cndmask_b32_e64 v4, 0, v4, s0
	v_cmp_nlt_f32_e64 s0, 0x42b17218, v10
	v_trunc_f32_e32 v11, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_cndmask_b32_e64 v4, 0x7f800000, v4, s0
	v_cmp_eq_f32_e64 s0, v7, v6
	v_cmp_neq_f32_e64 s1, v11, v8
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_fma_f32 v3, v4, v3, v4
	v_cmp_eq_f32_e64 s2, 0x7f800000, v4
	s_and_b32 s1, s0, s1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_cndmask_b32_e64 v6, 1.0, v2, s1
	v_cndmask_b32_e64 v7, 0, v2, s1
	v_cndmask_b32_e64 v3, v3, v4, s2
	v_cmp_eq_f32_e64 s2, 0, v2
	v_cmp_class_f32_e64 s1, v2, 0x204
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_bfi_b32 v3, 0x7fffffff, v3, v6
	v_cndmask_b32_e64 v4, 0x7f800000, 0, s2
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_cndmask_b32_e64 v6, 0x7fc00000, v3, s0
	v_cmp_gt_f32_e64 s0, 0, v2
	v_bfi_b32 v4, 0x7fffffff, v4, v7
	s_delay_alu instid0(VALU_DEP_2)
	v_cndmask_b32_e64 v3, v3, v6, s0
	s_or_b32 s0, s2, s1
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_cndmask_b32_e64 v3, v3, v4, s0
	v_cmp_o_f32_e64 s0, v2, v2
	s_delay_alu instid0(VALU_DEP_1)
	v_cndmask_b32_e64 v4, 0x7fc00000, v3, s0
.LBB1_6:
	s_or_b32 exec_lo, exec_lo, s3
	v_lshlrev_b32_e32 v2, 5, v1
	s_cmp_lt_u32 s9, 2
	s_delay_alu instid0(VALU_DEP_1)
	v_add_lshl_u32 v2, v2, v0, 2
	ds_store_b32 v2, v4
	s_cbranch_scc1 .LBB1_10
.LBB1_7:                                ; =>This Inner Loop Header: Depth=1
	s_lshr_b32 s1, s9, 1
	s_mov_b32 s2, exec_lo
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	v_cmpx_gt_u32_e64 s1, v1
	s_cbranch_execz .LBB1_9
; %bb.8:                                ;   in Loop: Header=BB1_7 Depth=1
	v_lshl_add_u32 v3, s1, 2, v2
	ds_load_b32 v3, v3
	ds_load_b32 v4, v2
	s_waitcnt lgkmcnt(0)
	v_add_f32_e32 v3, v3, v4
	ds_store_b32 v2, v3
.LBB1_9:                                ;   in Loop: Header=BB1_7 Depth=1
	s_or_b32 exec_lo, exec_lo, s2
	s_cmp_lt_u32 s9, 4
	s_mov_b32 s9, s1
	s_cbranch_scc0 .LBB1_7
.LBB1_10:                               ; %._crit_edge
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB1_12
; %bb.11:
	v_mad_u64_u32 v[1:2], null, s14, s8, v[0:1]
	ds_load_b32 v3, v5
	v_mov_b32_e32 v2, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 2, v[1:2]
	v_add_co_u32 v0, vcc_lo, s6, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v1, vcc_lo, s7, v1, vcc_lo
	s_waitcnt lgkmcnt(0)
	global_store_b32 v[0:1], v3, off
.LBB1_12:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z8BlockDevPfS_S_i
		.amdhsa_group_segment_fixed_size 4224
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 288
		.amdhsa_user_sgpr_count 14
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 1
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 1
		.amdhsa_next_free_vgpr 14
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end1:
	.size	_Z8BlockDevPfS_S_i, .Lfunc_end1-_Z8BlockDevPfS_S_i
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 1712
; NumSgprs: 18
; NumVgprs: 14
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 4224 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 14
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 14
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 1
	.text
	.protected	_Z9NormalizePfS_S_i     ; -- Begin function _Z9NormalizePfS_S_i
	.globl	_Z9NormalizePfS_S_i
	.p2align	8
	.type	_Z9NormalizePfS_S_i,@function
_Z9NormalizePfS_S_i:                    ; @_Z9NormalizePfS_S_i
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x2c
	s_load_b32 s4, s[0:1], 0x18
	v_and_b32_e32 v1, 0x3ff, v0
	v_bfe_u32 v4, v0, 10, 10
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s3, s2, 16
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[2:3], null, s14, s2, v[1:2]
	v_mad_u64_u32 v[0:1], null, s15, s3, v[4:5]
	s_mov_b32 s2, exec_lo
	v_max_u32_e32 v1, v2, v0
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_gt_u32_e64 s4, v1
	s_cbranch_execz .LBB2_4
; %bb.1:
	s_load_b64 s[2:3], s[0:1], 0x10
	v_mov_b32_e32 v1, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[4:5], 2, v[0:1]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v6, vcc_lo, s2, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_4) | instid1(VALU_DEP_1)
	v_add_co_ci_u32_e32 v7, vcc_lo, s3, v5, vcc_lo
	s_load_b128 s[0:3], s[0:1], 0x0
	global_load_b32 v6, v[6:7], off
	v_mad_u64_u32 v[7:8], null, v0, s4, v[2:3]
	v_mov_b32_e32 v8, v1
	v_lshlrev_b64 v[2:3], 2, v[7:8]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_co_u32 v2, s0, s0, v2
	v_add_co_ci_u32_e64 v3, s0, s1, v3, s0
	s_mov_b32 s0, exec_lo
	s_waitcnt vmcnt(0)
	v_cmpx_neq_f32_e32 0, v6
	s_cbranch_execz .LBB2_3
; %bb.2:
	v_add_co_u32 v0, vcc_lo, s2, v4
	v_add_co_ci_u32_e32 v1, vcc_lo, s3, v5, vcc_lo
	global_load_b32 v4, v[2:3], off
	global_load_b32 v0, v[0:1], off
	s_waitcnt vmcnt(0)
	v_sub_f32_e32 v0, v4, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_scale_f32 v1, null, v6, v6, v0
	v_rcp_f32_e32 v4, v1
	s_waitcnt_depctr 0xfff
	v_fma_f32 v5, -v1, v4, 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v4, v5, v4
	v_div_scale_f32 v5, vcc_lo, v0, v6, v0
	v_mul_f32_e32 v7, v5, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v8, -v1, v7, v5
	v_fmac_f32_e32 v7, v8, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v1, -v1, v7, v5
	v_div_fmas_f32 v1, v1, v4, v7
	s_delay_alu instid0(VALU_DEP_1)
	v_div_fixup_f32 v1, v1, v6, v0
.LBB2_3:                                ; %.sink.split
	s_or_b32 exec_lo, exec_lo, s0
	global_store_b32 v[2:3], v1, off
.LBB2_4:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z9NormalizePfS_S_i
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 288
		.amdhsa_user_sgpr_count 14
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 1
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 1
		.amdhsa_next_free_vgpr 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_end2:
	.size	_Z9NormalizePfS_S_i, .Lfunc_end2-_Z9NormalizePfS_S_i
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 372
; 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: 1
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .offset:         16
        .size:           4
        .value_kind:     by_value
      - .offset:         24
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         28
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         36
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         38
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         40
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         42
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         44
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         46
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         88
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 4096
    .kernarg_segment_align: 8
    .kernarg_segment_size: 280
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z9BlockMeanPfS_i
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z9BlockMeanPfS_i.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     5
    .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
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         36
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         44
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         46
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         48
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         50
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         52
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         54
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         96
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 4224
    .kernarg_segment_align: 8
    .kernarg_segment_size: 288
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z8BlockDevPfS_S_i
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z8BlockDevPfS_S_i.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     14
    .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
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         36
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         44
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         46
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         48
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         50
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         52
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         54
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         96
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 288
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z9NormalizePfS_S_i
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z9NormalizePfS_S_i.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
