#include "segment.cuh"

// helper function for determining if t in [a, b] where order of a, b is unknown

__device__ bool isBetween(const double &t, const double &a, const double &b) {
	return (a <= t && t <= b) || (b <= t && t <= a);
}

__device__ Segment::Segment() {}

Segment::Segment(Point *a, Point *b) : endpoint1(a), endpoint2(b) {}

double Segment::length() {
	double x1 = endpoint1->getX();
	double y1 = endpoint1->getY();
	double x2 = endpoint2->getX();
	double y2 = endpoint2->getY();
	return pow(pow(x1-x2,2) + pow(y1-y2,2), 0.5);
}

__device__ void parametrize(Segment &e, Segment &f, double *t1, double *t2, double *det) {
	// parametrize and represent as matrix equation to be solved: 
	// t1 * x0 + (1-t1) * x1 = t2 * x2 + (1-t2) * x3
	// (x0-x1) * t1 + (x3-x2) * t2 = x3 - x1

	// note this originally was passed in to create a matrix
	// but dynamic matrix memory allocation costs time;
	// instead, use in this array form
	double arr[4];
	// first column (matches t1)
	arr[0] = e.endpoint1->getX() - e.endpoint2->getX();
	arr[2] = e.endpoint1->getY() - e.endpoint2->getY();
	// second column (matches t2)
	arr[1] = f.endpoint2->getX() - f.endpoint1->getX();
	arr[3] = f.endpoint2->getY() - f.endpoint1->getY();
	
	double determinant = arr[0] * arr[3] - arr[1] * arr[2];
	// target vector
	double targX = f.endpoint2->getX() - e.endpoint2->getX();
	double targY = f.endpoint2->getY() - e.endpoint2->getY();
	// scaled solution is adjugate of arr multiplied by target
	// adjugate is arr[3], -arr[1], -arr[2], arr[0]
	*t1 = arr[3] * targX - arr[1] * targY;
	*t2 = -arr[2] * targX + arr[0] * targY;
	*det = determinant;
}

__device__ bool Segment::intersection(Segment &other, Point *pt) {
	double t1 = 0;
	double t2 = 0;
	double det = 0;
	parametrize(*this, other, &t1, &t2, &det);
	bool detect = (det != 0) && isBetween(t1, 0, det) && isBetween(t2, 0, det);
	if(detect && pt) {
		double x = (endpoint1->getX() * t1 + endpoint2->getX() * (det - t1)) / det;
		double y = (endpoint1->getY() * t1 + endpoint2->getY() * (det - t1)) / det;
		*pt = Point(x, y);
	}
	return detect;
}

__device__ void Segment::unitNormal(double *nx, double *ny) {
	double deltaX = endpoint2->getX() - endpoint1->getX();
	double deltaY = endpoint2->getY() - endpoint1->getY();
	double unitX = deltaX / length();
	double unitY = deltaY / length();
	*nx = unitY;
	*ny = -unitX;
}

void Segment::scaledNormal(double *nx, double *ny) {
	double deltaX = endpoint2->getX() - endpoint1->getX();
	double deltaY = endpoint2->getY() - endpoint1->getY();
	*nx = deltaY/2;
	*ny = -deltaX/2;
}

	.file	"tmpxft_0027d01d_00000000-6_segment.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB3929:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE3929:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z9isBetweenRKdS0_S0_
	.type	_Z9isBetweenRKdS0_S0_, @function
_Z9isBetweenRKdS0_S0_:
.LFB3915:
	.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
.LFE3915:
	.size	_Z9isBetweenRKdS0_S0_, .-_Z9isBetweenRKdS0_S0_
	.align 2
	.globl	_ZN7SegmentC2Ev
	.type	_ZN7SegmentC2Ev, @function
_ZN7SegmentC2Ev:
.LFB3917:
	.cfi_startproc
	endbr64
	xorl	%eax, %eax
	movq	%rax, -16(%rsp)
	movq	-16(%rsp), %rdi
	jmp	free@PLT
	.cfi_endproc
.LFE3917:
	.size	_ZN7SegmentC2Ev, .-_ZN7SegmentC2Ev
	.globl	_ZN7SegmentC1Ev
	.set	_ZN7SegmentC1Ev,_ZN7SegmentC2Ev
	.align 2
	.globl	_ZN7SegmentC2EP5PointS1_
	.type	_ZN7SegmentC2EP5PointS1_, @function
_ZN7SegmentC2EP5PointS1_:
.LFB3920:
	.cfi_startproc
	endbr64
	movq	%rsi, (%rdi)
	movq	%rdx, 8(%rdi)
	ret
	.cfi_endproc
.LFE3920:
	.size	_ZN7SegmentC2EP5PointS1_, .-_ZN7SegmentC2EP5PointS1_
	.globl	_ZN7SegmentC1EP5PointS1_
	.set	_ZN7SegmentC1EP5PointS1_,_ZN7SegmentC2EP5PointS1_
	.align 2
	.globl	_ZN7Segment6lengthEv
	.type	_ZN7Segment6lengthEv, @function
_ZN7Segment6lengthEv:
.LFB3922:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movq	%rdi, %rbx
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	movq	(%rdi), %rdi
	call	_ZNK5Point4getXEv@PLT
	movq	(%rbx), %rdi
	movsd	%xmm0, 24(%rsp)
	call	_ZNK5Point4getYEv@PLT
	movq	8(%rbx), %rdi
	movsd	%xmm0, 16(%rsp)
	call	_ZNK5Point4getXEv@PLT
	movq	8(%rbx), %rdi
	movsd	%xmm0, 8(%rsp)
	call	_ZNK5Point4getYEv@PLT
	movsd	16(%rsp), %xmm2
	movsd	24(%rsp), %xmm1
	movaps	%xmm0, %xmm3
	subsd	8(%rsp), %xmm1
	subsd	%xmm3, %xmm2
	movaps	%xmm1, %xmm0
	mulsd	%xmm0, %xmm0
	movaps	%xmm2, %xmm1
	mulsd	%xmm2, %xmm1
	addsd	%xmm1, %xmm0
	movsd	.LC0(%rip), %xmm1
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	pow@PLT
	.cfi_endproc
.LFE3922:
	.size	_ZN7Segment6lengthEv, .-_ZN7Segment6lengthEv
	.globl	_Z11parametrizeR7SegmentS0_PdS1_S1_
	.type	_Z11parametrizeR7SegmentS0_PdS1_S1_, @function
_Z11parametrizeR7SegmentS0_PdS1_S1_:
.LFB3923:
	.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
.LFE3923:
	.size	_Z11parametrizeR7SegmentS0_PdS1_S1_, .-_Z11parametrizeR7SegmentS0_PdS1_S1_
	.align 2
	.globl	_ZN7Segment12intersectionERS_P5Point
	.type	_ZN7Segment12intersectionERS_P5Point, @function
_ZN7Segment12intersectionERS_P5Point:
.LFB3924:
	.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
.LFE3924:
	.size	_ZN7Segment12intersectionERS_P5Point, .-_ZN7Segment12intersectionERS_P5Point
	.align 2
	.globl	_ZN7Segment10unitNormalEPdS0_
	.type	_ZN7Segment10unitNormalEPdS0_, @function
_ZN7Segment10unitNormalEPdS0_:
.LFB3925:
	.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
.LFE3925:
	.size	_ZN7Segment10unitNormalEPdS0_, .-_ZN7Segment10unitNormalEPdS0_
	.align 2
	.globl	_ZN7Segment12scaledNormalEPdS0_
	.type	_ZN7Segment12scaledNormalEPdS0_, @function
_ZN7Segment12scaledNormalEPdS0_:
.LFB3926:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	movq	%rsi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	movq	%rdx, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	movq	%rdi, %rbx
	subq	$16, %rsp
	.cfi_def_cfa_offset 48
	movq	8(%rdi), %rdi
	call	_ZNK5Point4getXEv@PLT
	movq	(%rbx), %rdi
	movsd	%xmm0, (%rsp)
	call	_ZNK5Point4getXEv@PLT
	movsd	(%rsp), %xmm1
	movq	8(%rbx), %rdi
	subsd	%xmm0, %xmm1
	movsd	%xmm1, 8(%rsp)
	call	_ZNK5Point4getYEv@PLT
	movq	(%rbx), %rdi
	movsd	%xmm0, (%rsp)
	call	_ZNK5Point4getYEv@PLT
	movsd	(%rsp), %xmm2
	movsd	8(%rsp), %xmm1
	xorps	.LC1(%rip), %xmm1
	movsd	.LC0(%rip), %xmm3
	subsd	%xmm0, %xmm2
	movaps	%xmm1, %xmm0
	mulsd	%xmm3, %xmm0
	mulsd	%xmm3, %xmm2
	movsd	%xmm2, (%r12)
	movsd	%xmm0, 0(%rbp)
	addq	$16, %rsp
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3926:
	.size	_ZN7Segment12scaledNormalEPdS0_, .-_ZN7Segment12scaledNormalEPdS0_
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB3952:
	.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
.LFE3952:
	.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
.LC0:
	.long	0
	.long	1071644672
	.section	.rodata.cst16,"aM",@progbits,16
	.align 16
.LC1:
	.long	0
	.long	-2147483648
	.long	0
	.long	0
	.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 segment_cuh
#define segment_cuh

#ifdef __HIPCC__
#define CUDA_HOSTDEV __host__ __device__
#define CUDA_DEV __device__
#else
#define CUDA_HOSTDEV
#define CUDA_DEV
#endif

#include <math.h>
#include "point.cuh"

/**
 * represent a point (point1, point2) in the plane
 */

class Segment {
	private:
		Point *endpoint1, *endpoint2;
	public:
		CUDA_DEV Segment(); // this just exists to create arrays
		CUDA_HOSTDEV Segment(Point *a, Point *b);
		CUDA_HOSTDEV double length();

		// combine intersection determination and finding point of intersection;
		// return true if this segment intersects other in exactly one point,
		// store intersection point in pt if they intersect (undefined behavior if
		// intersection does not exist or if segments overlap with positive length)
		CUDA_DEV bool intersection(Segment &other, Point *pt = NULL);
		
		// get the 2x1 unit normal to this segment which lies on the right
		// going from endpoint 1 to endpoint 2, setting values in nx, ny
		CUDA_DEV void unitNormal(double *nx, double *ny);
		// return the 2x1 normal to this segment that has length |segment|/2
		void scaledNormal(double *nx, double *ny);

		// helper function that determines parameters at intersection point of e and f,
		// storing as t1, t2; intersection is x0 * (t1 / det) + x1 * (1 - t1/det)
		friend CUDA_DEV void parametrize(Segment &e, Segment &f, double *t1, double *t2, double *det);
		friend class Pixel;
};

#endif 

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

                                        # End of file scope inline assembly
	.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
