/*! \file
*  \brief Custom GPU kernels used in the SAP algorithm
*
*/

#include <stdio.h>
#include <stdlib.h>
#include <cmath>
#include "kernels.cuh"
using namespace std;


__global__ void readVec(float * d_vector){
	/**
	Reads some number (specified by the number of threads at execution)
	of entries from an array of floats on the device. Useful for debugging.

	@param d_vector The name of the array of floats on the device
	which the kernel prints entries from. 
	*/
	int idx = blockIdx.x*blockDim.x + threadIdx.x;
	printf("The entry at %d is %f\n",idx,d_vector[idx]);
}

__global__ void readVecInt(int * d_vector){
	/**
	Reads some number (specified by the number of threads at execution)
	of entries from an array of floats on the device. Useful for debugging.

	@param d_vector The name of the array of floats on the device
	which the kernel prints entries from. 
	*/
	int idx = blockIdx.x*blockDim.x + threadIdx.x;
	printf("The entry at %d is %i\n",idx,d_vector[idx]);
}

// Calculate the secants for a collection of vectors stored as columns in a matrix, then normalize
__global__ void calculate_secants(float * d_secants, float * dpoints_in, int * d_int_constants){
	/** 
	Calculates the normalized secant set for a set of points.

	@param d_secants The secant set for d_points. Secants are stored as column
	vectors so the dimension of this matrix is (dim. of input data x number of
	secants). This matrix is the output of this kernel.
	@param d_points The input points stored as column vectors. The dimension
	of this matrix is (dim. of input data x number of points in data set)..
	@param d_int_constants An integer array which holds the input dimension 
	(the second entry in the array) and number of points (the second 
	entry in the array)

	*/

	int idx = blockIdx.x*blockDim.x + threadIdx.x;
	// Number of points
	int n = d_int_constants[1];
	// Input dimension
	int input_dim = d_int_constants[0];
	// Parameters used to pair points to calculate secants
	int i = idx % (n-1);
	int j = (idx - i)/(n-1);
	// Number of points for given pair
	int pair1;
	int pair2;
	if (i >= j){
		pair1 = i+1;
		pair2 = j;
	}else{
		pair1 = n-i-1;
		pair2 = n-j-1;
	}
	// For loop calculates secant coordinate by coordinate
	for (int p = 0; p < input_dim; p++){
		d_secants[idx*input_dim + p] = dpoints_in[input_dim*pair1 + p] - dpoints_in[input_dim*pair2 + p];
	}
	// Variable to store the norm of the secant
	float norm = 0;
	// Iterate through entries of the secant to calculate its norm
	for (int p = 0; p < input_dim; p++){
		norm = norm + powf(d_secants[idx*input_dim + p],2);
	}
	norm = sqrtf(norm);
	// As long as the norm is not zero, normalize the secant
	if (norm != 0.0){
		for (int p = 0; p < input_dim; p++){
			d_secants[idx*input_dim + p] = (1/norm)*d_secants[idx*input_dim + p];
		}
	}
}

// Take a matrix and return a vector whose entries are the l2 norms of the matrix
__global__ void calculate_col_norms(float * d_matrix, float * d_col_norms, int * d_int_constants){
	/** 
	Takes a matrix and returns a row vector whose entries are the L2 norms of each
	entry in the matrix.

	@param d_matrix The matrix whose column norms are being calculated 
	@param d_col_norms The row vector which stores the calculated column norms
	@param d_int_constants An integer array that stores the dimensions of the matrix. The third
	entry is the length of each column.

	*/

	int idx = blockIdx.x*blockDim.x + threadIdx.x;
	// Get the length of columns
	int proj_dim = d_int_constants[2];
	// Summation variable to store column norms
	float sum = 0;
	// Iterate down the columns, summing the squares of the entries
	for (int i = 0; i < proj_dim; i++){
		sum = sum + powf(d_matrix[idx*proj_dim + i],2);
	}
	// Take square root of resulting sum
	d_col_norms[idx] = sqrtf(sum);
}

__global__ void switch_columns(float * d_matrix, int * d_column_switch_indices){
	/** 
	Takes a matrix and switches two specified columns.

	@param d_matrix The matrix whose columns are being switched.
	@param d_column_switch_indices The indices of the columns to be switched.

	*/

	int idx = blockIdx.x*blockDim.x + threadIdx.x;
	// Initialize float to hold values from was column while the other is 
	// copied over
	float temp;
	// Number of rows
	int rows = d_column_switch_indices[0];
	// Index of first column to be swapped
	int index1 = d_column_switch_indices[1];
	// Index of second column to be swapped
	int index2 = d_column_switch_indices[2];
	// Save first column as temp
	temp = d_matrix[index1*rows + idx];
	// Copy second column into first index position
	d_matrix[index1*rows + idx] =  d_matrix[index2*rows + idx];
	// Copy original first column into second index position
	d_matrix[index2*rows + idx] = temp;
}

__global__ void shift_first_column(float * d_new_matrix, float * d_old_matrix, float * d_new_column, float * d_algo_constants){
	/** 
	Takes as input a matrix and shifted the first column by replacing it by
	a convex combination of the first column and a new column.

	@param d_new_matrix The output matrix with shifted column. 
	@param d_old_matrix The input matrix whose first column is being shifted.
	@param d_new_column The column vector that will be used to shift the first
	column of d_matrix.
	@param d_algo_constants A float array used to store the parameter controling the 
	convex combination.

	*/

	int idx = blockIdx.x*blockDim.x + threadIdx.x;
	// The parameter that controls how much the new column will be shifted by
	float alpha = d_algo_constants[0];
	// Shift entries of the first column of the matrix
	d_new_matrix[idx] = (1.0-alpha)*d_old_matrix[idx] + alpha*(d_new_column[idx] - d_old_matrix[idx]);
}

__global__ void normalize_first_column(float * d_matrix, int * d_int_constants){
	/** 
	Kernel that normalizes the first column vector of a matrix.
	
	@param d_matrix The matrix whose first column will be normalized. 
	@param d_int_constants An integer array whose first entry is the number of rows
	of the matrix.

	*/
	
	// Variable to hold sum of squares of entries of the first column
	float sum = 0;
	// Variable to hold the height of the column to be normalized.
	int column_height = d_int_constants[0];
	// Sum squares of entries of the first column
	for (int i = 0; i < column_height; i++){
		sum = sum + powf(d_matrix[i],2);
	}
	sum = sqrtf(sum);
	// As long as norm is not zero, divide all entries in first column by 
	// the first column norm
	if (sum != 0.0){
		for (int i = 0; i < column_height; i++){
			d_matrix[i] = (1/sum)*d_matrix[i];
		}
	}	
}

__global__ void make_identity(float * d_matrix, int * d_int_constants){
	/** 
	Kernel that takes a matrix and sets it as the identity matrix.
	
	@param d_matrix The matrix which will be set as the identity matrix.. 
	@param d_int_constants An integer array that holds the dimenison of d_matrix.

	*/
	
	int idx = blockIdx.x*blockDim.x + threadIdx.x;
	// The column number of the matrix
	int col = idx/d_int_constants[2];
	// The row number of the matrix
	int row = idx - col*d_int_constants[2];
	// If the column/row number correspond to a diagonal element
	// then set existing entry to 1, otherwise set it to 0.
	if (row == col){
		d_matrix[idx] = 1.0;
	}else{
		d_matrix[idx] = 0.0;
	}
}










	.file	"tmpxft_002a43fe_00000000-6_kernels.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2029:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2029:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z26__device_stub__Z7readVecPfPf
	.type	_Z26__device_stub__Z7readVecPfPf, @function
_Z26__device_stub__Z7readVecPfPf:
.LFB2051:
	.cfi_startproc
	endbr64
	subq	$104, %rsp
	.cfi_def_cfa_offset 112
	movq	%rdi, 8(%rsp)
	leaq	24(%rsp), %rcx
	leaq	16(%rsp), %rdx
	leaq	44(%rsp), %rsi
	leaq	32(%rsp), %rdi
	movq	%fs:40, %rax
	movq	%rax, 88(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 40(%rsp)
	movq	%rax, 80(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 32(%rsp)
	movq	%rax, 44(%rsp)
	movl	$1, 52(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	24(%rsp)
	.cfi_def_cfa_offset 120
	leaq	_Z7readVecPf(%rip), %rdi
	pushq	24(%rsp)
	.cfi_def_cfa_offset 128
	movq	60(%rsp), %rcx
	movl	68(%rsp), %r8d
	movq	48(%rsp), %rsi
	movl	56(%rsp), %edx
	leaq	96(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 120
	popq	%rdx
	.cfi_def_cfa_offset 112
.L2:
	movq	88(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$104, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2051:
	.size	_Z26__device_stub__Z7readVecPfPf, .-_Z26__device_stub__Z7readVecPfPf
	.globl	_Z7readVecPf
	.type	_Z7readVecPf, @function
_Z7readVecPf:
.LFB2052:
	.cfi_startproc
	endbr64
	jmp	_Z26__device_stub__Z7readVecPfPf
	.cfi_endproc
.LFE2052:
	.size	_Z7readVecPf, .-_Z7readVecPf
	.globl	_Z30__device_stub__Z10readVecIntPiPi
	.type	_Z30__device_stub__Z10readVecIntPiPi, @function
_Z30__device_stub__Z10readVecIntPiPi:
.LFB2053:
	.cfi_startproc
	endbr64
	subq	$104, %rsp
	.cfi_def_cfa_offset 112
	movq	%rdi, 8(%rsp)
	leaq	24(%rsp), %rcx
	leaq	16(%rsp), %rdx
	leaq	44(%rsp), %rsi
	leaq	32(%rsp), %rdi
	movq	%fs:40, %rax
	movq	%rax, 88(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 40(%rsp)
	movq	%rax, 80(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 32(%rsp)
	movq	%rax, 44(%rsp)
	movl	$1, 52(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L8
	pushq	24(%rsp)
	.cfi_def_cfa_offset 120
	leaq	_Z10readVecIntPi(%rip), %rdi
	pushq	24(%rsp)
	.cfi_def_cfa_offset 128
	movq	60(%rsp), %rcx
	movl	68(%rsp), %r8d
	movq	48(%rsp), %rsi
	movl	56(%rsp), %edx
	leaq	96(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 120
	popq	%rdx
	.cfi_def_cfa_offset 112
.L8:
	movq	88(%rsp), %rax
	subq	%fs:40, %rax
	je	.L10
	call	__stack_chk_fail@PLT
.L10:
	addq	$104, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2053:
	.size	_Z30__device_stub__Z10readVecIntPiPi, .-_Z30__device_stub__Z10readVecIntPiPi
	.globl	_Z10readVecIntPi
	.type	_Z10readVecIntPi, @function
_Z10readVecIntPi:
.LFB2054:
	.cfi_startproc
	endbr64
	jmp	_Z30__device_stub__Z10readVecIntPiPi
	.cfi_endproc
.LFE2054:
	.size	_Z10readVecIntPi, .-_Z10readVecIntPi
	.globl	_Z41__device_stub__Z17calculate_secantsPfS_PiPfS_Pi
	.type	_Z41__device_stub__Z17calculate_secantsPfS_PiPfS_Pi, @function
_Z41__device_stub__Z17calculate_secantsPfS_PiPfS_Pi:
.LFB2055:
	.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
	movq	%rdx, 8(%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	8(%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	.L13
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z17calculate_secantsPfS_Pi(%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
.L13:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L15
	call	__stack_chk_fail@PLT
.L15:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2055:
	.size	_Z41__device_stub__Z17calculate_secantsPfS_PiPfS_Pi, .-_Z41__device_stub__Z17calculate_secantsPfS_PiPfS_Pi
	.globl	_Z17calculate_secantsPfS_Pi
	.type	_Z17calculate_secantsPfS_Pi, @function
_Z17calculate_secantsPfS_Pi:
.LFB2056:
	.cfi_startproc
	endbr64
	jmp	_Z41__device_stub__Z17calculate_secantsPfS_PiPfS_Pi
	.cfi_endproc
.LFE2056:
	.size	_Z17calculate_secantsPfS_Pi, .-_Z17calculate_secantsPfS_Pi
	.globl	_Z43__device_stub__Z19calculate_col_normsPfS_PiPfS_Pi
	.type	_Z43__device_stub__Z19calculate_col_normsPfS_PiPfS_Pi, @function
_Z43__device_stub__Z19calculate_col_normsPfS_PiPfS_Pi:
.LFB2057:
	.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
	movq	%rdx, 8(%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	8(%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	.L18
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z19calculate_col_normsPfS_Pi(%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
.L18:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L20
	call	__stack_chk_fail@PLT
.L20:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2057:
	.size	_Z43__device_stub__Z19calculate_col_normsPfS_PiPfS_Pi, .-_Z43__device_stub__Z19calculate_col_normsPfS_PiPfS_Pi
	.globl	_Z19calculate_col_normsPfS_Pi
	.type	_Z19calculate_col_normsPfS_Pi, @function
_Z19calculate_col_normsPfS_Pi:
.LFB2058:
	.cfi_startproc
	endbr64
	jmp	_Z43__device_stub__Z19calculate_col_normsPfS_PiPfS_Pi
	.cfi_endproc
.LFE2058:
	.size	_Z19calculate_col_normsPfS_Pi, .-_Z19calculate_col_normsPfS_Pi
	.globl	_Z36__device_stub__Z14switch_columnsPfPiPfPi
	.type	_Z36__device_stub__Z14switch_columnsPfPiPfPi, @function
_Z36__device_stub__Z14switch_columnsPfPiPfPi:
.LFB2059:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 8(%rsp)
	leaq	32(%rsp), %rcx
	leaq	24(%rsp), %rdx
	movq	%rsi, (%rsp)
	leaq	40(%rsp), %rdi
	leaq	52(%rsp), %rsi
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 48(%rsp)
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 40(%rsp)
	movq	%rax, 52(%rsp)
	movl	$1, 60(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L23
	pushq	32(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z14switch_columnsPfPi(%rip), %rdi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 144
	movq	68(%rsp), %rcx
	movl	76(%rsp), %r8d
	movq	56(%rsp), %rsi
	movl	64(%rsp), %edx
	leaq	104(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L23:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L25
	call	__stack_chk_fail@PLT
.L25:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2059:
	.size	_Z36__device_stub__Z14switch_columnsPfPiPfPi, .-_Z36__device_stub__Z14switch_columnsPfPiPfPi
	.globl	_Z14switch_columnsPfPi
	.type	_Z14switch_columnsPfPi, @function
_Z14switch_columnsPfPi:
.LFB2060:
	.cfi_startproc
	endbr64
	jmp	_Z36__device_stub__Z14switch_columnsPfPiPfPi
	.cfi_endproc
.LFE2060:
	.size	_Z14switch_columnsPfPi, .-_Z14switch_columnsPfPi
	.globl	_Z44__device_stub__Z18shift_first_columnPfS_S_S_PfS_S_S_
	.type	_Z44__device_stub__Z18shift_first_columnPfS_S_S_PfS_S_S_, @function
_Z44__device_stub__Z18shift_first_columnPfS_S_S_PfS_S_S_:
.LFB2061:
	.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
	movq	%rcx, (%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)
	movq	%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	.L28
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z18shift_first_columnPfS_S_S_(%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
.L28:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L30
	call	__stack_chk_fail@PLT
.L30:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2061:
	.size	_Z44__device_stub__Z18shift_first_columnPfS_S_S_PfS_S_S_, .-_Z44__device_stub__Z18shift_first_columnPfS_S_S_PfS_S_S_
	.globl	_Z18shift_first_columnPfS_S_S_
	.type	_Z18shift_first_columnPfS_S_S_, @function
_Z18shift_first_columnPfS_S_S_:
.LFB2062:
	.cfi_startproc
	endbr64
	jmp	_Z44__device_stub__Z18shift_first_columnPfS_S_S_PfS_S_S_
	.cfi_endproc
.LFE2062:
	.size	_Z18shift_first_columnPfS_S_S_, .-_Z18shift_first_columnPfS_S_S_
	.globl	_Z44__device_stub__Z22normalize_first_columnPfPiPfPi
	.type	_Z44__device_stub__Z22normalize_first_columnPfPiPfPi, @function
_Z44__device_stub__Z22normalize_first_columnPfPiPfPi:
.LFB2063:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 8(%rsp)
	leaq	32(%rsp), %rcx
	leaq	24(%rsp), %rdx
	movq	%rsi, (%rsp)
	leaq	40(%rsp), %rdi
	leaq	52(%rsp), %rsi
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 48(%rsp)
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 40(%rsp)
	movq	%rax, 52(%rsp)
	movl	$1, 60(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L33
	pushq	32(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z22normalize_first_columnPfPi(%rip), %rdi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 144
	movq	68(%rsp), %rcx
	movl	76(%rsp), %r8d
	movq	56(%rsp), %rsi
	movl	64(%rsp), %edx
	leaq	104(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L33:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L35
	call	__stack_chk_fail@PLT
.L35:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2063:
	.size	_Z44__device_stub__Z22normalize_first_columnPfPiPfPi, .-_Z44__device_stub__Z22normalize_first_columnPfPiPfPi
	.globl	_Z22normalize_first_columnPfPi
	.type	_Z22normalize_first_columnPfPi, @function
_Z22normalize_first_columnPfPi:
.LFB2064:
	.cfi_startproc
	endbr64
	jmp	_Z44__device_stub__Z22normalize_first_columnPfPiPfPi
	.cfi_endproc
.LFE2064:
	.size	_Z22normalize_first_columnPfPi, .-_Z22normalize_first_columnPfPi
	.globl	_Z35__device_stub__Z13make_identityPfPiPfPi
	.type	_Z35__device_stub__Z13make_identityPfPiPfPi, @function
_Z35__device_stub__Z13make_identityPfPiPfPi:
.LFB2065:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 8(%rsp)
	leaq	32(%rsp), %rcx
	leaq	24(%rsp), %rdx
	movq	%rsi, (%rsp)
	leaq	40(%rsp), %rdi
	leaq	52(%rsp), %rsi
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 48(%rsp)
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 40(%rsp)
	movq	%rax, 52(%rsp)
	movl	$1, 60(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L38
	pushq	32(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z13make_identityPfPi(%rip), %rdi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 144
	movq	68(%rsp), %rcx
	movl	76(%rsp), %r8d
	movq	56(%rsp), %rsi
	movl	64(%rsp), %edx
	leaq	104(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L38:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L40
	call	__stack_chk_fail@PLT
.L40:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2065:
	.size	_Z35__device_stub__Z13make_identityPfPiPfPi, .-_Z35__device_stub__Z13make_identityPfPiPfPi
	.globl	_Z13make_identityPfPi
	.type	_Z13make_identityPfPi, @function
_Z13make_identityPfPi:
.LFB2066:
	.cfi_startproc
	endbr64
	jmp	_Z35__device_stub__Z13make_identityPfPiPfPi
	.cfi_endproc
.LFE2066:
	.size	_Z13make_identityPfPi, .-_Z13make_identityPfPi
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"_Z13make_identityPfPi"
.LC1:
	.string	"_Z22normalize_first_columnPfPi"
.LC2:
	.string	"_Z18shift_first_columnPfS_S_S_"
.LC3:
	.string	"_Z14switch_columnsPfPi"
.LC4:
	.string	"_Z19calculate_col_normsPfS_Pi"
.LC5:
	.string	"_Z17calculate_secantsPfS_Pi"
.LC6:
	.string	"_Z10readVecIntPi"
.LC7:
	.string	"_Z7readVecPf"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2068:
	.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	.LC0(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z13make_identityPfPi(%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
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC1(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z22normalize_first_columnPfPi(%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
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC2(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z18shift_first_columnPfS_S_S_(%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
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC3(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z14switch_columnsPfPi(%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
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC4(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z19calculate_col_normsPfS_Pi(%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
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC5(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z17calculate_secantsPfS_Pi(%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
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC6(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z10readVecIntPi(%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
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC7(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z7readVecPf(%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
.LFE2068:
	.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
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:


#include <hip/hip_runtime.h>
/*! \file
*  \brief Header file for custom GPU kernels used in the SAP algorithm
*
*/

__global__ void readVec(float * d_vector);

__global__ void readVecInt(int * d_vector);

__global__ void calculate_secants(float * d_secants, float * dpoints_in, int * d_int_constants);

__global__ void calculate_col_norms(float * d_matrix, float * d_col_norms, int * d_int_constants);

__global__ void switch_columns(float * d_matrix, int * d_column_switch_indices);

__global__ void shift_first_column(float * d_matrix, float * projectioni, float * d_new_column, float * d_algo_constants);

__global__ void normalize_first_column(float * d_matrix, int * d_int_constants);

__global__ void make_identity(float * d_matrix, int * d_int_constants);










	.text
	.file	"kernels.hip"
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_

	.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_
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.amdgpu_metadata
---
amdhsa.kernels:  []
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
