#include <math.h>
#include <stdio.h>
#include <cuda_runtime.h>

#define PI 3.141592
/* 
 *  Return number of hyperparameters
 */
__device__ int retNumParams(){
	return 4;
}

/*
 * Compute the covariance matrix in column oriented vector format using X
 * The output matrix is upper triangular
 * Output matrix formula k(x,z) = sf^2 * exp( -2*sin^2( pi*||x-z||/p )/ell^2 )	
 */
__device__ void covMat(double *X, int nX, double *cM, double *lh){
	double l = exp(lh[0]);
	double sf2 = exp(2 * lh[1]);
	double p = exp(lh[2]);
	double sn2 = exp(2 * lh[3]);
	int i, j;
	for(i = 0; i < nX; i++){
		for(j = i; j < nX; j++){
			double _int = fabs(X[i] - X[j]);
			cM[i + j * nX] = sf2 * exp(-2 * pow(sin((PI * _int) / p), 2) / pow(l, 2));
			if(j == i) cM[i + j * nX] += sn2;
		}

	}
}


/*
 * Compute the covariance matrix in column oriented vector format using X
 * The output matrix is upper triangular
 * Output matrix formula k(x,z) = sf^2 * exp( -2*sin^2( pi*||x-z||/p )/ell^2 )
 * same as above except doesn't add sn2 to main diagonal	
 * Divides all elements by sn2 (when noise is high) and then adds an identity matrix
 */
__device__ void covMatSp(double *X, int nX, double *cM, double *lh){
	double l = exp(lh[0]);
	double sf2 = exp(2 * lh[1]);
	double p = exp(lh[2]);
	double sn2 = exp(2 * lh[3]);
	for(int i = 0; i < nX; i++)
	{
		for(int j = i; j < nX; j++)
		{
			double _int = fabs(X[i] - X[j]);
			cM[i + j * nX] = sf2 * exp(-2 * pow(sin((PI * _int) / p), 2) / pow(l, 2)) / sn2;
			if (j == i)
			{
				cM[i + j * nX] += 1;
			}
		}
	}
}

/*
 * Compute the derivative of the covariance matrix in column oriented vector format using X with respect to hyperparameter #n
 * The output matrix is upper triangular
 */
__device__ void dcovMat(double *X, int nX, double *dcM, double *lh, int n){
	double l = exp(lh[0]);
	double sf2 = exp(2*lh[1]);
	double p = exp(lh[2]);
	double sn2 = exp(2*lh[3]);
	int i, j;
	double tmp100;
	for(i = 0; i < nX; i++)
	{
		for(j = i; j < nX; j++)
		{
			double _int = fabs(X[i]-X[j]);
			switch(n){
				case 1://log(l)
					dcM[i+j*nX] = (sf2*exp(-2*pow(sin((PI*_int)/p),2)/pow(l,2)))*4*(pow(sin((PI*_int)/p),2)/pow(l,2));
					break;
				case 2://log(sf)
					dcM[i+j*nX] = 2*sf2*exp(-2*pow(sin((PI*_int)/p),2)/pow(l,2));
					break;
				case 3://log(p)
					//	dcM[i+j*nX] = sf2*exp(-2*pow(sin((PI*_int)/p),2)/pow(l,2))*(4*sin((PI*_int)/p))*((PI*_int)/(pow(p,2)*pow(l,2)))*(cos((PI*_int)/p));
					tmp100 = (PI*_int)/p;
					dcM[i+j*nX]= ((sf2*4*tmp100)/(pow(l,2)))*exp((-2*pow(sin(tmp100),2))/pow(l,2))*sin(tmp100)*cos(tmp100);
					break;
				case 4://log(sn)
					if(i == j)
						dcM[i+j*nX] = 2*sn2;
					else
						dcM[i+j*nX] = 0;
					break;
				default:
					return;
			}
		}
	}
}

/* 
 * Compute element wise cross-covariance matrix between a vector of training observations, X and a vector test observations, Y.
 * The output matrix cM is a nX x nY matrix in column oriented format
 */

__device__ void covMatCross(double *X, double *Y, int nX, int nY, double *cM, double *lh){
	double l = exp(lh[0]);
	double sf2 = exp(2 * lh[1]);
	double p = exp(lh[2]);
	int i, j;
	for(j = 0; j < nY; j++)
	{
		for(i = 0; i < nX; i++)
		{
			double _int = fabs(X[i] - Y[j]);
			cM[i + j * nX] = sf2 * exp(-2 * pow(sin((PI * _int) / p), 2) / pow(l, 2));
		}
	}
}
	.file	"tmpxft_00380d60_00000000-6_covPNoise.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2034:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2034:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z12retNumParamsv
	.type	_Z12retNumParamsv, @function
_Z12retNumParamsv:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2027:
	.size	_Z12retNumParamsv, .-_Z12retNumParamsv
	.globl	_Z6covMatPdiS_S_
	.type	_Z6covMatPdiS_S_, @function
_Z6covMatPdiS_S_:
.LFB2028:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2028:
	.size	_Z6covMatPdiS_S_, .-_Z6covMatPdiS_S_
	.globl	_Z8covMatSpPdiS_S_
	.type	_Z8covMatSpPdiS_S_, @function
_Z8covMatSpPdiS_S_:
.LFB2173:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2173:
	.size	_Z8covMatSpPdiS_S_, .-_Z8covMatSpPdiS_S_
	.globl	_Z7dcovMatPdiS_S_i
	.type	_Z7dcovMatPdiS_S_i, @function
_Z7dcovMatPdiS_S_i:
.LFB2030:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2030:
	.size	_Z7dcovMatPdiS_S_i, .-_Z7dcovMatPdiS_S_i
	.globl	_Z11covMatCrossPdS_iiS_S_
	.type	_Z11covMatCrossPdS_iiS_S_, @function
_Z11covMatCrossPdS_iiS_S_:
.LFB2031:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2031:
	.size	_Z11covMatCrossPdS_iiS_S_, .-_Z11covMatCrossPdS_iiS_S_
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2057:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	movq	%rax, %rdi
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2057:
	.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:

#ifndef COVPNOISE_CUH
#define COVPNOISE_CUH

/* covariance function related routines */
__device__ int retNumParams();
__device__ void covVec(double *X, int nX, double *cV, double *lh);
__device__ void dcovVec(double *X, int nX, double *dcV, double *lh, int n);
__device__ void covMat(double *X, int nX, double *cM, double *lh);
__device__ void covMatSp(double *X, int nX, double *cM, double *lh);
__device__ void dcovMat(double *X, int nX, double *dcM, double *lh, int n);
__device__ void covMatCross(double *X, double *Y, int nX, int nY, double *cM, double *lh);
__device__ void dcovMatCross(double *X, double *Y, int nX, int nY, double *dcM, double *lh, int n);

#endif

	.text
	.file	"covPNoise.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
