//=====================================================================
//	MAIN FUNCTION
//=====================================================================
__device__ void kernel_fin_2(int timeinst, float* d_initvalu, float* d_finavalu, int offset_ecc,
    int offset_Dyad, int offset_SL, int offset_Cyt, float* d_params, float* d_com) {

//=====================================================================
//	VARIABLES
//=====================================================================

// input parameters
  float BtotDyad;
  float CaMKIItotDyad;

  // compute variables
  float Vmyo;																			// [L]
  float Vdyad;																			// [L]
  float VSL;																				// [L]
  // float kDyadSL;																			// [L/msec]
  float kSLmyo;																			// [L/msec]
  float k0Boff;																			// [s^-1]
  float k0Bon;																			// [uM^-1 s^-1] kon = koff/Kd
  float k2Boff;																			// [s^-1]
  float k2Bon;																			// [uM^-1 s^-1]
  // float k4Boff;																			// [s^-1]
  float k4Bon;																			// [uM^-1 s^-1]
  float CaMtotDyad;
  float Bdyad;																			// [uM dyad]
  float J_cam_dyadSL;																	// [uM/msec dyad]
  float J_ca2cam_dyadSL;																	// [uM/msec dyad]
  float J_ca4cam_dyadSL;																	// [uM/msec dyad]
  float J_cam_SLmyo;																		// [umol/msec]
  float J_ca2cam_SLmyo;																	// [umol/msec]
  float J_ca4cam_SLmyo;																	// [umol/msec]

//=====================================================================
//	COMPUTATION
//=====================================================================

  // input parameters
  BtotDyad = d_params[1];
  CaMKIItotDyad = d_params[2];

  // ADJUST ECC incorporate Ca buffering from CaM, convert JCaCyt from uM/msec to mM/msec
  d_finavalu[offset_ecc + 35] = d_finavalu[offset_ecc + 35] + 1e-3 * d_com[0];
  d_finavalu[offset_ecc + 36] = d_finavalu[offset_ecc + 36] + 1e-3 * d_com[1];
  d_finavalu[offset_ecc + 37] = d_finavalu[offset_ecc + 37] + 1e-3 * d_com[2];

  // incorporate CaM diffusion between compartments
  Vmyo = 2.1454e-11;																// [L]
  Vdyad = 1.7790e-14;																// [L]
  VSL = 6.6013e-13;																// [L]
  // kDyadSL = 3.6363e-16;															// [L/msec]
  kSLmyo = 8.587e-15;																// [L/msec]
  k0Boff = 0.0014;																// [s^-1]
  k0Bon = k0Boff / 0.2;																// [uM^-1 s^-1] kon = koff/Kd
  k2Boff = k0Boff / 100;															// [s^-1]
  k2Bon = k0Bon;																	// [uM^-1 s^-1]
  // k4Boff = k2Boff;																// [s^-1]
  k4Bon = k0Bon;																	// [uM^-1 s^-1]
  CaMtotDyad = d_initvalu[offset_Dyad + 0] + d_initvalu[offset_Dyad + 1]
      + d_initvalu[offset_Dyad + 2] + d_initvalu[offset_Dyad + 3] + d_initvalu[offset_Dyad + 4]
      + d_initvalu[offset_Dyad + 5]
      + CaMKIItotDyad
          * (d_initvalu[offset_Dyad + 6] + d_initvalu[offset_Dyad + 7] + d_initvalu[offset_Dyad + 8]
              + d_initvalu[offset_Dyad + 9]) + d_initvalu[offset_Dyad + 12]
      + d_initvalu[offset_Dyad + 13] + d_initvalu[offset_Dyad + 14];
  Bdyad = BtotDyad - CaMtotDyad;																				// [uM dyad]
  J_cam_dyadSL = 1e-3
      * (k0Boff * d_initvalu[offset_Dyad + 0] - k0Bon * Bdyad * d_initvalu[offset_SL + 0]);	// [uM/msec dyad]
  J_ca2cam_dyadSL = 1e-3
      * (k2Boff * d_initvalu[offset_Dyad + 1] - k2Bon * Bdyad * d_initvalu[offset_SL + 1]);	// [uM/msec dyad]
  J_ca4cam_dyadSL = 1e-3
      * (k2Boff * d_initvalu[offset_Dyad + 2] - k4Bon * Bdyad * d_initvalu[offset_SL + 2]);	// [uM/msec dyad]

  J_cam_SLmyo = kSLmyo * (d_initvalu[offset_SL + 0] - d_initvalu[offset_Cyt + 0]);		// [umol/msec]
  J_ca2cam_SLmyo = kSLmyo * (d_initvalu[offset_SL + 1] - d_initvalu[offset_Cyt + 1]);	// [umol/msec]
  J_ca4cam_SLmyo = kSLmyo * (d_initvalu[offset_SL + 2] - d_initvalu[offset_Cyt + 2]);	// [umol/msec]

  // ADJUST CAM Dyad
  d_finavalu[offset_Dyad + 0] = d_finavalu[offset_Dyad + 0] - J_cam_dyadSL;
  d_finavalu[offset_Dyad + 1] = d_finavalu[offset_Dyad + 1] - J_ca2cam_dyadSL;
  d_finavalu[offset_Dyad + 2] = d_finavalu[offset_Dyad + 2] - J_ca4cam_dyadSL;

  // ADJUST CAM Sl
  d_finavalu[offset_SL + 0] = d_finavalu[offset_SL + 0] + J_cam_dyadSL * Vdyad / VSL
      - J_cam_SLmyo / VSL;
  d_finavalu[offset_SL + 1] = d_finavalu[offset_SL + 1] + J_ca2cam_dyadSL * Vdyad / VSL
      - J_ca2cam_SLmyo / VSL;
  d_finavalu[offset_SL + 2] = d_finavalu[offset_SL + 2] + J_ca4cam_dyadSL * Vdyad / VSL
      - J_ca4cam_SLmyo / VSL;

  // ADJUST CAM Cyt
  d_finavalu[offset_Cyt + 0] = d_finavalu[offset_Cyt + 0] + J_cam_SLmyo / Vmyo;
  d_finavalu[offset_Cyt + 1] = d_finavalu[offset_Cyt + 1] + J_ca2cam_SLmyo / Vmyo;
  d_finavalu[offset_Cyt + 2] = d_finavalu[offset_Cyt + 2] + J_ca4cam_SLmyo / Vmyo;

}

	.file	"tmpxft_00270f17_00000000-6_kernel_fin_2.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2011:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2011:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z12kernel_fin_2iPfS_iiiiS_S_
	.type	_Z12kernel_fin_2iPfS_iiiiS_S_, @function
_Z12kernel_fin_2iPfS_iiiiS_S_:
.LFB2008:
	.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
.LFE2008:
	.size	_Z12kernel_fin_2iPfS_iiiiS_S_, .-_Z12kernel_fin_2iPfS_iiiiS_S_
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2034:
	.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
.LFE2034:
	.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>
//=====================================================================
//	MAIN FUNCTION
//=====================================================================
__device__ void kernel_fin_2(int timeinst, float* d_initvalu, float* d_finavalu, int offset_ecc,
    int offset_Dyad, int offset_SL, int offset_Cyt, float* d_params, float* d_com) {

//=====================================================================
//	VARIABLES
//=====================================================================

// input parameters
  float BtotDyad;
  float CaMKIItotDyad;

  // compute variables
  float Vmyo;																			// [L]
  float Vdyad;																			// [L]
  float VSL;																				// [L]
  // float kDyadSL;																			// [L/msec]
  float kSLmyo;																			// [L/msec]
  float k0Boff;																			// [s^-1]
  float k0Bon;																			// [uM^-1 s^-1] kon = koff/Kd
  float k2Boff;																			// [s^-1]
  float k2Bon;																			// [uM^-1 s^-1]
  // float k4Boff;																			// [s^-1]
  float k4Bon;																			// [uM^-1 s^-1]
  float CaMtotDyad;
  float Bdyad;																			// [uM dyad]
  float J_cam_dyadSL;																	// [uM/msec dyad]
  float J_ca2cam_dyadSL;																	// [uM/msec dyad]
  float J_ca4cam_dyadSL;																	// [uM/msec dyad]
  float J_cam_SLmyo;																		// [umol/msec]
  float J_ca2cam_SLmyo;																	// [umol/msec]
  float J_ca4cam_SLmyo;																	// [umol/msec]

//=====================================================================
//	COMPUTATION
//=====================================================================

  // input parameters
  BtotDyad = d_params[1];
  CaMKIItotDyad = d_params[2];

  // ADJUST ECC incorporate Ca buffering from CaM, convert JCaCyt from uM/msec to mM/msec
  d_finavalu[offset_ecc + 35] = d_finavalu[offset_ecc + 35] + 1e-3 * d_com[0];
  d_finavalu[offset_ecc + 36] = d_finavalu[offset_ecc + 36] + 1e-3 * d_com[1];
  d_finavalu[offset_ecc + 37] = d_finavalu[offset_ecc + 37] + 1e-3 * d_com[2];

  // incorporate CaM diffusion between compartments
  Vmyo = 2.1454e-11;																// [L]
  Vdyad = 1.7790e-14;																// [L]
  VSL = 6.6013e-13;																// [L]
  // kDyadSL = 3.6363e-16;															// [L/msec]
  kSLmyo = 8.587e-15;																// [L/msec]
  k0Boff = 0.0014;																// [s^-1]
  k0Bon = k0Boff / 0.2;																// [uM^-1 s^-1] kon = koff/Kd
  k2Boff = k0Boff / 100;															// [s^-1]
  k2Bon = k0Bon;																	// [uM^-1 s^-1]
  // k4Boff = k2Boff;																// [s^-1]
  k4Bon = k0Bon;																	// [uM^-1 s^-1]
  CaMtotDyad = d_initvalu[offset_Dyad + 0] + d_initvalu[offset_Dyad + 1]
      + d_initvalu[offset_Dyad + 2] + d_initvalu[offset_Dyad + 3] + d_initvalu[offset_Dyad + 4]
      + d_initvalu[offset_Dyad + 5]
      + CaMKIItotDyad
          * (d_initvalu[offset_Dyad + 6] + d_initvalu[offset_Dyad + 7] + d_initvalu[offset_Dyad + 8]
              + d_initvalu[offset_Dyad + 9]) + d_initvalu[offset_Dyad + 12]
      + d_initvalu[offset_Dyad + 13] + d_initvalu[offset_Dyad + 14];
  Bdyad = BtotDyad - CaMtotDyad;																				// [uM dyad]
  J_cam_dyadSL = 1e-3
      * (k0Boff * d_initvalu[offset_Dyad + 0] - k0Bon * Bdyad * d_initvalu[offset_SL + 0]);	// [uM/msec dyad]
  J_ca2cam_dyadSL = 1e-3
      * (k2Boff * d_initvalu[offset_Dyad + 1] - k2Bon * Bdyad * d_initvalu[offset_SL + 1]);	// [uM/msec dyad]
  J_ca4cam_dyadSL = 1e-3
      * (k2Boff * d_initvalu[offset_Dyad + 2] - k4Bon * Bdyad * d_initvalu[offset_SL + 2]);	// [uM/msec dyad]

  J_cam_SLmyo = kSLmyo * (d_initvalu[offset_SL + 0] - d_initvalu[offset_Cyt + 0]);		// [umol/msec]
  J_ca2cam_SLmyo = kSLmyo * (d_initvalu[offset_SL + 1] - d_initvalu[offset_Cyt + 1]);	// [umol/msec]
  J_ca4cam_SLmyo = kSLmyo * (d_initvalu[offset_SL + 2] - d_initvalu[offset_Cyt + 2]);	// [umol/msec]

  // ADJUST CAM Dyad
  d_finavalu[offset_Dyad + 0] = d_finavalu[offset_Dyad + 0] - J_cam_dyadSL;
  d_finavalu[offset_Dyad + 1] = d_finavalu[offset_Dyad + 1] - J_ca2cam_dyadSL;
  d_finavalu[offset_Dyad + 2] = d_finavalu[offset_Dyad + 2] - J_ca4cam_dyadSL;

  // ADJUST CAM Sl
  d_finavalu[offset_SL + 0] = d_finavalu[offset_SL + 0] + J_cam_dyadSL * Vdyad / VSL
      - J_cam_SLmyo / VSL;
  d_finavalu[offset_SL + 1] = d_finavalu[offset_SL + 1] + J_ca2cam_dyadSL * Vdyad / VSL
      - J_ca2cam_SLmyo / VSL;
  d_finavalu[offset_SL + 2] = d_finavalu[offset_SL + 2] + J_ca4cam_dyadSL * Vdyad / VSL
      - J_ca4cam_SLmyo / VSL;

  // ADJUST CAM Cyt
  d_finavalu[offset_Cyt + 0] = d_finavalu[offset_Cyt + 0] + J_cam_SLmyo / Vmyo;
  d_finavalu[offset_Cyt + 1] = d_finavalu[offset_Cyt + 1] + J_ca2cam_SLmyo / Vmyo;
  d_finavalu[offset_Cyt + 2] = d_finavalu[offset_Cyt + 2] + J_ca4cam_SLmyo / Vmyo;

}

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