#ifndef __CUDA_RUNTIME_H__
#include "cuda_runtime.h"
#endif // !"cuda_runtime.h"

#ifndef __DEVICE_LAUNCH_PARAMETERS_H__
#include "device_launch_parameters.h"
#endif // !__DEVICE_LAUNCH_PARAMETERS_H__

#include <stdio.h>
#include <string>
#include <iostream>
#include <chrono>
#include <ctime>
#include <vector>

#include "main.cuh"

using namespace std;

// Clear screen helper
void cls() {
	std::cout << "\033[2J\033[1;1H";
}

int main()
{
	cudaError_t cudaStatus;

	// Store pointers to GPU memory locally
	int* cflag_ptr = getFlagAddr();
	int* board_ptr = getBoardAddr();
	int MAX_N = getMaxN();
	
	// Program start time
	auto global_start = chrono::system_clock::now(); 

	// Initialize GPU
	cudaStatus = cudaSetDevice(0);
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
		goto Error;
	}

	// Loop until user issues quit command
	while (1) {

		cout << "Interactive GPU-Accelerated N-Queens Solver" << endl;
		cout << "Please select an option: " << endl;
		cout << "1 - Solve for N" << endl;
		cout << "2 - Solve for range of N" << endl;
		cout << "3 - Quit" << endl;

		int resp = 0;
		char _;

		cin >> resp;

		// Quit out of program
		if (resp == 3) { 
			break;
		}

		// Solve for range of N
		else if (resp == 2) { 

			int lower, upper;

			cout << "Enter lower bound (4 <= N <= " << MAX_N << "): ";
			cin >> lower;
			cout << "Enter upper bound (" << lower << " < N <= " << MAX_N << "): ";
			cin >> upper;
			cls();

			if (lower >= upper || lower < 4 || upper > MAX_N) {
				continue;
			}

			global_start = chrono::system_clock::now();

			cudaStatus = rangeSolve(lower, upper, cflag_ptr, board_ptr);

			auto global_end = chrono::system_clock::now();
			chrono::duration<double> elapsed_seconds = (global_end - global_start);
			cout << "Total exec time (s): " << elapsed_seconds.count() << endl;

			cout << endl;
			cout << "Press any key to continue." << endl;
			cin >> _;
			cls();

		}

		// Solve for single value of N
		else if (resp == 1) { 

			int Nq = 0;

			cout << "Enter number of queens (4 <= N <= " << MAX_N << "): ";
			cin >> Nq;
			cls();

			if (Nq < 4 || Nq > MAX_N) {
				continue;
			}

			global_start = chrono::system_clock::now();

			cudaStatus = singleSolve(Nq, cflag_ptr, board_ptr);

			// Display total execution time
			auto global_end = chrono::system_clock::now();
			chrono::duration<double> elapsed_mseconds = 1000*(global_end - global_start);
			cout << "Total exec time (ms): " << elapsed_mseconds.count() << endl;

			cout << "Press any key to continue." << endl;
			cin >> _;
			cls();
		}
	}

	// Free up all GPU memory
Error:
	memPurge();

	// Ensure no errors on the status flag
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "queens died :(");
		return 1;
	}

	// cudaDeviceReset must be called before exiting
	cudaStatus = cudaDeviceReset();
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "cudaDeviceReset failed!");
		return 1;
	}

	return 0;
}
	.file	"tmpxft_002ac3a0_00000000-6_main.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB4105:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE4105:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.align 2
	.type	_ZNSolsEPFRSoS_E.isra.0, @function
_ZNSolsEPFRSoS_E.isra.0:
.LFB4812:
	.cfi_startproc
	jmp	*%rsi
	.cfi_endproc
.LFE4812:
	.size	_ZNSolsEPFRSoS_E.isra.0, .-_ZNSolsEPFRSoS_E.isra.0
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"\033[2J\033[1;1H"
	.text
	.globl	_Z3clsv
	.type	_Z3clsv, @function
_Z3clsv:
.LFB4097:
	.cfi_startproc
	endbr64
	leaq	.LC0(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	jmp	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	.cfi_endproc
.LFE4097:
	.size	_Z3clsv, .-_Z3clsv
	.section	.rodata.str1.1
.LC1:
	.string	"cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?"
.LC2:
	.string	"Interactive GPU-Accelerated N-Queens Solver"
.LC3:
	.string	"Please select an option: "
.LC4:
	.string	"1 - Solve for N"
.LC5:
	.string	"2 - Solve for range of N"
.LC6:
	.string	"3 - Quit"
.LC7:
	.string	"Enter lower bound (4 <= N <= "
.LC8:
	.string	"): "
.LC9:
	.string	"Enter upper bound ("
.LC10:
	.string	" < N <= "
.LC12:
	.string	"Total exec time (s): "
.LC13:
	.string	"Press any key to continue."
.LC14:
	.string	"Enter number of queens (4 <= N <= "
.LC15:
	.string	"Total exec time (ms): "
.LC16:
	.string	"queens died :("
.LC17:
	.string	"cudaDeviceReset failed!"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB4098:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	leaq	.LC2(%rip), %r15
	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	$56, %rsp
	.cfi_def_cfa_offset 112
	movq	%fs:40, %rax
	movq	%rax, 40(%rsp)
	xorl	%eax, %eax
	call	_Z11getFlagAddrv@PLT
	movq	%rax, %r12
	call	_Z12getBoardAddrv@PLT
	movq	%rax, %r13
	call	_Z7getMaxNv@PLT
	movl	%eax, %ebp
	call	_ZNSt6chrono3_V212system_clock3nowEv@PLT
	xorl	%edi, %edi
	call	cudaSetDevice@PLT
	movl	%eax, %ebx
	testl	%eax, %eax
	je	.L5
	movq	stderr(%rip), %rdi
	leaq	.LC1(%rip), %rdx
	movl	$2, %esi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
	call	_Z8memPurgev@PLT
	jmp	.L6
.L30:
	cmpl	$2, %eax
	jne	.L8
	leaq	.LC7(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%ebp, %esi
	leaq	.LC8(%rip), %r14
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	movq	%r14, %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	leaq	32(%rsp), %rsi
	leaq	_ZSt3cin(%rip), %rdi
	call	_ZNSirsERi@PLT
	leaq	.LC9(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	32(%rsp), %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	leaq	.LC10(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%ebp, %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	movq	%r14, %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	leaq	36(%rsp), %rsi
	leaq	_ZSt3cin(%rip), %rdi
	call	_ZNSirsERi@PLT
	call	_Z3clsv
	movl	32(%rsp), %ecx
	movl	36(%rsp), %eax
	cmpl	$3, %ecx
	setle	%dl
	cmpl	%eax, %ecx
	setge	%cl
	orb	%cl, %dl
	je	.L29
.L5:
	movq	%r15, %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	.LC3(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	.LC4(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	.LC5(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	.LC6(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	xorl	%edx, %edx
	leaq	28(%rsp), %rsi
	leaq	_ZSt3cin(%rip), %rdi
	movl	%edx, 28(%rsp)
	call	_ZNSirsERi@PLT
	movl	28(%rsp), %eax
	cmpl	$3, %eax
	jne	.L30
	jmp	.L7
.L29:
	cmpl	%ebp, %eax
	jg	.L5
	call	_ZNSt6chrono3_V212system_clock3nowEv@PLT
	movl	36(%rsp), %esi
	movl	32(%rsp), %edi
	movq	%r13, %rcx
	movq	%r12, %rdx
	movq	%rax, %r14
	call	_Z10rangeSolveiiPiS_@PLT
	movl	%eax, %ebx
	call	_ZNSt6chrono3_V212system_clock3nowEv@PLT
	leaq	.LC12(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	subq	%r14, %rax
	cvtsi2sdq	%rax, %xmm0
	divsd	.LC11(%rip), %xmm0
	movsd	%xmm0, 8(%rsp)
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movsd	8(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	jmp	.L26
.L8:
	decl	%eax
	jne	.L5
	xorl	%eax, %eax
	leaq	.LC14(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	movl	%eax, 36(%rsp)
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%ebp, %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	leaq	.LC8(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	leaq	36(%rsp), %rsi
	leaq	_ZSt3cin(%rip), %rdi
	call	_ZNSirsERi@PLT
	call	_Z3clsv
	movl	36(%rsp), %eax
	cmpl	%ebp, %eax
	jg	.L5
	cmpl	$3, %eax
	jle	.L5
	call	_ZNSt6chrono3_V212system_clock3nowEv@PLT
	movl	36(%rsp), %edi
	movq	%r13, %rdx
	movq	%r12, %rsi
	movq	%rax, %r14
	call	_Z11singleSolveiPiS_@PLT
	movl	%eax, %ebx
	call	_ZNSt6chrono3_V212system_clock3nowEv@PLT
	leaq	.LC15(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	subq	%r14, %rax
	imulq	$1000, %rax, %rax
	cvtsi2sdq	%rax, %xmm0
	divsd	.LC11(%rip), %xmm0
	movsd	%xmm0, 8(%rsp)
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movsd	8(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
.L26:
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	.LC13(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	27(%rsp), %rsi
	leaq	_ZSt3cin(%rip), %rdi
	call	_ZStrsIcSt11char_traitsIcEERSt13basic_istreamIT_T0_ES6_RS3_@PLT
	call	_Z3clsv
	jmp	.L5
.L7:
	call	_Z8memPurgev@PLT
	testl	%ebx, %ebx
	je	.L24
.L6:
	movq	stderr(%rip), %rdi
	leaq	.LC16(%rip), %rdx
	movl	$2, %esi
	xorl	%eax, %eax
	jmp	.L27
.L24:
	call	cudaDeviceReset@PLT
	movl	%eax, %edx
	xorl	%eax, %eax
	testl	%edx, %edx
	je	.L4
	movq	stderr(%rip), %rdi
	leaq	.LC17(%rip), %rdx
	movl	$2, %esi
.L27:
	call	__fprintf_chk@PLT
	movl	$1, %eax
.L4:
	movq	40(%rsp), %rdx
	subq	%fs:40, %rdx
	je	.L19
	call	__stack_chk_fail@PLT
.L19:
	addq	$56, %rsp
	.cfi_def_cfa_offset 56
	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
.LFE4098:
	.size	main, .-main
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB4128:
	.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
.LFE4128:
	.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.cst8,"aM",@progbits,8
	.align 8
.LC11:
	.long	0
	.long	1104006501
	.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 __CUDA_RUNTIME_H__
#include "hip/hip_runtime.h"
#endif // !"cuda_runtime.h"

#ifndef __DEVICE_LAUNCH_PARAMETERS_H__

#endif // !__DEVICE_LAUNCH_PARAMETERS_H__

using namespace std;

// GPU functions
__device__ void register_q(int x, int y, int num_queens);
__device__ void case1(int i, int N);
__device__ void case2(int i, int N);
__global__ void N_Queens_Kernel(int num_queens);
__global__ void clearBuffers(int num_queens);

// Host interface functions
void memPurge();
int* getBoardAddr();
int* getFlagAddr();
int getMaxN();

// Solver functions
hipError_t singleSolve(int Nq, int* cflag_ptr, int* board_ptr);
hipError_t rangeSolve(int lower, int upper, int* cflag_ptr, int* board_ptr);
	.text
	.file	"main.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
