#include "../include/camera.cuh"
#include "../include/math_utils.cuh"

__host__ __device__
Camera::Camera(int width, int height, float fov, const vec3 &origin, const vec3 &dir)
    : width{width}, height{height}, fov{fov * M_PI_F / 180.0f},
    origin{origin}, dir{dir}
    {
        inv_width = 1.0f / width;
        inv_height = 1.0f / height;
        aspect_ratio = width / float(height);
        angle = tan(0.5f * fov);

        calc_axes();
    }

__host__ __device__
void Camera::move(const vec3 &new_origin, const vec3 &new_dir) {
    origin = new_origin;
    dir = new_dir;
    dir.normalize();
    calc_axes();
}

__host__ __device__
void Camera::move_from_to(const vec3 &from, const vec3 &to) {
    move(from, to - from);
}

__host__ __device__
vec3 Camera::ray_dir_at_pixel(float x, float y) const {
    float xx = (2 * x * inv_width - 1) * angle * aspect_ratio; 
    float yy = (1 - 2 * y * inv_height) * angle;

    vec3 img_pt = right * xx + up * yy + dir;
    img_pt.normalize();
    return img_pt;
}

__host__ __device__
void Camera::calc_axes() {
    // approximately "up"
    vec3 tmp_up = {0, 1, 0};
    tmp_up.normalize();
    right = cross(dir, tmp_up);
    right.normalize();
    up = cross(right, dir);
}
	.file	"tmpxft_002e1be0_00000000-6_camera.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB3681:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE3681:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.align 2
	.type	_ZN4vec39normalizeEv.isra.0, @function
_ZN4vec39normalizeEv.isra.0:
.LFB4327:
	.cfi_startproc
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movq	%rdi, %rbx
	subq	$16, %rsp
	.cfi_def_cfa_offset 32
	movss	(%rdi), %xmm3
	movss	4(%rdi), %xmm2
	movss	8(%rdi), %xmm1
	movaps	%xmm3, %xmm0
	movaps	%xmm2, %xmm4
	movss	%xmm3, 12(%rsp)
	mulss	%xmm2, %xmm4
	movss	%xmm2, 8(%rsp)
	mulss	%xmm3, %xmm0
	movss	%xmm1, 4(%rsp)
	addss	%xmm4, %xmm0
	movaps	%xmm1, %xmm4
	mulss	%xmm1, %xmm4
	addss	%xmm4, %xmm0
	call	sqrtf@PLT
	movss	12(%rsp), %xmm3
	movss	8(%rsp), %xmm2
	movaps	%xmm0, %xmm4
	movss	.LC0(%rip), %xmm0
	movss	4(%rsp), %xmm1
	divss	%xmm4, %xmm0
	mulss	%xmm0, %xmm3
	mulss	%xmm0, %xmm2
	mulss	%xmm0, %xmm1
	movss	%xmm3, (%rbx)
	movss	%xmm2, 4(%rbx)
	movss	%xmm1, 8(%rbx)
	addq	$16, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4327:
	.size	_ZN4vec39normalizeEv.isra.0, .-_ZN4vec39normalizeEv.isra.0
	.section	.text._Z5crossRK4vec3S1_,"axG",@progbits,_Z5crossRK4vec3S1_,comdat
	.weak	_Z5crossRK4vec3S1_
	.type	_Z5crossRK4vec3S1_, @function
_Z5crossRK4vec3S1_:
.LFB3667:
	.cfi_startproc
	endbr64
	subq	$56, %rsp
	.cfi_def_cfa_offset 64
	movss	4(%rsi), %xmm6
	movss	4(%rdi), %xmm2
	movss	(%rdi), %xmm1
	movss	(%rsi), %xmm5
	movss	8(%rsi), %xmm0
	movss	8(%rdi), %xmm3
	movaps	%xmm2, %xmm4
	movaps	%xmm6, %xmm7
	movq	%fs:40, %rax
	movq	%rax, 40(%rsp)
	xorl	%eax, %eax
	mulss	%xmm5, %xmm2
	mulss	%xmm0, %xmm4
	mulss	%xmm3, %xmm7
	mulss	%xmm1, %xmm0
	mulss	%xmm6, %xmm1
	mulss	%xmm5, %xmm3
	subss	%xmm7, %xmm4
	subss	%xmm2, %xmm1
	subss	%xmm3, %xmm0
	xorps	.LC1(%rip), %xmm0
	movss	%xmm4, 28(%rsp)
	movss	%xmm0, 32(%rsp)
	movq	28(%rsp), %xmm0
	movss	%xmm1, 36(%rsp)
	movd	36(%rsp), %xmm1
	movq	40(%rsp), %rax
	subq	%fs:40, %rax
	je	.L5
	call	__stack_chk_fail@PLT
.L5:
	addq	$56, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3667:
	.size	_Z5crossRK4vec3S1_, .-_Z5crossRK4vec3S1_
	.text
	.align 2
	.globl	_ZNK6Camera16ray_dir_at_pixelEff
	.type	_ZNK6Camera16ray_dir_at_pixelEff, @function
_ZNK6Camera16ray_dir_at_pixelEff:
.LFB3677:
	.cfi_startproc
	endbr64
	subq	$56, %rsp
	.cfi_def_cfa_offset 64
	addss	%xmm0, %xmm0
	addss	%xmm1, %xmm1
	movss	.LC0(%rip), %xmm2
	mulss	12(%rdi), %xmm1
	movss	24(%rdi), %xmm3
	movq	%fs:40, %rax
	movq	%rax, 40(%rsp)
	xorl	%eax, %eax
	mulss	8(%rdi), %xmm0
	movss	56(%rdi), %xmm4
	subss	%xmm2, %xmm0
	subss	%xmm1, %xmm2
	mulss	%xmm3, %xmm2
	mulss	%xmm3, %xmm0
	movss	60(%rdi), %xmm3
	mulss	16(%rdi), %xmm0
	movaps	%xmm2, %xmm1
	movss	72(%rdi), %xmm2
	mulss	%xmm1, %xmm3
	mulss	%xmm0, %xmm2
	mulss	%xmm1, %xmm4
	mulss	52(%rdi), %xmm1
	addss	%xmm3, %xmm2
	movss	68(%rdi), %xmm3
	addss	48(%rdi), %xmm2
	mulss	%xmm0, %xmm3
	mulss	64(%rdi), %xmm0
	movss	%xmm2, 24(%rsp)
	addss	%xmm4, %xmm3
	addss	44(%rdi), %xmm3
	addss	%xmm1, %xmm0
	addss	40(%rdi), %xmm0
	leaq	16(%rsp), %rdi
	movss	%xmm3, 20(%rsp)
	movss	%xmm0, 16(%rsp)
	call	_ZN4vec39normalizeEv.isra.0
	movq	16(%rsp), %xmm0
	movd	24(%rsp), %xmm1
	movq	40(%rsp), %rax
	subq	%fs:40, %rax
	je	.L9
	call	__stack_chk_fail@PLT
.L9:
	addq	$56, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3677:
	.size	_ZNK6Camera16ray_dir_at_pixelEff, .-_ZNK6Camera16ray_dir_at_pixelEff
	.align 2
	.globl	_ZN6Camera9calc_axesEv
	.type	_ZN6Camera9calc_axesEv, @function
_ZN6Camera9calc_axesEv:
.LFB3678:
	.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
	leaq	40(%rbx), %r12
	subq	$32, %rsp
	.cfi_def_cfa_offset 64
	movq	%fs:40, %rax
	movq	%rax, 24(%rsp)
	movabsq	$4575657221408423936, %rax
	leaq	12(%rsp), %rbp
	movl	$0x00000000, 20(%rsp)
	movq	%rbp, %rdi
	movq	%rax, 12(%rsp)
	call	_ZN4vec39normalizeEv.isra.0
	movq	%rbp, %rsi
	movq	%r12, %rdi
	leaq	64(%rbx), %rbp
	call	_Z5crossRK4vec3S1_
	movq	%rbp, %rdi
	movq	%xmm0, 64(%rbx)
	movd	%xmm1, 72(%rbx)
	call	_ZN4vec39normalizeEv.isra.0
	movq	%r12, %rsi
	movq	%rbp, %rdi
	call	_Z5crossRK4vec3S1_
	movq	%xmm0, 52(%rbx)
	movd	%xmm1, 60(%rbx)
	movq	24(%rsp), %rax
	subq	%fs:40, %rax
	je	.L12
	call	__stack_chk_fail@PLT
.L12:
	addq	$32, %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
.LFE3678:
	.size	_ZN6Camera9calc_axesEv, .-_ZN6Camera9calc_axesEv
	.align 2
	.globl	_ZN6CameraC2EiifRK4vec3S2_
	.type	_ZN6CameraC2EiifRK4vec3S2_, @function
_ZN6CameraC2EiifRK4vec3S2_:
.LFB3673:
	.cfi_startproc
	endbr64
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movss	.LC3(%rip), %xmm1
	movl	%edx, %eax
	movl	%esi, (%rdi)
	movss	.LC0(%rip), %xmm2
	movl	%edx, 4(%rdi)
	mulss	%xmm0, %xmm1
	movq	(%rcx), %rdx
	movq	%rdi, 8(%rsp)
	divss	.LC4(%rip), %xmm1
	movss	%xmm1, 20(%rdi)
	cvtsi2ssl	%esi, %xmm1
	movaps	%xmm2, %xmm3
	movq	%rdx, 28(%rdi)
	movl	8(%rcx), %edx
	mulss	.LC5(%rip), %xmm0
	movl	%edx, 36(%rdi)
	movq	(%r8), %rdx
	divss	%xmm1, %xmm3
	movq	%rdx, 40(%rdi)
	movl	8(%r8), %edx
	movl	%edx, 48(%rdi)
	movss	%xmm3, 8(%rdi)
	cvtsi2ssl	%eax, %xmm3
	divss	%xmm3, %xmm2
	divss	%xmm3, %xmm1
	movss	%xmm2, 12(%rdi)
	movss	%xmm1, 16(%rdi)
	call	tanf@PLT
	movq	8(%rsp), %rdi
	movss	%xmm0, 24(%rdi)
	addq	$24, %rsp
	.cfi_def_cfa_offset 8
	jmp	_ZN6Camera9calc_axesEv
	.cfi_endproc
.LFE3673:
	.size	_ZN6CameraC2EiifRK4vec3S2_, .-_ZN6CameraC2EiifRK4vec3S2_
	.globl	_ZN6CameraC1EiifRK4vec3S2_
	.set	_ZN6CameraC1EiifRK4vec3S2_,_ZN6CameraC2EiifRK4vec3S2_
	.align 2
	.globl	_ZN6Camera4moveERK4vec3S2_
	.type	_ZN6Camera4moveERK4vec3S2_, @function
_ZN6Camera4moveERK4vec3S2_:
.LFB3675:
	.cfi_startproc
	endbr64
	movq	(%rsi), %rax
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movq	%rdi, %rbx
	leaq	40(%rdi), %rdi
	movq	%rax, -12(%rdi)
	movl	8(%rsi), %eax
	movl	%eax, -4(%rdi)
	movq	(%rdx), %rax
	movq	%rax, (%rdi)
	movl	8(%rdx), %eax
	movl	%eax, 8(%rdi)
	call	_ZN4vec39normalizeEv.isra.0
	movq	%rbx, %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	_ZN6Camera9calc_axesEv
	.cfi_endproc
.LFE3675:
	.size	_ZN6Camera4moveERK4vec3S2_, .-_ZN6Camera4moveERK4vec3S2_
	.align 2
	.globl	_ZN6Camera12move_from_toERK4vec3S2_
	.type	_ZN6Camera12move_from_toERK4vec3S2_, @function
_ZN6Camera12move_from_toERK4vec3S2_:
.LFB3676:
	.cfi_startproc
	endbr64
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movss	8(%rdx), %xmm0
	movss	4(%rdx), %xmm1
	movq	%fs:40, %rax
	movq	%rax, 24(%rsp)
	xorl	%eax, %eax
	movss	(%rdx), %xmm2
	subss	(%rsi), %xmm2
	leaq	12(%rsp), %rdx
	subss	8(%rsi), %xmm0
	subss	4(%rsi), %xmm1
	movss	%xmm2, 12(%rsp)
	movss	%xmm1, 16(%rsp)
	movss	%xmm0, 20(%rsp)
	call	_ZN6Camera4moveERK4vec3S2_
	movq	24(%rsp), %rax
	subq	%fs:40, %rax
	je	.L19
	call	__stack_chk_fail@PLT
.L19:
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3676:
	.size	_ZN6Camera12move_from_toERK4vec3S2_, .-_ZN6Camera12move_from_toERK4vec3S2_
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB3704:
	.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
.LFE3704:
	.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	1065353216
	.section	.rodata.cst16,"aM",@progbits,16
	.align 16
.LC1:
	.long	-2147483648
	.long	0
	.long	0
	.long	0
	.section	.rodata.cst4
	.align 4
.LC3:
	.long	1078530011
	.align 4
.LC4:
	.long	1127481344
	.align 4
.LC5:
	.long	1056964608
	.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>
#include "../include/camera.cuh"
#include "../include/math_utils.cuh"

__host__ __device__
Camera::Camera(int width, int height, float fov, const vec3 &origin, const vec3 &dir)
    : width{width}, height{height}, fov{fov * M_PI_F / 180.0f},
    origin{origin}, dir{dir}
    {
        inv_width = 1.0f / width;
        inv_height = 1.0f / height;
        aspect_ratio = width / float(height);
        angle = tan(0.5f * fov);

        calc_axes();
    }

__host__ __device__
void Camera::move(const vec3 &new_origin, const vec3 &new_dir) {
    origin = new_origin;
    dir = new_dir;
    dir.normalize();
    calc_axes();
}

__host__ __device__
void Camera::move_from_to(const vec3 &from, const vec3 &to) {
    move(from, to - from);
}

__host__ __device__
vec3 Camera::ray_dir_at_pixel(float x, float y) const {
    float xx = (2 * x * inv_width - 1) * angle * aspect_ratio; 
    float yy = (1 - 2 * y * inv_height) * angle;

    vec3 img_pt = right * xx + up * yy + dir;
    img_pt.normalize();
    return img_pt;
}

__host__ __device__
void Camera::calc_axes() {
    // approximately "up"
    vec3 tmp_up = {0, 1, 0};
    tmp_up.normalize();
    right = cross(dir, tmp_up);
    right.normalize();
    up = cross(right, dir);
}
	.text
	.file	"camera.hip"
                                        # Start of file scope inline assembly
	.globl	_ZSt21ios_base_library_initv

                                        # End of file scope inline assembly
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function _ZN6CameraC2EiifRK4vec3S2_
.LCPI0_0:
	.long	0x40490fdb                      # float 3.14159274
.LCPI0_1:
	.long	0x43340000                      # float 180
.LCPI0_2:
	.long	0x3f800000                      # float 1
.LCPI0_3:
	.long	0x3f000000                      # float 0.5
	.text
	.globl	_ZN6CameraC2EiifRK4vec3S2_
	.p2align	1, 0x90
	.type	_ZN6CameraC2EiifRK4vec3S2_,@function
_ZN6CameraC2EiifRK4vec3S2_:             # @_ZN6CameraC2EiifRK4vec3S2_
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset %rbx, -16
	movq	%rdi, %rbx
	movl	%esi, (%rdi)
	movl	%edx, 4(%rdi)
	movss	.LCPI0_0(%rip), %xmm1           # xmm1 = mem[0],zero,zero,zero
	mulss	%xmm0, %xmm1
	divss	.LCPI0_1(%rip), %xmm1
	movss	%xmm1, 20(%rdi)
	movq	(%rcx), %rax
	movq	%rax, 28(%rdi)
	movl	8(%rcx), %eax
	movl	%eax, 36(%rdi)
	movq	(%r8), %rax
	movq	%rax, 40(%rdi)
	movl	8(%r8), %eax
	movl	%eax, 48(%rdi)
	xorps	%xmm1, %xmm1
	cvtsi2ss	%esi, %xmm1
	movss	.LCPI0_2(%rip), %xmm2           # xmm2 = mem[0],zero,zero,zero
	movaps	%xmm2, %xmm3
	cvtsi2ss	%edx, %xmm4
	divss	%xmm1, %xmm3
	movss	%xmm3, 8(%rdi)
	divss	%xmm4, %xmm2
	movss	%xmm2, 12(%rdi)
	divss	%xmm4, %xmm1
	movss	%xmm1, 16(%rdi)
	mulss	.LCPI0_3(%rip), %xmm0
	callq	tanf
	movss	%xmm0, 24(%rbx)
	movq	%rbx, %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	_ZN6Camera9calc_axesEv          # TAILCALL
.Lfunc_end0:
	.size	_ZN6CameraC2EiifRK4vec3S2_, .Lfunc_end0-_ZN6CameraC2EiifRK4vec3S2_
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst16,"aM",@progbits,16
	.p2align	4, 0x0                          # -- Begin function _ZN6Camera9calc_axesEv
.LCPI1_0:
	.long	0x80000000                      # float -0
	.long	0x80000000                      # float -0
	.long	0x80000000                      # float -0
	.long	0x80000000                      # float -0
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0
.LCPI1_1:
	.long	0x3f800000                      # float 1
	.text
	.globl	_ZN6Camera9calc_axesEv
	.p2align	1, 0x90
	.type	_ZN6Camera9calc_axesEv,@function
_ZN6Camera9calc_axesEv:                 # @_ZN6Camera9calc_axesEv
	.cfi_startproc
# %bb.0:
	movss	40(%rdi), %xmm1                 # xmm1 = mem[0],zero,zero,zero
	movss	44(%rdi), %xmm3                 # xmm3 = mem[0],zero,zero,zero
	xorps	%xmm4, %xmm4
	movaps	%xmm3, %xmm0
	mulss	%xmm4, %xmm0
	movss	48(%rdi), %xmm2                 # xmm2 = mem[0],zero,zero,zero
	movaps	%xmm1, %xmm6
	mulss	%xmm4, %xmm6
	mulss	%xmm2, %xmm4
	subss	%xmm4, %xmm6
	movaps	.LCPI1_0(%rip), %xmm5           # xmm5 = [-0.0E+0,-0.0E+0,-0.0E+0,-0.0E+0]
	movaps	%xmm6, %xmm4
	xorps	%xmm5, %xmm4
	movaps	%xmm1, %xmm7
	subss	%xmm0, %xmm7
	subss	%xmm2, %xmm0
	movaps	%xmm0, %xmm8
	mulss	%xmm0, %xmm8
	mulss	%xmm6, %xmm6
	addss	%xmm8, %xmm6
	movaps	%xmm7, %xmm8
	mulss	%xmm7, %xmm8
	addss	%xmm6, %xmm8
	sqrtss	%xmm8, %xmm8
	movss	.LCPI1_1(%rip), %xmm6           # xmm6 = mem[0],zero,zero,zero
	divss	%xmm8, %xmm6
	mulss	%xmm6, %xmm0
	movss	%xmm0, 64(%rdi)
	mulss	%xmm6, %xmm4
	movss	%xmm4, 68(%rdi)
	mulss	%xmm7, %xmm6
	movss	%xmm6, 72(%rdi)
	movaps	%xmm2, %xmm7
	mulss	%xmm4, %xmm7
	mulss	%xmm0, %xmm2
	mulss	%xmm3, %xmm0
	mulss	%xmm6, %xmm3
	subss	%xmm3, %xmm7
	mulss	%xmm1, %xmm6
	subss	%xmm6, %xmm2
	xorps	%xmm5, %xmm2
	unpcklps	%xmm2, %xmm7                    # xmm7 = xmm7[0],xmm2[0],xmm7[1],xmm2[1]
	mulss	%xmm1, %xmm4
	subss	%xmm4, %xmm0
	movlps	%xmm7, 52(%rdi)
	movss	%xmm0, 60(%rdi)
	retq
.Lfunc_end1:
	.size	_ZN6Camera9calc_axesEv, .Lfunc_end1-_ZN6Camera9calc_axesEv
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function _ZN6Camera4moveERK4vec3S2_
.LCPI2_0:
	.long	0x3f800000                      # float 1
	.text
	.globl	_ZN6Camera4moveERK4vec3S2_
	.p2align	1, 0x90
	.type	_ZN6Camera4moveERK4vec3S2_,@function
_ZN6Camera4moveERK4vec3S2_:             # @_ZN6Camera4moveERK4vec3S2_
	.cfi_startproc
# %bb.0:
	movl	8(%rsi), %eax
	movl	%eax, 36(%rdi)
	movq	(%rsi), %rax
	movq	%rax, 28(%rdi)
	movq	(%rdx), %rax
	movq	%rax, 40(%rdi)
	movl	8(%rdx), %eax
	movl	%eax, 48(%rdi)
	movss	40(%rdi), %xmm0                 # xmm0 = mem[0],zero,zero,zero
	movss	44(%rdi), %xmm1                 # xmm1 = mem[0],zero,zero,zero
	movaps	%xmm0, %xmm2
	mulss	%xmm0, %xmm2
	movaps	%xmm1, %xmm3
	mulss	%xmm1, %xmm3
	addss	%xmm2, %xmm3
	movss	48(%rdi), %xmm2                 # xmm2 = mem[0],zero,zero,zero
	movaps	%xmm2, %xmm4
	mulss	%xmm2, %xmm4
	addss	%xmm3, %xmm4
	xorps	%xmm3, %xmm3
	sqrtss	%xmm4, %xmm3
	movss	.LCPI2_0(%rip), %xmm4           # xmm4 = mem[0],zero,zero,zero
	divss	%xmm3, %xmm4
	mulss	%xmm4, %xmm0
	movss	%xmm0, 40(%rdi)
	mulss	%xmm4, %xmm1
	movss	%xmm1, 44(%rdi)
	mulss	%xmm2, %xmm4
	movss	%xmm4, 48(%rdi)
	jmp	_ZN6Camera9calc_axesEv          # TAILCALL
.Lfunc_end2:
	.size	_ZN6Camera4moveERK4vec3S2_, .Lfunc_end2-_ZN6Camera4moveERK4vec3S2_
	.cfi_endproc
                                        # -- End function
	.globl	_ZN6Camera12move_from_toERK4vec3S2_ # -- Begin function _ZN6Camera12move_from_toERK4vec3S2_
	.p2align	1, 0x90
	.type	_ZN6Camera12move_from_toERK4vec3S2_,@function
_ZN6Camera12move_from_toERK4vec3S2_:    # @_ZN6Camera12move_from_toERK4vec3S2_
	.cfi_startproc
# %bb.0:
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movss	(%rdx), %xmm0                   # xmm0 = mem[0],zero,zero,zero
	movss	4(%rdx), %xmm1                  # xmm1 = mem[0],zero,zero,zero
	subss	(%rsi), %xmm0
	subss	4(%rsi), %xmm1
	movss	8(%rdx), %xmm2                  # xmm2 = mem[0],zero,zero,zero
	subss	8(%rsi), %xmm2
	unpcklps	%xmm1, %xmm0                    # xmm0 = xmm0[0],xmm1[0],xmm0[1],xmm1[1]
	leaq	8(%rsp), %rdx
	movlps	%xmm0, (%rdx)
	movss	%xmm2, 8(%rdx)
	callq	_ZN6Camera4moveERK4vec3S2_
	addq	$24, %rsp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end3:
	.size	_ZN6Camera12move_from_toERK4vec3S2_, .Lfunc_end3-_ZN6Camera12move_from_toERK4vec3S2_
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function _ZNK6Camera16ray_dir_at_pixelEff
.LCPI4_0:
	.long	0xbf800000                      # float -1
.LCPI4_1:
	.long	0x3f800000                      # float 1
	.text
	.globl	_ZNK6Camera16ray_dir_at_pixelEff
	.p2align	1, 0x90
	.type	_ZNK6Camera16ray_dir_at_pixelEff,@function
_ZNK6Camera16ray_dir_at_pixelEff:       # @_ZNK6Camera16ray_dir_at_pixelEff
	.cfi_startproc
# %bb.0:
	movaps	%xmm1, %xmm2
	addss	%xmm0, %xmm0
	mulss	8(%rdi), %xmm0
	addss	.LCPI4_0(%rip), %xmm0
	movss	24(%rdi), %xmm4                 # xmm4 = mem[0],zero,zero,zero
	mulss	%xmm4, %xmm0
	mulss	16(%rdi), %xmm0
	addss	%xmm1, %xmm2
	mulss	12(%rdi), %xmm2
	movss	.LCPI4_1(%rip), %xmm1           # xmm1 = mem[0],zero,zero,zero
	movaps	%xmm1, %xmm3
	subss	%xmm2, %xmm3
	movss	64(%rdi), %xmm5                 # xmm5 = mem[0],zero,zero,zero
	mulss	%xmm0, %xmm5
	movss	68(%rdi), %xmm6                 # xmm6 = mem[0],zero,zero,zero
	mulss	%xmm0, %xmm6
	mulss	72(%rdi), %xmm0
	mulss	%xmm4, %xmm3
	movss	52(%rdi), %xmm2                 # xmm2 = mem[0],zero,zero,zero
	mulss	%xmm3, %xmm2
	addss	%xmm5, %xmm2
	movss	56(%rdi), %xmm4                 # xmm4 = mem[0],zero,zero,zero
	mulss	%xmm3, %xmm4
	mulss	60(%rdi), %xmm3
	addss	%xmm6, %xmm4
	addss	%xmm0, %xmm3
	addss	40(%rdi), %xmm2
	addss	44(%rdi), %xmm4
	addss	48(%rdi), %xmm3
	movaps	%xmm2, %xmm0
	mulss	%xmm2, %xmm0
	movaps	%xmm4, %xmm5
	mulss	%xmm4, %xmm5
	addss	%xmm0, %xmm5
	movaps	%xmm3, %xmm0
	mulss	%xmm3, %xmm0
	addss	%xmm5, %xmm0
	sqrtss	%xmm0, %xmm0
	divss	%xmm0, %xmm1
	mulss	%xmm1, %xmm2
	mulss	%xmm1, %xmm4
	unpcklps	%xmm4, %xmm2                    # xmm2 = xmm2[0],xmm4[0],xmm2[1],xmm4[1]
	mulss	%xmm3, %xmm1
	movaps	%xmm2, %xmm0
	retq
.Lfunc_end4:
	.size	_ZNK6Camera16ray_dir_at_pixelEff, .Lfunc_end4-_ZNK6Camera16ray_dir_at_pixelEff
	.cfi_endproc
                                        # -- End function
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.globl	_ZN6CameraC1EiifRK4vec3S2_
	.type	_ZN6CameraC1EiifRK4vec3S2_,@function
.set _ZN6CameraC1EiifRK4vec3S2_, _ZN6CameraC2EiifRK4vec3S2_
	.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
