#include "credit.cuh"

__device__ double RationalApproximation(double t)
{
    double c[] = {2.515517, 0.802853, 0.010328};
    double d[] = {1.432788, 0.189269, 0.001308};
    return t - ((c[2]*t + c[1])*t + c[0]) /
               (((d[2]*t + d[1])*t + d[0])*t + 1.0);
}

__device__ double qNorm(double p)
{
    if ( (p < 0.5) & (p > 0))
    {
        return -RationalApproximation( sqrt(-2.0*log(p)) );
    }
    else if ( (p >= 0.5) & (p < 1))
    {
        return RationalApproximation( sqrt(-2.0*log(1-p)) );
    }
}

Credit::Credit() {}

Credit::~Credit() {}

Credit::Credit(double PD, double EAD, double LGD, double FG, double FL) : PD(PD), EAD(EAD), LGD(LGD), FG(FG), FL(FL)
{
    this->rho = pow(FG, 2) + pow(FL, 2);
    this->FI  = sqrt( 1 - this->rho );
}

__device__ double Credit::loss(double rG, double rL, curandState * state)
{
    double Y = sqrt(this->rho) * ( this->FG * rG + FL * rL ) + sqrt( 1 - this->rho ) * curand_normal(state);
    return (qNorm(this->PD) >= Y) * EAD * LGD;
}

__device__ Portfolio::Portfolio(Credit *cartera, int n): carte(cartera), n(n) {}

__device__ Portfolio::~Portfolio() {}

__device__ double Portfolio::loss(double fg, double fl, curandState *state)
{
    double loss = 0;
    for (int i = 0; i < this->n; i++)
    {
        loss += this->carte[i].loss(fg, fl, state);
    }
    return loss;
}

	.file	"tmpxft_0026dfe7_00000000-6_credit.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB4343:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE4343:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z21RationalApproximationd
	.type	_Z21RationalApproximationd, @function
_Z21RationalApproximationd:
.LFB4322:
	.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
.LFE4322:
	.size	_Z21RationalApproximationd, .-_Z21RationalApproximationd
	.globl	_Z5qNormd
	.type	_Z5qNormd, @function
_Z5qNormd:
.LFB5013:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_Z21RationalApproximationd
	.cfi_endproc
.LFE5013:
	.size	_Z5qNormd, .-_Z5qNormd
	.align 2
	.globl	_ZN6CreditC2Ev
	.type	_ZN6CreditC2Ev, @function
_ZN6CreditC2Ev:
.LFB4325:
	.cfi_startproc
	endbr64
	ret
	.cfi_endproc
.LFE4325:
	.size	_ZN6CreditC2Ev, .-_ZN6CreditC2Ev
	.globl	_ZN6CreditC1Ev
	.set	_ZN6CreditC1Ev,_ZN6CreditC2Ev
	.align 2
	.globl	_ZN6CreditD2Ev
	.type	_ZN6CreditD2Ev, @function
_ZN6CreditD2Ev:
.LFB4328:
	.cfi_startproc
	endbr64
	ret
	.cfi_endproc
.LFE4328:
	.size	_ZN6CreditD2Ev, .-_ZN6CreditD2Ev
	.globl	_ZN6CreditD1Ev
	.set	_ZN6CreditD1Ev,_ZN6CreditD2Ev
	.align 2
	.globl	_ZN6CreditC2Eddddd
	.type	_ZN6CreditC2Eddddd, @function
_ZN6CreditC2Eddddd:
.LFB4331:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movq	%rdi, %rbx
	movsd	%xmm3, 24(%rdi)
	mulsd	%xmm3, %xmm3
	movsd	%xmm4, 32(%rdi)
	mulsd	%xmm4, %xmm4
	movsd	%xmm0, (%rdi)
	movsd	.LC0(%rip), %xmm0
	movsd	%xmm1, 8(%rdi)
	movsd	%xmm2, 16(%rdi)
	addsd	%xmm4, %xmm3
	subsd	%xmm3, %xmm0
	movsd	%xmm3, 48(%rdi)
	call	sqrt@PLT
	movsd	%xmm0, 40(%rbx)
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4331:
	.size	_ZN6CreditC2Eddddd, .-_ZN6CreditC2Eddddd
	.globl	_ZN6CreditC1Eddddd
	.set	_ZN6CreditC1Eddddd,_ZN6CreditC2Eddddd
	.align 2
	.globl	_ZN6Credit4lossEddP17curandStateXORWOW
	.type	_ZN6Credit4lossEddP17curandStateXORWOW, @function
_ZN6Credit4lossEddP17curandStateXORWOW:
.LFB4333:
	.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
.LFE4333:
	.size	_ZN6Credit4lossEddP17curandStateXORWOW, .-_ZN6Credit4lossEddP17curandStateXORWOW
	.align 2
	.globl	_ZN9PortfolioC2EP6Crediti
	.type	_ZN9PortfolioC2EP6Crediti, @function
_ZN9PortfolioC2EP6Crediti:
.LFB4335:
	.cfi_startproc
	endbr64
	movq	%rsi, (%rdi)
	xorl	%eax, %eax
	movl	%edx, 8(%rdi)
	movq	%rax, -16(%rsp)
	movq	-16(%rsp), %rdi
	jmp	free@PLT
	.cfi_endproc
.LFE4335:
	.size	_ZN9PortfolioC2EP6Crediti, .-_ZN9PortfolioC2EP6Crediti
	.globl	_ZN9PortfolioC1EP6Crediti
	.set	_ZN9PortfolioC1EP6Crediti,_ZN9PortfolioC2EP6Crediti
	.align 2
	.globl	_ZN9PortfolioD2Ev
	.type	_ZN9PortfolioD2Ev, @function
_ZN9PortfolioD2Ev:
.LFB4338:
	.cfi_startproc
	endbr64
	xorl	%eax, %eax
	movq	%rax, -16(%rsp)
	movq	-16(%rsp), %rdi
	jmp	free@PLT
	.cfi_endproc
.LFE4338:
	.size	_ZN9PortfolioD2Ev, .-_ZN9PortfolioD2Ev
	.globl	_ZN9PortfolioD1Ev
	.set	_ZN9PortfolioD1Ev,_ZN9PortfolioD2Ev
	.align 2
	.globl	_ZN9Portfolio4lossEddP17curandStateXORWOW
	.type	_ZN9Portfolio4lossEddP17curandStateXORWOW, @function
_ZN9Portfolio4lossEddP17curandStateXORWOW:
.LFB5011:
	.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
.LFE5011:
	.size	_ZN9Portfolio4lossEddP17curandStateXORWOW, .-_ZN9Portfolio4lossEddP17curandStateXORWOW
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC1:
	.string	"precalc_xorwow_matrix"
.LC2:
	.string	"precalc_xorwow_offset_matrix"
.LC3:
	.string	"mrg32k3aM1"
.LC4:
	.string	"mrg32k3aM2"
.LC5:
	.string	"mrg32k3aM1SubSeq"
.LC6:
	.string	"mrg32k3aM2SubSeq"
.LC7:
	.string	"mrg32k3aM1Seq"
.LC8:
	.string	"mrg32k3aM2Seq"
.LC9:
	.string	"__cr_lgamma_table"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB4366:
	.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
	movl	$102400, %r9d
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC1(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	movq	%rdx, %rcx
	leaq	_ZL21precalc_xorwow_matrix(%rip), %rsi
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterVar@PLT
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rdx
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC2(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$102400, %r9d
	leaq	_ZL28precalc_xorwow_offset_matrix(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rcx
	.cfi_def_cfa_offset 24
	popq	%rsi
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC3(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$2304, %r9d
	leaq	_ZL10mrg32k3aM1(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rdi
	.cfi_def_cfa_offset 24
	popq	%r8
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC4(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$2304, %r9d
	leaq	_ZL10mrg32k3aM2(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%r9
	.cfi_def_cfa_offset 24
	popq	%r10
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC5(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$2016, %r9d
	leaq	_ZL16mrg32k3aM1SubSeq(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%r11
	.cfi_def_cfa_offset 24
	popq	%rax
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC6(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$2016, %r9d
	leaq	_ZL16mrg32k3aM2SubSeq(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rdx
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC7(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$2304, %r9d
	leaq	_ZL13mrg32k3aM1Seq(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rcx
	.cfi_def_cfa_offset 24
	popq	%rsi
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC8(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$2304, %r9d
	leaq	_ZL13mrg32k3aM2Seq(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rdi
	.cfi_def_cfa_offset 24
	popq	%r8
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$1
	.cfi_def_cfa_offset 32
	leaq	.LC9(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movl	$72, %r9d
	movq	%rdx, %rcx
	leaq	_ZL17__cr_lgamma_table(%rip), %rsi
	call	__cudaRegisterVar@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	popq	%r9
	.cfi_def_cfa_offset 24
	popq	%r10
	.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
.LFE4366:
	.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
	.local	_ZL17__cr_lgamma_table
	.comm	_ZL17__cr_lgamma_table,72,32
	.local	_ZL13mrg32k3aM2Seq
	.comm	_ZL13mrg32k3aM2Seq,2304,32
	.local	_ZL13mrg32k3aM1Seq
	.comm	_ZL13mrg32k3aM1Seq,2304,32
	.local	_ZL16mrg32k3aM2SubSeq
	.comm	_ZL16mrg32k3aM2SubSeq,2016,32
	.local	_ZL16mrg32k3aM1SubSeq
	.comm	_ZL16mrg32k3aM1SubSeq,2016,32
	.local	_ZL10mrg32k3aM2
	.comm	_ZL10mrg32k3aM2,2304,32
	.local	_ZL10mrg32k3aM1
	.comm	_ZL10mrg32k3aM1,2304,32
	.local	_ZL28precalc_xorwow_offset_matrix
	.comm	_ZL28precalc_xorwow_offset_matrix,102400,32
	.local	_ZL21precalc_xorwow_matrix
	.comm	_ZL21precalc_xorwow_matrix,102400,32
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC0:
	.long	0
	.long	1072693248
	.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 CREDIT_H
#define CREDIT_H

#include <hip/hip_runtime.h>
#include <math.h>
#include <hiprand/hiprand_kernel.h>
#include <random>

class Credit
{
    double PD, EAD, LGD, FG, FL, FI, rho;

public:
    Credit();
    Credit(double PD, double EAD, double LGD, double FG, double FL);
    ~Credit();

    __device__ double loss(double rG, double rL, hiprandState * state);
};

class Portfolio
{
    Credit * carte;
    int n;

public:
    __device__ Portfolio(Credit * cartera, int n);
    __device__ ~Portfolio();

    __device__ double loss(double fg, double fl, hiprandState * state);
};



#endif // CREDIT_H

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