/**
 * @file   : main_draft.cu
 * @brief  : main file draft for 2-dim. Ising in CUDA C++11/14, 
 * @details : 
 * 
 * @author : Ernest Yeung <ernestyalumni@gmail.com>
 * @date   : 20180108    
 * @ref    : M. Hjorth-Jensen, Computational Physics, University of Oslo (2015)
 * https://github.com/CompPhysics/ComputationalPhysics/blob/master/doc/Programs/LecturePrograms/programs/StatPhys/cpp/ising_2dim.cpp
 * https://www.paypal.com/cgi-bin/webscr?cmd=_donations&business=ernestsaveschristmas%2bpaypal%40gmail%2ecom&lc=US&item_name=ernestyalumni&currency_code=USD&bn=PP%2dDonationsBF%3abtn_donateCC_LG%2egif%3aNonHosted 
 * 
 * which won't go through a 3rd. party such as indiegogo, kickstarter, patreon.  
 * Otherwise, I receive emails and messages on how all my (free) material on 
 * physics, math, and engineering have helped students with their studies, 
 * and I know what it's like to not have money as a student, but love physics 
 * (or math, sciences, etc.), so I am committed to keeping all my material 
 * open-source and free, whether or not 
 * sufficiently crowdfunded, under the open-source MIT license: 
 * 	feel free to copy, edit, paste, make your own versions, share, use as you wish.  
 *  Just don't be an asshole and not give credit where credit is due.  
 * Peace out, never give up! -EY
 * 
 * */
/* 
 * COMPILATION TIP
 * nvcc main.cu ./devgrid2d/devgrid2d.cu -o main
 * 
 * */
#include <iostream>
#include <memory>  
#include <array> // std::array in TransProb struct
#include <cmath> // std::exp in TransProb


/* ********** device GPU structs ********** */

struct Sysparam {
	float E; // total energy E
	float M; // total magnetization H 
	float T; // temperature T of the system (it's kT; treat Boltzmann constant k as a unit conversion)  
}; 

struct Avg {
	// (data) members
	// average values of physical parameters  
	float Eavg; 
	float Mavg; 
	float Esq_avg; // Esq = E*E 
	float Msq_avg; // Msq = M*M
	float absM_avg; // absM = |M| 
	float M4_avg; // M4_avg = M*M*M*M
};

/** @struct TransProb
 *  @brief transition probabilities to new spin configuration for 2-dim. Ising model 
 * */
struct TransProb {
	// (data) members
	// transition probabilities data 
	std::array<float,17> transProb; 
	float J; // spin constant 
	
	// getting functions
	/** @fn get_by_DeltaE 
	 * @details given DeltaE (\Delta E), DeltaE = -8J, -4J,...8J, we want to get the 
	 * transition probability from std::unique_ptr transprob (but transprob indexed by 
	 * 0,1,...(17-1)
	 * */
	float get_by_DeltaE(const int DeltaE) {
		return transProb[DeltaE+8]; } ;
};

// custom deleters as structs   
struct del_Sysparam_struct { void operator()(Sysparam* ptr) { cudaFree(ptr); } };
struct del_Avg_struct { void operator()(Avg* ptr) { cudaFree(ptr); } };	
struct del_TransProb_struct { void operator()(TransProb* ptr) { cudaFree(ptr); } };  	

/* ********** END of device GPU structs ********** */


/* ********** host CPU structs ********** */
struct h_Sysparam { 
	float E; // total energy E 
	float M; // total magnetization M 
	float T; // temperature T of the system 

	// constructors
	// default constructor
	/** @fn Sysparam()
	 * @brief default constructor for struct Sysparam 
	 * @details set all E,M,T parameters to 0
	 * */
	h_Sysparam() : E {0.f}, M {0.f}, T {0.f} {}; 

	/** @fn Sysparam(double, double,double)
	 * @brief constructor for struct Sysparam 
	 * */
	h_Sysparam(float E, float M, float T) : E {E}, M {M}, T {T} {} ;

	/** @fn Sysparam(double)
	 * @brief constructor for struct Sysparam, when only given the system temperature (initially)
	 * */
	h_Sysparam(float T) : E {0.f}, M {0.f}, T {T} {}; 	

};

struct h_Avg {
	// (data) members
	// average values of physical parameters  
	float Eavg; 
	float Mavg; 
	float Esq_avg; // Esq = E*E 
	float Msq_avg; // Msq = M*M
	float absM_avg; // absM = |M| 
	float M4_avg; // M4_avg = M*M*M*M
	
	// constructors
	// default constructor
	h_Avg() : Eavg {0.f}, Mavg(0.f), Esq_avg(0.f), Msq_avg{0.f}, absM_avg{0.f},M4_avg{0.f} {};

	h_Avg(float Eavg, float Mavg, float Esq_avg, float Msq_avg, float absM_avg, float M4_avg) : 
		Eavg {Eavg}, Mavg{Mavg}, Esq_avg{Esq_avg}, Msq_avg{Msq_avg}, absM_avg{absM_avg}, M4_avg{M4_avg}  { } ;

};


/** @struct TransProb
 *  @brief transition probabilities to new spin configuration for 2-dim. Ising model 
 * */
struct h_TransProb {
	// (data) members
	// transition probabilities data 
	std::array<float,17> transProb; 
	float J; // spin constant 
	
	// constructors
	// default constructor 
	h_TransProb() : J {1.f} { };

	h_TransProb(float J, h_Sysparam& sysparams) : J {J} {
		float T = sysparams.T; // temperature
		for (int de = -8; de<= 8; de+=4) { 
			transProb[de+8] = std::exp(-((float) de)/T); }
	}; 

	// getting functions
	/** @fn get_by_DeltaE 
	 * @details given DeltaE (\Delta E), DeltaE = -8J, -4J,...8J, we want to get the 
	 * transition probability from std::unique_ptr transprob (but transprob indexed by 
	 * 0,1,...(17-1)
	 * */
	float get_by_DeltaE(int DeltaE) {
		return transProb[DeltaE+8]; 
	}; 
};

/** @fn calc_transProb
 * 
 * */
void calc_transProb(TransProb & transProb, const float T) {
	for (int de = -8; de <= 8; de+=4) {
		transProb.transProb[de+8] = std::exp( -((float) de)/T); 
	}
}
 
/* ********** END of host CPU structs ********** */



int main(int argc, char* argv[]) 
{
	std::unique_ptr<Sysparam,del_Sysparam_struct> d_sysparams(nullptr, del_Sysparam_struct()); 
	cudaMallocManaged((void **) &d_sysparams, 1*sizeof(Sysparam));  
	std::unique_ptr<Avg,del_Avg_struct> d_avgs(nullptr, del_Avg_struct()); 
	cudaMallocManaged((void **) &d_avgs, 1*sizeof(Avg));  
	std::unique_ptr<TransProb,del_TransProb_struct> d_transProb(nullptr, del_TransProb_struct()); 
	cudaMallocManaged((void **) &d_transProb, 1*sizeof(TransProb));  

	/* ****************************************************************************************** */
	/* ******************** Ways to initialize structs on device GPU; 2 ways ******************** */
	/* ****************************************************************************************** */
	/* *************** 1. cudaMemcpy from host to device *************** */

	// error: have to be of same type
//	h_Sysparam h_sysparams { 1.f, 3.f, 2.f };  
//	h_Avg h_avgs { 1.1f, 2.1f, 1.2f, 2.2f, 2.3f, 2.4f }; 
//	h_TransProb(1.f, h_sysparams);  

	// some host CPU values to input 
	Sysparam h_sysparams { 1.f, 3.f, 2.f };  
	Avg h_avgs { 1.1f, 2.1f, 1.2f, 2.2f, 2.3f, 2.4f }; 
	TransProb h_transProb; 
	calc_transProb(h_transProb, h_sysparams.T); 

	cudaMemcpy( d_sysparams.get(), &h_sysparams, 1*sizeof(Sysparam), cudaMemcpyHostToDevice);  // possible error have to be of same type
	cudaMemcpy( d_avgs.get(), &h_avgs, 1*sizeof(Avg), cudaMemcpyHostToDevice);  // possible error have to be of same type
	cudaMemcpy( d_transProb.get(), &h_transProb, 1*sizeof(TransProb), cudaMemcpyHostToDevice);  // possible error have to be of same type

	/* sanity check; copy back to host for sanity check */  
	Sysparam h_sysparams_out ;  
	Avg h_avgs_out ; 
	TransProb h_transProb_out; 

	cudaMemcpy(&h_sysparams_out, d_sysparams.get(), 1*sizeof(Sysparam), cudaMemcpyDeviceToHost);  // possible error have to be of same type
	cudaMemcpy(&h_avgs_out, d_avgs.get(), 1*sizeof(Avg), cudaMemcpyDeviceToHost);  // possible error have to be of same type
	cudaMemcpy(&h_transProb_out, d_transProb.get(), 1*sizeof(TransProb), cudaMemcpyDeviceToHost);  // possible error have to be of same type
	
	std::cout << " h_sysparams_out : " << h_sysparams_out.E << " " << h_sysparams_out.M << " " << h_sysparams_out.T << std::endl; 
	std::cout << std::endl << " h_avgs_out : " << h_avgs_out.Eavg << " " << h_avgs_out.Mavg 
		<< " " << h_avgs_out.Esq_avg << " " << h_avgs_out.Msq_avg << " " << h_avgs_out.absM_avg << " " << h_avgs_out.M4_avg  << std::endl; 
	for (int de =-8; de <= 8; de+=4) {
		std::cout << h_transProb_out.transProb[de+8] << " "; } std::cout << std::endl;
	/* END of sanity check */ 

	/* *************** 2. directly from host set values *************** */
	d_sysparams->E = 1.5f; 
	d_sysparams->M = 3.5f; 
	d_sysparams->T = 2.5f; 
	cudaMemcpy(&h_sysparams_out, d_sysparams.get(), 1*sizeof(Sysparam), cudaMemcpyDeviceToHost);  // possible error have to be of same type
	std::cout << std::endl << " h_sysparams_out : " << h_sysparams_out.E << " " << h_sysparams_out.M << " " << h_sysparams_out.T << std::endl; 
	
	
	
	
}

	.file	"tmpxft_002a0ef6_00000000-6_main_draft.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB4373:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE4373:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z14calc_transProbR9TransProbf
	.type	_Z14calc_transProbR9TransProbf, @function
_Z14calc_transProbR9TransProbf:
.LFB4345:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movq	%rdi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movq	$-8, %rbx
	subq	$24, %rsp
	.cfi_def_cfa_offset 48
	movss	%xmm0, 12(%rsp)
.L3:
	cvtsi2ssl	%ebx, %xmm0
	xorps	.LC0(%rip), %xmm0
	divss	12(%rsp), %xmm0
	call	expf@PLT
	movss	%xmm0, 32(%rbp,%rbx,4)
	addq	$4, %rbx
	cmpq	$12, %rbx
	jne	.L3
	addq	$24, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4345:
	.size	_Z14calc_transProbR9TransProbf, .-_Z14calc_transProbR9TransProbf
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC2:
	.string	" h_sysparams_out : "
.LC3:
	.string	" "
.LC4:
	.string	" h_avgs_out : "
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB4346:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA4346
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	xorl	%ecx, %ecx
	movl	$1, %edx
	movl	$12, %esi
	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	$264, %rsp
	.cfi_def_cfa_offset 320
	movq	%fs:40, %rax
	movq	%rax, 248(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rdi
	movq	%rcx, 8(%rsp)
.LEHB0:
	call	cudaMallocManaged@PLT
.LEHE0:
	xorl	%edx, %edx
	leaq	16(%rsp), %rdi
	movl	$24, %esi
	movq	%rdx, 16(%rsp)
	movl	$1, %edx
.LEHB1:
	call	cudaMallocManaged@PLT
.LEHE1:
	xorl	%eax, %eax
	leaq	24(%rsp), %rdi
	movl	$1, %edx
	movl	$72, %esi
	movq	%rax, 24(%rsp)
.LEHB2:
	call	cudaMallocManaged@PLT
	movabsq	$4629700418002223104, %rax
	movss	.LC1(%rip), %xmm0
	leaq	104(%rsp), %rbx
	movq	%rax, 32(%rsp)
	movq	%rbx, %rdi
	movabsq	$4613487457626541261, %rax
	movq	%rax, 56(%rsp)
	movabsq	$4615288900055308698, %rax
	movq	%rax, 64(%rsp)
	movabsq	$4618891780624167731, %rax
	movq	%rax, 72(%rsp)
	movss	%xmm0, 40(%rsp)
	call	_Z14calc_transProbR9TransProbf
	movq	8(%rsp), %rdi
	leaq	32(%rsp), %rsi
	movl	$1, %ecx
	movl	$12, %edx
	call	cudaMemcpy@PLT
	movq	16(%rsp), %rdi
	leaq	56(%rsp), %rsi
	movl	$1, %ecx
	movl	$24, %edx
	call	cudaMemcpy@PLT
	movq	24(%rsp), %rdi
	movl	$1, %ecx
	movl	$72, %edx
	movq	%rbx, %rsi
	call	cudaMemcpy@PLT
	leaq	44(%rsp), %r14
	movq	8(%rsp), %rsi
	movl	$2, %ecx
	movl	$12, %edx
	movq	%r14, %rdi
	call	cudaMemcpy@PLT
	movq	16(%rsp), %rsi
	leaq	80(%rsp), %rdi
	movl	$2, %ecx
	movl	$24, %edx
	call	cudaMemcpy@PLT
	movq	24(%rsp), %rsi
	movl	$2, %ecx
	movl	$72, %edx
	leaq	176(%rsp), %r15
	movq	%r15, %rdi
	call	cudaMemcpy@PLT
	leaq	.LC2(%rip), %r13
	leaq	_ZSt4cout(%rip), %rbp
	movq	%r13, %rsi
	movq	%rbp, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	cvtss2sd	44(%rsp), %xmm0
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	leaq	.LC3(%rip), %rbx
	movq	%rax, %rdi
	movq	%rbx, %rsi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	cvtss2sd	48(%rsp), %xmm0
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	%rax, %rdi
	movq	%rbx, %rsi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	cvtss2sd	52(%rsp), %xmm0
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rbp, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	movq	%rax, %rdi
	leaq	.LC4(%rip), %rsi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	cvtss2sd	80(%rsp), %xmm0
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	%rax, %rdi
	movq	%rbx, %rsi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	cvtss2sd	84(%rsp), %xmm0
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	%rax, %rdi
	movq	%rbx, %rsi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	cvtss2sd	88(%rsp), %xmm0
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	%rax, %rdi
	movq	%rbx, %rsi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	cvtss2sd	92(%rsp), %xmm0
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	%rax, %rdi
	movq	%rbx, %rsi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	cvtss2sd	96(%rsp), %xmm0
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	%rax, %rdi
	movq	%rbx, %rsi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	cvtss2sd	100(%rsp), %xmm0
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	movq	$-4, %r12
.L8:
	movq	%rbp, %rdi
	cvtss2sd	16(%r15,%r12,4), %xmm0
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	%rax, %rdi
	movq	%rbx, %rsi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	addq	$4, %r12
	cmpq	$16, %r12
	jne	.L8
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rbp, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	movq	8(%rsp), %rsi
	movl	$2, %ecx
	movl	$12, %edx
	movq	%r14, %rdi
	movabsq	$4638707617261158400, %rax
	movq	%rax, (%rsi)
	movl	$0x40200000, 8(%rsi)
	call	cudaMemcpy@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rbp, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	movq	%rax, %rdi
	movq	%r13, %rsi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	cvtss2sd	44(%rsp), %xmm0
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	%rax, %rdi
	movq	%rbx, %rsi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	cvtss2sd	48(%rsp), %xmm0
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	%rax, %rdi
	movq	%rbx, %rsi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	cvtss2sd	52(%rsp), %xmm0
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
.LEHE2:
	movq	24(%rsp), %rdi
	testq	%rdi, %rdi
	je	.L9
	call	cudaFree@PLT
.L9:
	movq	16(%rsp), %rdi
	testq	%rdi, %rdi
	je	.L10
	call	cudaFree@PLT
.L10:
	movq	8(%rsp), %rdi
	testq	%rdi, %rdi
	je	.L11
	call	cudaFree@PLT
.L11:
	movq	248(%rsp), %rax
	subq	%fs:40, %rax
	je	.L19
	jmp	.L42
.L22:
	endbr64
	movq	%rax, %rbx
.L12:
	movq	24(%rsp), %rdi
	testq	%rdi, %rdi
	je	.L14
	call	cudaFree@PLT
	jmp	.L14
.L21:
	endbr64
	movq	%rax, %rbx
.L14:
	movq	16(%rsp), %rdi
	testq	%rdi, %rdi
	je	.L16
	call	cudaFree@PLT
	jmp	.L16
.L20:
	endbr64
	movq	%rax, %rbx
.L16:
	movq	8(%rsp), %rdi
	testq	%rdi, %rdi
	je	.L17
	call	cudaFree@PLT
.L17:
	movq	248(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L42
	movq	%rbx, %rdi
.LEHB3:
	call	_Unwind_Resume@PLT
.LEHE3:
.L42:
	call	__stack_chk_fail@PLT
.L19:
	addq	$264, %rsp
	.cfi_def_cfa_offset 56
	xorl	%eax, %eax
	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
	ret
	.cfi_endproc
.LFE4346:
	.globl	__gxx_personality_v0
	.section	.gcc_except_table,"a",@progbits
.LLSDA4346:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE4346-.LLSDACSB4346
.LLSDACSB4346:
	.uleb128 .LEHB0-.LFB4346
	.uleb128 .LEHE0-.LEHB0
	.uleb128 .L20-.LFB4346
	.uleb128 0
	.uleb128 .LEHB1-.LFB4346
	.uleb128 .LEHE1-.LEHB1
	.uleb128 .L21-.LFB4346
	.uleb128 0
	.uleb128 .LEHB2-.LFB4346
	.uleb128 .LEHE2-.LEHB2
	.uleb128 .L22-.LFB4346
	.uleb128 0
	.uleb128 .LEHB3-.LFB4346
	.uleb128 .LEHE3-.LEHB3
	.uleb128 0
	.uleb128 0
.LLSDACSE4346:
	.section	.text.startup
	.size	main, .-main
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB4396:
	.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
.LFE4396:
	.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
	.section	.rodata.cst16,"aM",@progbits,16
	.align 16
.LC0:
	.long	-2147483648
	.long	0
	.long	0
	.long	0
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC1:
	.long	1073741824
	.hidden	DW.ref.__gxx_personality_v0
	.weak	DW.ref.__gxx_personality_v0
	.section	.data.rel.local.DW.ref.__gxx_personality_v0,"awG",@progbits,DW.ref.__gxx_personality_v0,comdat
	.align 8
	.type	DW.ref.__gxx_personality_v0, @object
	.size	DW.ref.__gxx_personality_v0, 8
DW.ref.__gxx_personality_v0:
	.quad	__gxx_personality_v0
	.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:

/**
 * @file   : main_draft.cu
 * @brief  : main file draft for 2-dim. Ising in CUDA C++11/14, 
 * @details : 
 * 
 * @author : Ernest Yeung <ernestyalumni@gmail.com>
 * @date   : 20180108    
 * @ref    : M. Hjorth-Jensen, Computational Physics, University of Oslo (2015)
 * https://github.com/CompPhysics/ComputationalPhysics/blob/master/doc/Programs/LecturePrograms/programs/StatPhys/cpp/ising_2dim.cpp
 * https://www.paypal.com/cgi-bin/webscr?cmd=_donations&business=ernestsaveschristmas%2bpaypal%40gmail%2ecom&lc=US&item_name=ernestyalumni&currency_code=USD&bn=PP%2dDonationsBF%3abtn_donateCC_LG%2egif%3aNonHosted 
 * 
 * which won't go through a 3rd. party such as indiegogo, kickstarter, patreon.  
 * Otherwise, I receive emails and messages on how all my (free) material on 
 * physics, math, and engineering have helped students with their studies, 
 * and I know what it's like to not have money as a student, but love physics 
 * (or math, sciences, etc.), so I am committed to keeping all my material 
 * open-source and free, whether or not 
 * sufficiently crowdfunded, under the open-source MIT license: 
 * 	feel free to copy, edit, paste, make your own versions, share, use as you wish.  
 *  Just don't be an asshole and not give credit where credit is due.  
 * Peace out, never give up! -EY
 * 
 * */
/* 
 * COMPILATION TIP
 * nvcc main.cu ./devgrid2d/devgrid2d.cu -o main
 * 
 * */

#include <hip/hip_runtime.h>
#include <iostream>
#include <memory>  
#include <array> // std::array in TransProb struct
#include <cmath> // std::exp in TransProb


/* ********** device GPU structs ********** */

struct Sysparam {
	float E; // total energy E
	float M; // total magnetization H 
	float T; // temperature T of the system (it's kT; treat Boltzmann constant k as a unit conversion)  
}; 

struct Avg {
	// (data) members
	// average values of physical parameters  
	float Eavg; 
	float Mavg; 
	float Esq_avg; // Esq = E*E 
	float Msq_avg; // Msq = M*M
	float absM_avg; // absM = |M| 
	float M4_avg; // M4_avg = M*M*M*M
};

/** @struct TransProb
 *  @brief transition probabilities to new spin configuration for 2-dim. Ising model 
 * */
struct TransProb {
	// (data) members
	// transition probabilities data 
	std::array<float,17> transProb; 
	float J; // spin constant 
	
	// getting functions
	/** @fn get_by_DeltaE 
	 * @details given DeltaE (\Delta E), DeltaE = -8J, -4J,...8J, we want to get the 
	 * transition probability from std::unique_ptr transprob (but transprob indexed by 
	 * 0,1,...(17-1)
	 * */
	float get_by_DeltaE(const int DeltaE) {
		return transProb[DeltaE+8]; } ;
};

// custom deleters as structs   
struct del_Sysparam_struct { void operator()(Sysparam* ptr) { hipFree(ptr); } };
struct del_Avg_struct { void operator()(Avg* ptr) { hipFree(ptr); } };	
struct del_TransProb_struct { void operator()(TransProb* ptr) { hipFree(ptr); } };  	

/* ********** END of device GPU structs ********** */


/* ********** host CPU structs ********** */
struct h_Sysparam { 
	float E; // total energy E 
	float M; // total magnetization M 
	float T; // temperature T of the system 

	// constructors
	// default constructor
	/** @fn Sysparam()
	 * @brief default constructor for struct Sysparam 
	 * @details set all E,M,T parameters to 0
	 * */
	h_Sysparam() : E {0.f}, M {0.f}, T {0.f} {}; 

	/** @fn Sysparam(double, double,double)
	 * @brief constructor for struct Sysparam 
	 * */
	h_Sysparam(float E, float M, float T) : E {E}, M {M}, T {T} {} ;

	/** @fn Sysparam(double)
	 * @brief constructor for struct Sysparam, when only given the system temperature (initially)
	 * */
	h_Sysparam(float T) : E {0.f}, M {0.f}, T {T} {}; 	

};

struct h_Avg {
	// (data) members
	// average values of physical parameters  
	float Eavg; 
	float Mavg; 
	float Esq_avg; // Esq = E*E 
	float Msq_avg; // Msq = M*M
	float absM_avg; // absM = |M| 
	float M4_avg; // M4_avg = M*M*M*M
	
	// constructors
	// default constructor
	h_Avg() : Eavg {0.f}, Mavg(0.f), Esq_avg(0.f), Msq_avg{0.f}, absM_avg{0.f},M4_avg{0.f} {};

	h_Avg(float Eavg, float Mavg, float Esq_avg, float Msq_avg, float absM_avg, float M4_avg) : 
		Eavg {Eavg}, Mavg{Mavg}, Esq_avg{Esq_avg}, Msq_avg{Msq_avg}, absM_avg{absM_avg}, M4_avg{M4_avg}  { } ;

};


/** @struct TransProb
 *  @brief transition probabilities to new spin configuration for 2-dim. Ising model 
 * */
struct h_TransProb {
	// (data) members
	// transition probabilities data 
	std::array<float,17> transProb; 
	float J; // spin constant 
	
	// constructors
	// default constructor 
	h_TransProb() : J {1.f} { };

	h_TransProb(float J, h_Sysparam& sysparams) : J {J} {
		float T = sysparams.T; // temperature
		for (int de = -8; de<= 8; de+=4) { 
			transProb[de+8] = std::exp(-((float) de)/T); }
	}; 

	// getting functions
	/** @fn get_by_DeltaE 
	 * @details given DeltaE (\Delta E), DeltaE = -8J, -4J,...8J, we want to get the 
	 * transition probability from std::unique_ptr transprob (but transprob indexed by 
	 * 0,1,...(17-1)
	 * */
	float get_by_DeltaE(int DeltaE) {
		return transProb[DeltaE+8]; 
	}; 
};

/** @fn calc_transProb
 * 
 * */
void calc_transProb(TransProb & transProb, const float T) {
	for (int de = -8; de <= 8; de+=4) {
		transProb.transProb[de+8] = std::exp( -((float) de)/T); 
	}
}
 
/* ********** END of host CPU structs ********** */



int main(int argc, char* argv[]) 
{
	std::unique_ptr<Sysparam,del_Sysparam_struct> d_sysparams(nullptr, del_Sysparam_struct()); 
	hipMallocManaged((void **) &d_sysparams, 1*sizeof(Sysparam));  
	std::unique_ptr<Avg,del_Avg_struct> d_avgs(nullptr, del_Avg_struct()); 
	hipMallocManaged((void **) &d_avgs, 1*sizeof(Avg));  
	std::unique_ptr<TransProb,del_TransProb_struct> d_transProb(nullptr, del_TransProb_struct()); 
	hipMallocManaged((void **) &d_transProb, 1*sizeof(TransProb));  

	/* ****************************************************************************************** */
	/* ******************** Ways to initialize structs on device GPU; 2 ways ******************** */
	/* ****************************************************************************************** */
	/* *************** 1. cudaMemcpy from host to device *************** */

	// error: have to be of same type
//	h_Sysparam h_sysparams { 1.f, 3.f, 2.f };  
//	h_Avg h_avgs { 1.1f, 2.1f, 1.2f, 2.2f, 2.3f, 2.4f }; 
//	h_TransProb(1.f, h_sysparams);  

	// some host CPU values to input 
	Sysparam h_sysparams { 1.f, 3.f, 2.f };  
	Avg h_avgs { 1.1f, 2.1f, 1.2f, 2.2f, 2.3f, 2.4f }; 
	TransProb h_transProb; 
	calc_transProb(h_transProb, h_sysparams.T); 

	hipMemcpy( d_sysparams.get(), &h_sysparams, 1*sizeof(Sysparam), hipMemcpyHostToDevice);  // possible error have to be of same type
	hipMemcpy( d_avgs.get(), &h_avgs, 1*sizeof(Avg), hipMemcpyHostToDevice);  // possible error have to be of same type
	hipMemcpy( d_transProb.get(), &h_transProb, 1*sizeof(TransProb), hipMemcpyHostToDevice);  // possible error have to be of same type

	/* sanity check; copy back to host for sanity check */  
	Sysparam h_sysparams_out ;  
	Avg h_avgs_out ; 
	TransProb h_transProb_out; 

	hipMemcpy(&h_sysparams_out, d_sysparams.get(), 1*sizeof(Sysparam), hipMemcpyDeviceToHost);  // possible error have to be of same type
	hipMemcpy(&h_avgs_out, d_avgs.get(), 1*sizeof(Avg), hipMemcpyDeviceToHost);  // possible error have to be of same type
	hipMemcpy(&h_transProb_out, d_transProb.get(), 1*sizeof(TransProb), hipMemcpyDeviceToHost);  // possible error have to be of same type
	
	std::cout << " h_sysparams_out : " << h_sysparams_out.E << " " << h_sysparams_out.M << " " << h_sysparams_out.T << std::endl; 
	std::cout << std::endl << " h_avgs_out : " << h_avgs_out.Eavg << " " << h_avgs_out.Mavg 
		<< " " << h_avgs_out.Esq_avg << " " << h_avgs_out.Msq_avg << " " << h_avgs_out.absM_avg << " " << h_avgs_out.M4_avg  << std::endl; 
	for (int de =-8; de <= 8; de+=4) {
		std::cout << h_transProb_out.transProb[de+8] << " "; } std::cout << std::endl;
	/* END of sanity check */ 

	/* *************** 2. directly from host set values *************** */
	d_sysparams->E = 1.5f; 
	d_sysparams->M = 3.5f; 
	d_sysparams->T = 2.5f; 
	hipMemcpy(&h_sysparams_out, d_sysparams.get(), 1*sizeof(Sysparam), hipMemcpyDeviceToHost);  // possible error have to be of same type
	std::cout << std::endl << " h_sysparams_out : " << h_sysparams_out.E << " " << h_sysparams_out.M << " " << h_sysparams_out.T << std::endl; 
	
	
	
	
}

	.text
	.file	"main_draft.hip"
                                        # Start of file scope inline assembly
	.globl	_ZSt21ios_base_library_initv

                                        # End of file scope inline assembly
	.section	.rodata.cst16,"aM",@progbits,16
	.p2align	4, 0x0                          # -- Begin function _Z14calc_transProbR9TransProbf
.LCPI0_0:
	.long	0x80000000                      # float -0
	.long	0x80000000                      # float -0
	.long	0x80000000                      # float -0
	.long	0x80000000                      # float -0
	.text
	.globl	_Z14calc_transProbR9TransProbf
	.type	_Z14calc_transProbR9TransProbf,@function
_Z14calc_transProbR9TransProbf:         # @_Z14calc_transProbR9TransProbf
	.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	$16, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movss	%xmm0, 12(%rsp)                 # 4-byte Spill
	movq	%rdi, %rbx
	xorl	%r14d, %r14d
.LBB0_1:                                # =>This Inner Loop Header: Depth=1
	leaq	-8(%r14), %r15
	xorps	%xmm0, %xmm0
	cvtsi2ss	%r15d, %xmm0
	xorps	.LCPI0_0(%rip), %xmm0
	divss	12(%rsp), %xmm0                 # 4-byte Folded Reload
	callq	expf
	movss	%xmm0, (%rbx,%r14,4)
	addq	$4, %r14
	cmpq	$5, %r15
	jl	.LBB0_1
# %bb.2:
	addq	$16, %rsp
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z14calc_transProbR9TransProbf, .Lfunc_end0-_Z14calc_transProbR9TransProbf
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst16,"aM",@progbits,16
	.p2align	4, 0x0                          # -- Begin function main
.LCPI1_0:
	.long	0x3f8ccccd                      # float 1.10000002
	.long	0x40066666                      # float 2.0999999
	.long	0x3f99999a                      # float 1.20000005
	.long	0x400ccccd                      # float 2.20000005
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0
.LCPI1_1:
	.long	0xbf000000                      # float -0.5
	.text
	.globl	main
	.type	main,@function
main:                                   # @main
.Lfunc_begin0:
	.cfi_startproc
	.cfi_personality 3, __gxx_personality_v0
	.cfi_lsda 3, .Lexception0
# %bb.0:
	pushq	%r14
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	subq	$248, %rsp
	.cfi_def_cfa_offset 272
	.cfi_offset %rbx, -24
	.cfi_offset %r14, -16
	movq	%rsp, %rdi
	movq	$0, (%rdi)
.Ltmp0:
	movl	$12, %esi
	movl	$1, %edx
	callq	hipMallocManaged
.Ltmp1:
# %bb.1:
	leaq	32(%rsp), %rdi
	movq	$0, (%rdi)
.Ltmp3:
	movl	$24, %esi
	movl	$1, %edx
	callq	hipMallocManaged
.Ltmp4:
# %bb.2:
	leaq	24(%rsp), %rdi
	movq	$0, (%rdi)
.Ltmp6:
	movl	$72, %esi
	movl	$1, %edx
	callq	hipMallocManaged
.Ltmp7:
# %bb.3:
	movabsq	$4629700418002223104, %rax      # imm = 0x404000003F800000
	movq	%rax, 68(%rsp)
	movl	$1073741824, 76(%rsp)           # imm = 0x40000000
	movaps	.LCPI1_0(%rip), %xmm0           # xmm0 = [1.10000002E+0,2.0999999E+0,1.20000005E+0,2.20000005E+0]
	movups	%xmm0, 80(%rsp)
	movabsq	$4618891780624167731, %rax      # imm = 0x4019999A40133333
	movq	%rax, 96(%rsp)
	xorl	%ebx, %ebx
.LBB1_4:                                # =>This Inner Loop Header: Depth=1
	leaq	-8(%rbx), %r14
	xorps	%xmm0, %xmm0
	cvtsi2ss	%r14d, %xmm0
	mulss	.LCPI1_1(%rip), %xmm0
	callq	expf
	movss	%xmm0, 176(%rsp,%rbx,4)
	addq	$4, %rbx
	cmpq	$5, %r14
	jl	.LBB1_4
# %bb.5:                                # %_Z14calc_transProbR9TransProbf.exit
	movq	(%rsp), %rdi
.Ltmp9:
	leaq	68(%rsp), %rsi
	movl	$12, %edx
	movl	$1, %ecx
	callq	hipMemcpy
.Ltmp10:
# %bb.6:
	movq	32(%rsp), %rdi
.Ltmp11:
	leaq	80(%rsp), %rsi
	movl	$24, %edx
	movl	$1, %ecx
	callq	hipMemcpy
.Ltmp12:
# %bb.7:
	movq	24(%rsp), %rdi
.Ltmp13:
	leaq	176(%rsp), %rsi
	movl	$72, %edx
	movl	$1, %ecx
	callq	hipMemcpy
.Ltmp14:
# %bb.8:
	movq	(%rsp), %rsi
.Ltmp16:
	leaq	12(%rsp), %rdi
	movl	$12, %edx
	movl	$2, %ecx
	callq	hipMemcpy
.Ltmp17:
# %bb.9:
	movq	32(%rsp), %rsi
.Ltmp18:
	leaq	44(%rsp), %rdi
	movl	$24, %edx
	movl	$2, %ecx
	callq	hipMemcpy
.Ltmp19:
# %bb.10:
	movq	24(%rsp), %rsi
.Ltmp20:
	leaq	104(%rsp), %rdi
	movl	$72, %edx
	movl	$2, %ecx
	callq	hipMemcpy
.Ltmp21:
# %bb.11:
.Ltmp22:
	movl	$_ZSt4cout, %edi
	movl	$.L.str, %esi
	movl	$19, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp23:
# %bb.12:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit
	xorps	%xmm0, %xmm0
	cvtss2sd	12(%rsp), %xmm0
.Ltmp24:
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIdEERSoT_
.Ltmp25:
# %bb.13:                               # %_ZNSolsEf.exit
.Ltmp26:
	movq	%rax, %rbx
	movl	$.L.str.1, %esi
	movl	$1, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp27:
# %bb.14:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit16
	xorps	%xmm0, %xmm0
	cvtss2sd	16(%rsp), %xmm0
.Ltmp28:
	movq	%rbx, %rdi
	callq	_ZNSo9_M_insertIdEERSoT_
.Ltmp29:
# %bb.15:                               # %_ZNSolsEf.exit17
.Ltmp30:
	movq	%rax, %rbx
	movl	$.L.str.1, %esi
	movl	$1, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp31:
# %bb.16:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit18
	xorps	%xmm0, %xmm0
	cvtss2sd	20(%rsp), %xmm0
.Ltmp32:
	movq	%rbx, %rdi
	callq	_ZNSo9_M_insertIdEERSoT_
.Ltmp33:
# %bb.17:                               # %_ZNSolsEf.exit19
	movq	%rax, %rbx
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
.Ltmp34:
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
.Ltmp35:
# %bb.18:                               # %.noexc
.Ltmp36:
	movsbl	%al, %esi
	movq	%rbx, %rdi
	callq	_ZNSo3putEc
.Ltmp37:
# %bb.19:                               # %.noexc45
.Ltmp38:
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
.Ltmp39:
# %bb.20:                               # %_ZNSolsEPFRSoS_E.exit
	movq	_ZSt4cout(%rip), %rax
	movl	$_ZSt4cout, %edi
	addq	-24(%rax), %rdi
.Ltmp40:
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
.Ltmp41:
# %bb.21:                               # %.noexc47
.Ltmp42:
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
.Ltmp43:
# %bb.22:                               # %.noexc48
.Ltmp44:
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
.Ltmp45:
# %bb.23:                               # %_ZNSolsEPFRSoS_E.exit20
.Ltmp46:
	movq	%rax, %rbx
	movl	$.L.str.2, %esi
	movl	$14, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp47:
# %bb.24:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit21
	xorps	%xmm0, %xmm0
	cvtss2sd	44(%rsp), %xmm0
.Ltmp48:
	movq	%rbx, %rdi
	callq	_ZNSo9_M_insertIdEERSoT_
.Ltmp49:
# %bb.25:                               # %_ZNSolsEf.exit22
.Ltmp50:
	movq	%rax, %rbx
	movl	$.L.str.1, %esi
	movl	$1, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp51:
# %bb.26:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit23
	xorps	%xmm0, %xmm0
	cvtss2sd	48(%rsp), %xmm0
.Ltmp52:
	movq	%rbx, %rdi
	callq	_ZNSo9_M_insertIdEERSoT_
.Ltmp53:
# %bb.27:                               # %_ZNSolsEf.exit24
.Ltmp54:
	movq	%rax, %rbx
	movl	$.L.str.1, %esi
	movl	$1, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp55:
# %bb.28:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit25
	xorps	%xmm0, %xmm0
	cvtss2sd	52(%rsp), %xmm0
.Ltmp56:
	movq	%rbx, %rdi
	callq	_ZNSo9_M_insertIdEERSoT_
.Ltmp57:
# %bb.29:                               # %_ZNSolsEf.exit26
.Ltmp58:
	movq	%rax, %rbx
	movl	$.L.str.1, %esi
	movl	$1, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp59:
# %bb.30:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit27
	xorps	%xmm0, %xmm0
	cvtss2sd	56(%rsp), %xmm0
.Ltmp60:
	movq	%rbx, %rdi
	callq	_ZNSo9_M_insertIdEERSoT_
.Ltmp61:
# %bb.31:                               # %_ZNSolsEf.exit28
.Ltmp62:
	movq	%rax, %rbx
	movl	$.L.str.1, %esi
	movl	$1, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp63:
# %bb.32:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit29
	xorps	%xmm0, %xmm0
	cvtss2sd	60(%rsp), %xmm0
.Ltmp64:
	movq	%rbx, %rdi
	callq	_ZNSo9_M_insertIdEERSoT_
.Ltmp65:
# %bb.33:                               # %_ZNSolsEf.exit30
.Ltmp66:
	movq	%rax, %rbx
	movl	$.L.str.1, %esi
	movl	$1, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp67:
# %bb.34:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit31
	xorps	%xmm0, %xmm0
	cvtss2sd	64(%rsp), %xmm0
.Ltmp68:
	movq	%rbx, %rdi
	callq	_ZNSo9_M_insertIdEERSoT_
.Ltmp69:
# %bb.35:                               # %_ZNSolsEf.exit32
	movq	%rax, %rbx
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
.Ltmp70:
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
.Ltmp71:
# %bb.36:                               # %.noexc51
.Ltmp72:
	movsbl	%al, %esi
	movq	%rbx, %rdi
	callq	_ZNSo3putEc
.Ltmp73:
# %bb.37:                               # %.noexc52
.Ltmp74:
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
.Ltmp75:
# %bb.38:                               # %_ZNSolsEPFRSoS_E.exit33.preheader.preheader
	movq	$-12, %rbx
.LBB1_39:                               # %_ZNSolsEPFRSoS_E.exit33.preheader
                                        # =>This Inner Loop Header: Depth=1
	xorps	%xmm0, %xmm0
	cvtss2sd	152(%rsp,%rbx,4), %xmm0
.Ltmp76:
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIdEERSoT_
.Ltmp77:
# %bb.40:                               # %_ZNSolsEf.exit35
                                        #   in Loop: Header=BB1_39 Depth=1
.Ltmp78:
	movl	$.L.str.1, %esi
	movl	$1, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp79:
# %bb.41:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit36
                                        #   in Loop: Header=BB1_39 Depth=1
	addq	$4, %rbx
	cmpq	$5, %rbx
	jl	.LBB1_39
# %bb.42:
	movq	_ZSt4cout(%rip), %rax
	movl	$_ZSt4cout, %edi
	addq	-24(%rax), %rdi
.Ltmp81:
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
.Ltmp82:
# %bb.43:                               # %.noexc55
.Ltmp83:
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
.Ltmp84:
# %bb.44:                               # %.noexc56
.Ltmp85:
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
.Ltmp86:
# %bb.45:                               # %_ZNSolsEPFRSoS_E.exit34
	movq	(%rsp), %rsi
	movabsq	$4638707617261158400, %rax      # imm = 0x406000003FC00000
	movq	%rax, (%rsi)
	movl	$1075838976, 8(%rsi)            # imm = 0x40200000
.Ltmp87:
	leaq	12(%rsp), %rdi
	movl	$12, %edx
	movl	$2, %ecx
	callq	hipMemcpy
.Ltmp88:
# %bb.46:
	movq	_ZSt4cout(%rip), %rax
	movl	$_ZSt4cout, %edi
	addq	-24(%rax), %rdi
.Ltmp89:
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
.Ltmp90:
# %bb.47:                               # %.noexc59
.Ltmp91:
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
.Ltmp92:
# %bb.48:                               # %.noexc60
.Ltmp93:
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
.Ltmp94:
# %bb.49:                               # %_ZNSolsEPFRSoS_E.exit37
.Ltmp95:
	movq	%rax, %rbx
	movl	$.L.str, %esi
	movl	$19, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp96:
# %bb.50:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit38
	xorps	%xmm0, %xmm0
	cvtss2sd	12(%rsp), %xmm0
.Ltmp97:
	movq	%rbx, %rdi
	callq	_ZNSo9_M_insertIdEERSoT_
.Ltmp98:
# %bb.51:                               # %_ZNSolsEf.exit39
.Ltmp99:
	movq	%rax, %rbx
	movl	$.L.str.1, %esi
	movl	$1, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp100:
# %bb.52:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit40
	xorps	%xmm0, %xmm0
	cvtss2sd	16(%rsp), %xmm0
.Ltmp101:
	movq	%rbx, %rdi
	callq	_ZNSo9_M_insertIdEERSoT_
.Ltmp102:
# %bb.53:                               # %_ZNSolsEf.exit41
.Ltmp103:
	movq	%rax, %rbx
	movl	$.L.str.1, %esi
	movl	$1, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp104:
# %bb.54:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit42
	xorps	%xmm0, %xmm0
	cvtss2sd	20(%rsp), %xmm0
.Ltmp105:
	movq	%rbx, %rdi
	callq	_ZNSo9_M_insertIdEERSoT_
.Ltmp106:
# %bb.55:                               # %_ZNSolsEf.exit43
	movq	%rax, %rbx
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
.Ltmp107:
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
.Ltmp108:
# %bb.56:                               # %.noexc63
.Ltmp109:
	movsbl	%al, %esi
	movq	%rbx, %rdi
	callq	_ZNSo3putEc
.Ltmp110:
# %bb.57:                               # %.noexc64
.Ltmp111:
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
.Ltmp112:
# %bb.58:                               # %_ZNSolsEPFRSoS_E.exit44
	leaq	24(%rsp), %rdi
	callq	_ZNSt10unique_ptrI9TransProb20del_TransProb_structED2Ev
	leaq	32(%rsp), %rdi
	callq	_ZNSt10unique_ptrI3Avg14del_Avg_structED2Ev
	movq	%rsp, %rdi
	callq	_ZNSt10unique_ptrI8Sysparam19del_Sysparam_structED2Ev
	xorl	%eax, %eax
	addq	$248, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	retq
.LBB1_61:
	.cfi_def_cfa_offset 272
.Ltmp8:
	jmp	.LBB1_65
.LBB1_60:
.Ltmp5:
	movq	%rax, %rbx
	jmp	.LBB1_66
.LBB1_59:
.Ltmp2:
	movq	%rax, %rbx
	jmp	.LBB1_67
.LBB1_62:
.Ltmp15:
	jmp	.LBB1_65
.LBB1_64:
.Ltmp113:
	jmp	.LBB1_65
.LBB1_63:
.Ltmp80:
.LBB1_65:
	movq	%rax, %rbx
	leaq	24(%rsp), %rdi
	callq	_ZNSt10unique_ptrI9TransProb20del_TransProb_structED2Ev
.LBB1_66:
	leaq	32(%rsp), %rdi
	callq	_ZNSt10unique_ptrI3Avg14del_Avg_structED2Ev
.LBB1_67:
	movq	%rsp, %rdi
	callq	_ZNSt10unique_ptrI8Sysparam19del_Sysparam_structED2Ev
	movq	%rbx, %rdi
	callq	_Unwind_Resume@PLT
.Lfunc_end1:
	.size	main, .Lfunc_end1-main
	.cfi_endproc
	.section	.gcc_except_table,"a",@progbits
	.p2align	2, 0x0
GCC_except_table1:
.Lexception0:
	.byte	255                             # @LPStart Encoding = omit
	.byte	255                             # @TType Encoding = omit
	.byte	1                               # Call site Encoding = uleb128
	.uleb128 .Lcst_end0-.Lcst_begin0
.Lcst_begin0:
	.uleb128 .Ltmp0-.Lfunc_begin0           # >> Call Site 1 <<
	.uleb128 .Ltmp1-.Ltmp0                  #   Call between .Ltmp0 and .Ltmp1
	.uleb128 .Ltmp2-.Lfunc_begin0           #     jumps to .Ltmp2
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp3-.Lfunc_begin0           # >> Call Site 2 <<
	.uleb128 .Ltmp4-.Ltmp3                  #   Call between .Ltmp3 and .Ltmp4
	.uleb128 .Ltmp5-.Lfunc_begin0           #     jumps to .Ltmp5
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp6-.Lfunc_begin0           # >> Call Site 3 <<
	.uleb128 .Ltmp7-.Ltmp6                  #   Call between .Ltmp6 and .Ltmp7
	.uleb128 .Ltmp8-.Lfunc_begin0           #     jumps to .Ltmp8
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp9-.Lfunc_begin0           # >> Call Site 4 <<
	.uleb128 .Ltmp14-.Ltmp9                 #   Call between .Ltmp9 and .Ltmp14
	.uleb128 .Ltmp15-.Lfunc_begin0          #     jumps to .Ltmp15
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp16-.Lfunc_begin0          # >> Call Site 5 <<
	.uleb128 .Ltmp75-.Ltmp16                #   Call between .Ltmp16 and .Ltmp75
	.uleb128 .Ltmp113-.Lfunc_begin0         #     jumps to .Ltmp113
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp76-.Lfunc_begin0          # >> Call Site 6 <<
	.uleb128 .Ltmp79-.Ltmp76                #   Call between .Ltmp76 and .Ltmp79
	.uleb128 .Ltmp80-.Lfunc_begin0          #     jumps to .Ltmp80
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp81-.Lfunc_begin0          # >> Call Site 7 <<
	.uleb128 .Ltmp112-.Ltmp81               #   Call between .Ltmp81 and .Ltmp112
	.uleb128 .Ltmp113-.Lfunc_begin0         #     jumps to .Ltmp113
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp112-.Lfunc_begin0         # >> Call Site 8 <<
	.uleb128 .Lfunc_end1-.Ltmp112           #   Call between .Ltmp112 and .Lfunc_end1
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
.Lcst_end0:
	.p2align	2, 0x0
                                        # -- End function
	.section	.text._ZNSt10unique_ptrI9TransProb20del_TransProb_structED2Ev,"axG",@progbits,_ZNSt10unique_ptrI9TransProb20del_TransProb_structED2Ev,comdat
	.weak	_ZNSt10unique_ptrI9TransProb20del_TransProb_structED2Ev # -- Begin function _ZNSt10unique_ptrI9TransProb20del_TransProb_structED2Ev
	.p2align	1, 0x90
	.type	_ZNSt10unique_ptrI9TransProb20del_TransProb_structED2Ev,@function
_ZNSt10unique_ptrI9TransProb20del_TransProb_structED2Ev: # @_ZNSt10unique_ptrI9TransProb20del_TransProb_structED2Ev
.Lfunc_begin1:
	.cfi_startproc
	.cfi_personality 3, __gxx_personality_v0
	.cfi_lsda 3, .Lexception1
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset %rbx, -16
	movq	%rdi, %rbx
	movq	(%rdi), %rdi
	testq	%rdi, %rdi
	je	.LBB2_2
# %bb.1:
.Ltmp114:
	callq	hipFree
.Ltmp115:
.LBB2_2:                                # %_ZN20del_TransProb_structclEP9TransProb.exit
	movq	$0, (%rbx)
	popq	%rbx
	.cfi_def_cfa_offset 8
	retq
.LBB2_3:
	.cfi_def_cfa_offset 16
.Ltmp116:
	movq	%rax, %rdi
	callq	__clang_call_terminate
.Lfunc_end2:
	.size	_ZNSt10unique_ptrI9TransProb20del_TransProb_structED2Ev, .Lfunc_end2-_ZNSt10unique_ptrI9TransProb20del_TransProb_structED2Ev
	.cfi_endproc
	.section	.gcc_except_table._ZNSt10unique_ptrI9TransProb20del_TransProb_structED2Ev,"aG",@progbits,_ZNSt10unique_ptrI9TransProb20del_TransProb_structED2Ev,comdat
	.p2align	2, 0x0
GCC_except_table2:
.Lexception1:
	.byte	255                             # @LPStart Encoding = omit
	.byte	3                               # @TType Encoding = udata4
	.uleb128 .Lttbase0-.Lttbaseref0
.Lttbaseref0:
	.byte	1                               # Call site Encoding = uleb128
	.uleb128 .Lcst_end1-.Lcst_begin1
.Lcst_begin1:
	.uleb128 .Ltmp114-.Lfunc_begin1         # >> Call Site 1 <<
	.uleb128 .Ltmp115-.Ltmp114              #   Call between .Ltmp114 and .Ltmp115
	.uleb128 .Ltmp116-.Lfunc_begin1         #     jumps to .Ltmp116
	.byte	1                               #   On action: 1
.Lcst_end1:
	.byte	1                               # >> Action Record 1 <<
                                        #   Catch TypeInfo 1
	.byte	0                               #   No further actions
	.p2align	2, 0x0
                                        # >> Catch TypeInfos <<
	.long	0                               # TypeInfo 1
.Lttbase0:
	.p2align	2, 0x0
                                        # -- End function
	.section	.text._ZNSt10unique_ptrI3Avg14del_Avg_structED2Ev,"axG",@progbits,_ZNSt10unique_ptrI3Avg14del_Avg_structED2Ev,comdat
	.weak	_ZNSt10unique_ptrI3Avg14del_Avg_structED2Ev # -- Begin function _ZNSt10unique_ptrI3Avg14del_Avg_structED2Ev
	.p2align	1, 0x90
	.type	_ZNSt10unique_ptrI3Avg14del_Avg_structED2Ev,@function
_ZNSt10unique_ptrI3Avg14del_Avg_structED2Ev: # @_ZNSt10unique_ptrI3Avg14del_Avg_structED2Ev
.Lfunc_begin2:
	.cfi_startproc
	.cfi_personality 3, __gxx_personality_v0
	.cfi_lsda 3, .Lexception2
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset %rbx, -16
	movq	%rdi, %rbx
	movq	(%rdi), %rdi
	testq	%rdi, %rdi
	je	.LBB3_2
# %bb.1:
.Ltmp117:
	callq	hipFree
.Ltmp118:
.LBB3_2:                                # %_ZN14del_Avg_structclEP3Avg.exit
	movq	$0, (%rbx)
	popq	%rbx
	.cfi_def_cfa_offset 8
	retq
.LBB3_3:
	.cfi_def_cfa_offset 16
.Ltmp119:
	movq	%rax, %rdi
	callq	__clang_call_terminate
.Lfunc_end3:
	.size	_ZNSt10unique_ptrI3Avg14del_Avg_structED2Ev, .Lfunc_end3-_ZNSt10unique_ptrI3Avg14del_Avg_structED2Ev
	.cfi_endproc
	.section	.gcc_except_table._ZNSt10unique_ptrI3Avg14del_Avg_structED2Ev,"aG",@progbits,_ZNSt10unique_ptrI3Avg14del_Avg_structED2Ev,comdat
	.p2align	2, 0x0
GCC_except_table3:
.Lexception2:
	.byte	255                             # @LPStart Encoding = omit
	.byte	3                               # @TType Encoding = udata4
	.uleb128 .Lttbase1-.Lttbaseref1
.Lttbaseref1:
	.byte	1                               # Call site Encoding = uleb128
	.uleb128 .Lcst_end2-.Lcst_begin2
.Lcst_begin2:
	.uleb128 .Ltmp117-.Lfunc_begin2         # >> Call Site 1 <<
	.uleb128 .Ltmp118-.Ltmp117              #   Call between .Ltmp117 and .Ltmp118
	.uleb128 .Ltmp119-.Lfunc_begin2         #     jumps to .Ltmp119
	.byte	1                               #   On action: 1
.Lcst_end2:
	.byte	1                               # >> Action Record 1 <<
                                        #   Catch TypeInfo 1
	.byte	0                               #   No further actions
	.p2align	2, 0x0
                                        # >> Catch TypeInfos <<
	.long	0                               # TypeInfo 1
.Lttbase1:
	.p2align	2, 0x0
                                        # -- End function
	.section	.text._ZNSt10unique_ptrI8Sysparam19del_Sysparam_structED2Ev,"axG",@progbits,_ZNSt10unique_ptrI8Sysparam19del_Sysparam_structED2Ev,comdat
	.weak	_ZNSt10unique_ptrI8Sysparam19del_Sysparam_structED2Ev # -- Begin function _ZNSt10unique_ptrI8Sysparam19del_Sysparam_structED2Ev
	.p2align	1, 0x90
	.type	_ZNSt10unique_ptrI8Sysparam19del_Sysparam_structED2Ev,@function
_ZNSt10unique_ptrI8Sysparam19del_Sysparam_structED2Ev: # @_ZNSt10unique_ptrI8Sysparam19del_Sysparam_structED2Ev
.Lfunc_begin3:
	.cfi_startproc
	.cfi_personality 3, __gxx_personality_v0
	.cfi_lsda 3, .Lexception3
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset %rbx, -16
	movq	%rdi, %rbx
	movq	(%rdi), %rdi
	testq	%rdi, %rdi
	je	.LBB4_2
# %bb.1:
.Ltmp120:
	callq	hipFree
.Ltmp121:
.LBB4_2:                                # %_ZN19del_Sysparam_structclEP8Sysparam.exit
	movq	$0, (%rbx)
	popq	%rbx
	.cfi_def_cfa_offset 8
	retq
.LBB4_3:
	.cfi_def_cfa_offset 16
.Ltmp122:
	movq	%rax, %rdi
	callq	__clang_call_terminate
.Lfunc_end4:
	.size	_ZNSt10unique_ptrI8Sysparam19del_Sysparam_structED2Ev, .Lfunc_end4-_ZNSt10unique_ptrI8Sysparam19del_Sysparam_structED2Ev
	.cfi_endproc
	.section	.gcc_except_table._ZNSt10unique_ptrI8Sysparam19del_Sysparam_structED2Ev,"aG",@progbits,_ZNSt10unique_ptrI8Sysparam19del_Sysparam_structED2Ev,comdat
	.p2align	2, 0x0
GCC_except_table4:
.Lexception3:
	.byte	255                             # @LPStart Encoding = omit
	.byte	3                               # @TType Encoding = udata4
	.uleb128 .Lttbase2-.Lttbaseref2
.Lttbaseref2:
	.byte	1                               # Call site Encoding = uleb128
	.uleb128 .Lcst_end3-.Lcst_begin3
.Lcst_begin3:
	.uleb128 .Ltmp120-.Lfunc_begin3         # >> Call Site 1 <<
	.uleb128 .Ltmp121-.Ltmp120              #   Call between .Ltmp120 and .Ltmp121
	.uleb128 .Ltmp122-.Lfunc_begin3         #     jumps to .Ltmp122
	.byte	1                               #   On action: 1
.Lcst_end3:
	.byte	1                               # >> Action Record 1 <<
                                        #   Catch TypeInfo 1
	.byte	0                               #   No further actions
	.p2align	2, 0x0
                                        # >> Catch TypeInfos <<
	.long	0                               # TypeInfo 1
.Lttbase2:
	.p2align	2, 0x0
                                        # -- End function
	.section	.text.__clang_call_terminate,"axG",@progbits,__clang_call_terminate,comdat
	.hidden	__clang_call_terminate          # -- Begin function __clang_call_terminate
	.weak	__clang_call_terminate
	.type	__clang_call_terminate,@function
__clang_call_terminate:                 # @__clang_call_terminate
	.cfi_startproc
# %bb.0:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__cxa_begin_catch
	callq	_ZSt9terminatev
.Lfunc_end5:
	.size	__clang_call_terminate, .Lfunc_end5-__clang_call_terminate
	.cfi_endproc
                                        # -- End function
	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	" h_sysparams_out : "
	.size	.L.str, 20

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	" "
	.size	.L.str.1, 2

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	" h_avgs_out : "
	.size	.L.str.2, 15

	.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 __gxx_personality_v0
	.addrsig_sym _Unwind_Resume
	.addrsig_sym _ZSt4cout
	.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
