#include "plane.cuh"

#include <stdio.h>

__host__ __device__ Plane3::Plane3() {
    a = Vec3();
    b = Vec3();
    c = Vec3();
}

__host__ __device__ Plane3::Plane3(Vec3 a, Vec3 b, Vec3 c) {
    this->a = a;
    this->b = b;
    this->c = c;
}

__host__ __device__ float determinant(Vec3 a, Vec3 b, Vec3 c) {
    return (a.x * b.y * c.z) + (b.x * c.y * a.z) + (c.x * a.y * b.z) - (c.x * b.y * a.z) - (b.x * a.y * c.z) - (a.x * c.y * b.z);
}

__host__ __device__ float Plane3::outside(Vec3 x) {
    Vec3 b_prime = b - a;
    Vec3 c_prime = c - a;
    Vec3 x_prime = x - a;

    return determinant(b_prime, c_prime, x_prime);
}

__host__ __device__ Vec3 Plane3::intersection_point(Vec3 line_origin, Vec3 line_direction) {
    Vec3 plane_normal = ((c - a).cross(b - a)).normalize();
    float denom = plane_normal.dot(line_direction.normalize());
    float t = 99999999;


    if (denom > 0.0001) {
        Vec3 line_origin_to_point = a - line_origin;
        t = line_origin_to_point.dot(plane_normal) / denom;
    }

    return line_origin + (line_direction.normalize() * t);


    // float t  = (plane_normal.dot(a) - plane_normal.dot(line_origin)) / plane_normal.dot(line_direction.normalize());
    // if (t < 0) // Falling of screen
    //     t = 99999999.0;
    // return line_origin + ((line_direction).normalize() * t);
}
	.file	"tmpxft_002931f0_00000000-6_plane.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2049:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2049:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.text._ZmiRK4Vec3S1_,"axG",@progbits,_ZmiRK4Vec3S1_,comdat
	.weak	_ZmiRK4Vec3S1_
	.type	_ZmiRK4Vec3S1_, @function
_ZmiRK4Vec3S1_:
.LFB2011:
	.cfi_startproc
	endbr64
	subq	$56, %rsp
	.cfi_def_cfa_offset 64
	movss	8(%rdi), %xmm2
	movss	4(%rdi), %xmm1
	movq	%fs:40, %rax
	movq	%rax, 40(%rsp)
	xorl	%eax, %eax
	movss	(%rdi), %xmm0
	subss	(%rsi), %xmm0
	leaq	16(%rsp), %rdi
	subss	4(%rsi), %xmm1
	subss	8(%rsi), %xmm2
	call	_ZN4Vec3C1Efff@PLT
	movl	24(%rsp), %eax
	movq	16(%rsp), %xmm0
	movl	%eax, 36(%rsp)
	movss	36(%rsp), %xmm1
	movq	40(%rsp), %rax
	subq	%fs:40, %rax
	je	.L3
	call	__stack_chk_fail@PLT
.L3:
	addq	$56, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2011:
	.size	_ZmiRK4Vec3S1_, .-_ZmiRK4Vec3S1_
	.text
	.align 2
	.globl	_ZN6Plane3C2Ev
	.type	_ZN6Plane3C2Ev, @function
_ZN6Plane3C2Ev:
.LFB2039:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movq	%rdi, %rbx
	subq	$40, %rsp
	.cfi_def_cfa_offset 64
	movq	%fs:40, %rax
	movq	%rax, 24(%rsp)
	xorl	%eax, %eax
	leaq	12(%rsp), %rbp
	call	_ZN4Vec3C1Ev@PLT
	leaq	12(%rbx), %rdi
	call	_ZN4Vec3C1Ev@PLT
	leaq	24(%rbx), %rdi
	call	_ZN4Vec3C1Ev@PLT
	movq	%rbp, %rdi
	call	_ZN4Vec3C1Ev@PLT
	movq	12(%rsp), %rax
	movq	%rbp, %rdi
	movq	%rax, (%rbx)
	movl	20(%rsp), %eax
	movl	%eax, 8(%rbx)
	call	_ZN4Vec3C1Ev@PLT
	movq	12(%rsp), %rax
	movq	%rbp, %rdi
	movq	%rax, 12(%rbx)
	movl	20(%rsp), %eax
	movl	%eax, 20(%rbx)
	call	_ZN4Vec3C1Ev@PLT
	movq	12(%rsp), %rax
	movq	%rax, 24(%rbx)
	movl	20(%rsp), %eax
	movl	%eax, 32(%rbx)
	movq	24(%rsp), %rax
	subq	%fs:40, %rax
	je	.L7
	call	__stack_chk_fail@PLT
.L7:
	addq	$40, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2039:
	.size	_ZN6Plane3C2Ev, .-_ZN6Plane3C2Ev
	.globl	_ZN6Plane3C1Ev
	.set	_ZN6Plane3C1Ev,_ZN6Plane3C2Ev
	.align 2
	.globl	_ZN6Plane3C2E4Vec3S0_S0_
	.type	_ZN6Plane3C2E4Vec3S0_S0_, @function
_ZN6Plane3C2E4Vec3S0_S0_:
.LFB2042:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movq	%rdi, %rbx
	subq	$48, %rsp
	.cfi_def_cfa_offset 64
	movq	%xmm0, 32(%rsp)
	movss	%xmm1, 40(%rsp)
	movq	%xmm2, 16(%rsp)
	movss	%xmm3, 24(%rsp)
	movq	%xmm4, (%rsp)
	movss	%xmm5, 8(%rsp)
	call	_ZN4Vec3C1Ev@PLT
	leaq	12(%rbx), %rdi
	call	_ZN4Vec3C1Ev@PLT
	leaq	24(%rbx), %rdi
	call	_ZN4Vec3C1Ev@PLT
	movq	32(%rsp), %rax
	movq	%rax, (%rbx)
	movl	40(%rsp), %eax
	movl	%eax, 8(%rbx)
	movq	16(%rsp), %rax
	movq	%rax, 12(%rbx)
	movl	24(%rsp), %eax
	movl	%eax, 20(%rbx)
	movq	(%rsp), %rax
	movq	%rax, 24(%rbx)
	movl	8(%rsp), %eax
	movl	%eax, 32(%rbx)
	addq	$48, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2042:
	.size	_ZN6Plane3C2E4Vec3S0_S0_, .-_ZN6Plane3C2E4Vec3S0_S0_
	.globl	_ZN6Plane3C1E4Vec3S0_S0_
	.set	_ZN6Plane3C1E4Vec3S0_S0_,_ZN6Plane3C2E4Vec3S0_S0_
	.globl	_Z11determinant4Vec3S_S_
	.type	_Z11determinant4Vec3S_S_, @function
_Z11determinant4Vec3S_S_:
.LFB2044:
	.cfi_startproc
	endbr64
	movq	%xmm0, -16(%rsp)
	movq	%xmm2, -32(%rsp)
	movss	-16(%rsp), %xmm2
	movss	-32(%rsp), %xmm9
	movq	%xmm4, -48(%rsp)
	movss	-28(%rsp), %xmm6
	movss	-44(%rsp), %xmm7
	movaps	%xmm2, %xmm0
	movaps	%xmm9, %xmm8
	movss	-12(%rsp), %xmm4
	movss	-48(%rsp), %xmm10
	mulss	%xmm6, %xmm0
	mulss	%xmm7, %xmm8
	mulss	%xmm10, %xmm6
	mulss	%xmm7, %xmm2
	mulss	%xmm5, %xmm0
	mulss	%xmm1, %xmm8
	mulss	%xmm1, %xmm6
	mulss	%xmm3, %xmm2
	addss	%xmm8, %xmm0
	movaps	%xmm4, %xmm8
	mulss	%xmm10, %xmm8
	mulss	%xmm9, %xmm4
	mulss	%xmm3, %xmm8
	mulss	%xmm5, %xmm4
	addss	%xmm8, %xmm0
	subss	%xmm6, %xmm0
	subss	%xmm4, %xmm0
	subss	%xmm2, %xmm0
	ret
	.cfi_endproc
.LFE2044:
	.size	_Z11determinant4Vec3S_S_, .-_Z11determinant4Vec3S_S_
	.align 2
	.globl	_ZN6Plane37outsideE4Vec3
	.type	_ZN6Plane37outsideE4Vec3, @function
_ZN6Plane37outsideE4Vec3:
.LFB2045:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movq	%rdi, %rbx
	leaq	12(%rdi), %rdi
	movq	%rbx, %rsi
	subq	$64, %rsp
	.cfi_def_cfa_offset 80
	movq	%xmm0, (%rsp)
	movss	%xmm1, 8(%rsp)
	call	_ZmiRK4Vec3S1_
	leaq	24(%rbx), %rdi
	movq	%rbx, %rsi
	movq	%xmm0, 28(%rsp)
	movss	%xmm1, 36(%rsp)
	call	_ZmiRK4Vec3S1_
	movq	%rsp, %rdi
	movq	%rbx, %rsi
	movq	%xmm0, 40(%rsp)
	movss	%xmm1, 48(%rsp)
	call	_ZmiRK4Vec3S1_
	movq	40(%rsp), %xmm2
	movss	48(%rsp), %xmm3
	movaps	%xmm0, %xmm4
	movaps	%xmm1, %xmm5
	movq	28(%rsp), %xmm0
	movss	36(%rsp), %xmm1
	call	_Z11determinant4Vec3S_S_
	addq	$64, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2045:
	.size	_ZN6Plane37outsideE4Vec3, .-_ZN6Plane37outsideE4Vec3
	.align 2
	.globl	_ZN6Plane318intersection_pointE4Vec3S0_
	.type	_ZN6Plane318intersection_pointE4Vec3S0_, @function
_ZN6Plane318intersection_pointE4Vec3S0_:
.LFB2046:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	movq	%rdi, %rbx
	addq	$24, %rdi
	movq	%rbx, %rsi
	subq	$144, %rsp
	.cfi_def_cfa_offset 176
	movss	%xmm3, 40(%rsp)
	leaq	124(%rsp), %rbp
	leaq	32(%rsp), %r12
	movq	%xmm2, 32(%rsp)
	movq	%xmm0, 48(%rsp)
	movss	%xmm1, 56(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	call	_ZmiRK4Vec3S1_
	movq	%rbx, %rsi
	leaq	12(%rbx), %rdi
	movq	%xmm0, 112(%rsp)
	movss	%xmm1, 120(%rsp)
	call	_ZmiRK4Vec3S1_
	leaq	112(%rsp), %rdi
	call	_ZN4Vec35crossES_@PLT
	movq	%rbp, %rdi
	movq	%xmm0, 124(%rsp)
	movss	%xmm1, 132(%rsp)
	call	_ZN4Vec39normalizeEv@PLT
	movq	%r12, %rdi
	movq	%xmm0, 76(%rsp)
	movss	%xmm1, 84(%rsp)
	call	_ZN4Vec39normalizeEv@PLT
	leaq	76(%rsp), %rdi
	movq	%xmm0, 112(%rsp)
	movss	%xmm1, 120(%rsp)
	call	_ZN4Vec33dotES_@PLT
	movss	.LC0(%rip), %xmm3
	movaps	%xmm0, %xmm2
	cvtss2sd	%xmm0, %xmm0
	comisd	.LC1(%rip), %xmm0
	jbe	.L15
	leaq	48(%rsp), %rsi
	movq	%rbx, %rdi
	movss	%xmm2, 12(%rsp)
	call	_ZmiRK4Vec3S1_
	movq	%rbp, %rdi
	movq	%xmm0, 124(%rsp)
	movq	76(%rsp), %xmm0
	movss	%xmm1, 132(%rsp)
	movss	84(%rsp), %xmm1
	call	_ZN4Vec33dotES_@PLT
	movss	12(%rsp), %xmm2
	movaps	%xmm0, %xmm3
	divss	%xmm2, %xmm3
.L15:
	movq	%r12, %rdi
	movss	%xmm3, 12(%rsp)
	call	_ZN4Vec39normalizeEv@PLT
	movss	12(%rsp), %xmm3
	movq	%rbp, %rdi
	movq	%xmm0, 88(%rsp)
	movss	88(%rsp), %xmm0
	movaps	%xmm1, %xmm2
	mulss	%xmm3, %xmm2
	mulss	%xmm3, %xmm0
	mulss	92(%rsp), %xmm3
	movaps	%xmm3, %xmm1
	call	_ZN4Vec3C1Efff@PLT
	movss	124(%rsp), %xmm0
	movq	%rbp, %rdi
	addss	48(%rsp), %xmm0
	movss	132(%rsp), %xmm2
	addss	56(%rsp), %xmm2
	movss	128(%rsp), %xmm1
	addss	52(%rsp), %xmm1
	call	_ZN4Vec3C1Efff@PLT
	movq	124(%rsp), %xmm0
	movss	132(%rsp), %xmm1
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L17
	call	__stack_chk_fail@PLT
.L17:
	addq	$144, %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
.LFE2046:
	.size	_ZN6Plane318intersection_pointE4Vec3S0_, .-_ZN6Plane318intersection_pointE4Vec3S0_
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2072:
	.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
.LFE2072:
	.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.cst4,"aM",@progbits,4
	.align 4
.LC0:
	.long	1287568416
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC1:
	.long	-350469331
	.long	1058682594
	.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 HEADER_PLANE
#define HEADER_PLANE

#include "vector.cuh"

class Plane3 {
public:
    __host__ __device__ Plane3();
    __host__ __device__ Plane3(Vec3 a, Vec3 b, Vec3 c);

    Vec3 a;
    Vec3 b;
    Vec3 c;

    __host__ __device__ float outside(Vec3 point);
    __host__ __device__ Vec3 intersection_point(Vec3 line_origin, Vec3 line_direction);
};

#endif
	.text
	.file	"plane.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
