//
// Created by igor on 10.04.2021.
//

#include "Vector3.cuh"

__device__ __host__ Vector3::Vector3(float x, float y, float z) : x(x), y(y), z(z) {}

__device__ __host__ Vector3& Vector3::normalize() {
    float norm = this->norm();
    x /= norm;
    y /= norm;
    z /= norm;
    return *this;
}

__device__ __host__ float Vector3::norm() const{
    return sqrt(x*x+y*y+z*z);
}

__device__ __host__ float Vector3::dot(const Vector3 r) const {
    return x * r.x + y * r.y + z * r.z;
}

__device__ __host__ float Vector3::squared() const {
    return x*x+y*y+z*z;
}

__device__ __host__ Vector3 Vector3::cross(Vector3 b) const {
    const Vector3& a = *this;
    return {
        a.y * b.z - a.z * b.y,
        a.z * b.x - a.x * b.z,
        a.x * b.y - a.y * b.x
    };
}

__device__ __host__ Vector3 operator*(float s, const Vector3& v) {
    return {s*v.x, s*v.y, s*v.z};
}

__device__ __host__ Vector3 operator+(const Vector3 &l, const Vector3 &r) {
    return {l.x + r.x, l.y + r.y, l.z + r.z};
}

__device__ __host__ Vector3 operator-(const Vector3 &l, const Vector3 &r) {
    return {l.x - r.x, l.y - r.y, l.z - r.z};
}

__device__ __host__ Vector3 operator-(const Vector3 &l) {
    return -1 * l;
}

	.file	"tmpxft_00361ab3_00000000-6_Vector3.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2022:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2022:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.align 2
	.globl	_ZN7Vector3C2Efff
	.type	_ZN7Vector3C2Efff, @function
_ZN7Vector3C2Efff:
.LFB2009:
	.cfi_startproc
	endbr64
	movss	%xmm0, (%rdi)
	movss	%xmm1, 4(%rdi)
	movss	%xmm2, 8(%rdi)
	ret
	.cfi_endproc
.LFE2009:
	.size	_ZN7Vector3C2Efff, .-_ZN7Vector3C2Efff
	.globl	_ZN7Vector3C1Efff
	.set	_ZN7Vector3C1Efff,_ZN7Vector3C2Efff
	.align 2
	.globl	_ZNK7Vector34normEv
	.type	_ZNK7Vector34normEv, @function
_ZNK7Vector34normEv:
.LFB2012:
	.cfi_startproc
	endbr64
	movss	(%rdi), %xmm0
	movss	4(%rdi), %xmm2
	movss	8(%rdi), %xmm1
	mulss	%xmm0, %xmm0
	mulss	%xmm2, %xmm2
	mulss	%xmm1, %xmm1
	addss	%xmm2, %xmm0
	addss	%xmm1, %xmm0
	jmp	sqrtf@PLT
	.cfi_endproc
.LFE2012:
	.size	_ZNK7Vector34normEv, .-_ZNK7Vector34normEv
	.align 2
	.globl	_ZN7Vector39normalizeEv
	.type	_ZN7Vector39normalizeEv, @function
_ZN7Vector39normalizeEv:
.LFB2011:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movq	%rdi, %rbx
	call	_ZNK7Vector34normEv
	movss	(%rbx), %xmm1
	movq	%rbx, %rax
	divss	%xmm0, %xmm1
	movss	%xmm1, (%rbx)
	movss	4(%rbx), %xmm1
	divss	%xmm0, %xmm1
	movss	%xmm1, 4(%rbx)
	movss	8(%rbx), %xmm1
	divss	%xmm0, %xmm1
	movss	%xmm1, 8(%rbx)
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2011:
	.size	_ZN7Vector39normalizeEv, .-_ZN7Vector39normalizeEv
	.align 2
	.globl	_ZNK7Vector33dotES_
	.type	_ZNK7Vector33dotES_, @function
_ZNK7Vector33dotES_:
.LFB2013:
	.cfi_startproc
	endbr64
	movq	%xmm0, -16(%rsp)
	movss	4(%rdi), %xmm2
	movss	(%rdi), %xmm0
	mulss	-12(%rsp), %xmm2
	mulss	-16(%rsp), %xmm0
	mulss	8(%rdi), %xmm1
	addss	%xmm2, %xmm0
	addss	%xmm1, %xmm0
	ret
	.cfi_endproc
.LFE2013:
	.size	_ZNK7Vector33dotES_, .-_ZNK7Vector33dotES_
	.align 2
	.globl	_ZNK7Vector37squaredEv
	.type	_ZNK7Vector37squaredEv, @function
_ZNK7Vector37squaredEv:
.LFB2014:
	.cfi_startproc
	endbr64
	movss	(%rdi), %xmm0
	movss	4(%rdi), %xmm2
	movss	8(%rdi), %xmm1
	mulss	%xmm2, %xmm2
	mulss	%xmm0, %xmm0
	mulss	%xmm1, %xmm1
	addss	%xmm2, %xmm0
	addss	%xmm1, %xmm0
	ret
	.cfi_endproc
.LFE2014:
	.size	_ZNK7Vector37squaredEv, .-_ZNK7Vector37squaredEv
	.align 2
	.globl	_ZNK7Vector35crossES_
	.type	_ZNK7Vector35crossES_, @function
_ZNK7Vector35crossES_:
.LFB2015:
	.cfi_startproc
	endbr64
	movq	%xmm0, -40(%rsp)
	movss	4(%rdi), %xmm2
	movss	8(%rdi), %xmm0
	movss	-40(%rsp), %xmm5
	movss	-36(%rsp), %xmm6
	movss	(%rdi), %xmm3
	movaps	%xmm2, %xmm4
	movaps	%xmm0, %xmm7
	mulss	%xmm1, %xmm4
	mulss	%xmm6, %xmm7
	mulss	%xmm3, %xmm1
	mulss	%xmm5, %xmm0
	mulss	%xmm6, %xmm3
	mulss	%xmm5, %xmm2
	subss	%xmm7, %xmm4
	subss	%xmm1, %xmm0
	movaps	%xmm3, %xmm1
	movss	%xmm4, -20(%rsp)
	subss	%xmm2, %xmm1
	movss	%xmm0, -16(%rsp)
	movq	-20(%rsp), %xmm0
	ret
	.cfi_endproc
.LFE2015:
	.size	_ZNK7Vector35crossES_, .-_ZNK7Vector35crossES_
	.globl	_ZmlfRK7Vector3
	.type	_ZmlfRK7Vector3, @function
_ZmlfRK7Vector3:
.LFB2016:
	.cfi_startproc
	endbr64
	movaps	%xmm0, %xmm1
	movss	(%rdi), %xmm2
	movss	4(%rdi), %xmm0
	mulss	%xmm1, %xmm2
	mulss	%xmm1, %xmm0
	mulss	8(%rdi), %xmm1
	movss	%xmm2, -20(%rsp)
	movss	%xmm0, -16(%rsp)
	movq	-20(%rsp), %xmm0
	ret
	.cfi_endproc
.LFE2016:
	.size	_ZmlfRK7Vector3, .-_ZmlfRK7Vector3
	.globl	_ZplRK7Vector3S1_
	.type	_ZplRK7Vector3S1_, @function
_ZplRK7Vector3S1_:
.LFB2017:
	.cfi_startproc
	endbr64
	movss	4(%rdi), %xmm0
	movss	(%rdi), %xmm2
	addss	4(%rsi), %xmm0
	addss	(%rsi), %xmm2
	movss	8(%rdi), %xmm1
	addss	8(%rsi), %xmm1
	movss	%xmm2, -20(%rsp)
	movss	%xmm0, -16(%rsp)
	movq	-20(%rsp), %xmm0
	ret
	.cfi_endproc
.LFE2017:
	.size	_ZplRK7Vector3S1_, .-_ZplRK7Vector3S1_
	.globl	_ZmiRK7Vector3S1_
	.type	_ZmiRK7Vector3S1_, @function
_ZmiRK7Vector3S1_:
.LFB2018:
	.cfi_startproc
	endbr64
	movss	4(%rdi), %xmm0
	movss	(%rdi), %xmm2
	subss	4(%rsi), %xmm0
	subss	(%rsi), %xmm2
	movss	8(%rdi), %xmm1
	subss	8(%rsi), %xmm1
	movss	%xmm2, -20(%rsp)
	movss	%xmm0, -16(%rsp)
	movq	-20(%rsp), %xmm0
	ret
	.cfi_endproc
.LFE2018:
	.size	_ZmiRK7Vector3S1_, .-_ZmiRK7Vector3S1_
	.globl	_ZngRK7Vector3
	.type	_ZngRK7Vector3, @function
_ZngRK7Vector3:
.LFB2019:
	.cfi_startproc
	endbr64
	movss	.LC0(%rip), %xmm0
	jmp	_ZmlfRK7Vector3
	.cfi_endproc
.LFE2019:
	.size	_ZngRK7Vector3, .-_ZngRK7Vector3
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2045:
	.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
.LFE2045:
	.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	-1082130432
	.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:

//
// Created by igor on 10.04.2021.
//


#include <hip/hip_runtime.h>
#include "Vector3.cuh"

__device__ __host__ Vector3::Vector3(float x, float y, float z) : x(x), y(y), z(z) {}

__device__ __host__ Vector3& Vector3::normalize() {
    float norm = this->norm();
    x /= norm;
    y /= norm;
    z /= norm;
    return *this;
}

__device__ __host__ float Vector3::norm() const{
    return sqrt(x*x+y*y+z*z);
}

__device__ __host__ float Vector3::dot(const Vector3 r) const {
    return x * r.x + y * r.y + z * r.z;
}

__device__ __host__ float Vector3::squared() const {
    return x*x+y*y+z*z;
}

__device__ __host__ Vector3 Vector3::cross(Vector3 b) const {
    const Vector3& a = *this;
    return {
        a.y * b.z - a.z * b.y,
        a.z * b.x - a.x * b.z,
        a.x * b.y - a.y * b.x
    };
}

__device__ __host__ Vector3 operator*(float s, const Vector3& v) {
    return {s*v.x, s*v.y, s*v.z};
}

__device__ __host__ Vector3 operator+(const Vector3 &l, const Vector3 &r) {
    return {l.x + r.x, l.y + r.y, l.z + r.z};
}

__device__ __host__ Vector3 operator-(const Vector3 &l, const Vector3 &r) {
    return {l.x - r.x, l.y - r.y, l.z - r.z};
}

__device__ __host__ Vector3 operator-(const Vector3 &l) {
    return -1 * l;
}

	.text
	.file	"Vector3.hip"
	.globl	_ZN7Vector3C2Efff               # -- Begin function _ZN7Vector3C2Efff
	.p2align	1, 0x90
	.type	_ZN7Vector3C2Efff,@function
_ZN7Vector3C2Efff:                      # @_ZN7Vector3C2Efff
	.cfi_startproc
# %bb.0:
	movss	%xmm0, (%rdi)
	movss	%xmm1, 4(%rdi)
	movss	%xmm2, 8(%rdi)
	retq
.Lfunc_end0:
	.size	_ZN7Vector3C2Efff, .Lfunc_end0-_ZN7Vector3C2Efff
	.cfi_endproc
                                        # -- End function
	.globl	_ZN7Vector39normalizeEv         # -- Begin function _ZN7Vector39normalizeEv
	.p2align	1, 0x90
	.type	_ZN7Vector39normalizeEv,@function
_ZN7Vector39normalizeEv:                # @_ZN7Vector39normalizeEv
	.cfi_startproc
# %bb.0:
	movq	%rdi, %rax
	movss	(%rdi), %xmm0                   # xmm0 = mem[0],zero,zero,zero
	movss	4(%rdi), %xmm1                  # xmm1 = mem[0],zero,zero,zero
	movaps	%xmm0, %xmm2
	mulss	%xmm0, %xmm2
	movaps	%xmm1, %xmm3
	mulss	%xmm1, %xmm3
	addss	%xmm2, %xmm3
	movss	8(%rdi), %xmm2                  # xmm2 = mem[0],zero,zero,zero
	movaps	%xmm2, %xmm4
	mulss	%xmm2, %xmm4
	addss	%xmm3, %xmm4
	xorps	%xmm3, %xmm3
	sqrtss	%xmm4, %xmm3
	divss	%xmm3, %xmm0
	movss	%xmm0, (%rdi)
	divss	%xmm3, %xmm1
	movss	%xmm1, 4(%rdi)
	divss	%xmm3, %xmm2
	movss	%xmm2, 8(%rdi)
	retq
.Lfunc_end1:
	.size	_ZN7Vector39normalizeEv, .Lfunc_end1-_ZN7Vector39normalizeEv
	.cfi_endproc
                                        # -- End function
	.globl	_ZNK7Vector34normEv             # -- Begin function _ZNK7Vector34normEv
	.p2align	1, 0x90
	.type	_ZNK7Vector34normEv,@function
_ZNK7Vector34normEv:                    # @_ZNK7Vector34normEv
	.cfi_startproc
# %bb.0:
	movss	(%rdi), %xmm0                   # xmm0 = mem[0],zero,zero,zero
	movss	4(%rdi), %xmm1                  # xmm1 = mem[0],zero,zero,zero
	mulss	%xmm0, %xmm0
	mulss	%xmm1, %xmm1
	addss	%xmm0, %xmm1
	movss	8(%rdi), %xmm0                  # xmm0 = mem[0],zero,zero,zero
	mulss	%xmm0, %xmm0
	addss	%xmm1, %xmm0
	sqrtss	%xmm0, %xmm0
	retq
.Lfunc_end2:
	.size	_ZNK7Vector34normEv, .Lfunc_end2-_ZNK7Vector34normEv
	.cfi_endproc
                                        # -- End function
	.globl	_ZNK7Vector33dotES_             # -- Begin function _ZNK7Vector33dotES_
	.p2align	1, 0x90
	.type	_ZNK7Vector33dotES_,@function
_ZNK7Vector33dotES_:                    # @_ZNK7Vector33dotES_
	.cfi_startproc
# %bb.0:
	movss	(%rdi), %xmm2                   # xmm2 = mem[0],zero,zero,zero
	mulss	%xmm0, %xmm2
	shufps	$85, %xmm0, %xmm0               # xmm0 = xmm0[1,1,1,1]
	mulss	4(%rdi), %xmm0
	addss	%xmm2, %xmm0
	mulss	8(%rdi), %xmm1
	addss	%xmm1, %xmm0
	retq
.Lfunc_end3:
	.size	_ZNK7Vector33dotES_, .Lfunc_end3-_ZNK7Vector33dotES_
	.cfi_endproc
                                        # -- End function
	.globl	_ZNK7Vector37squaredEv          # -- Begin function _ZNK7Vector37squaredEv
	.p2align	1, 0x90
	.type	_ZNK7Vector37squaredEv,@function
_ZNK7Vector37squaredEv:                 # @_ZNK7Vector37squaredEv
	.cfi_startproc
# %bb.0:
	movss	(%rdi), %xmm0                   # xmm0 = mem[0],zero,zero,zero
	movss	4(%rdi), %xmm1                  # xmm1 = mem[0],zero,zero,zero
	mulss	%xmm0, %xmm0
	mulss	%xmm1, %xmm1
	addss	%xmm0, %xmm1
	movss	8(%rdi), %xmm0                  # xmm0 = mem[0],zero,zero,zero
	mulss	%xmm0, %xmm0
	addss	%xmm1, %xmm0
	retq
.Lfunc_end4:
	.size	_ZNK7Vector37squaredEv, .Lfunc_end4-_ZNK7Vector37squaredEv
	.cfi_endproc
                                        # -- End function
	.globl	_ZNK7Vector35crossES_           # -- Begin function _ZNK7Vector35crossES_
	.p2align	1, 0x90
	.type	_ZNK7Vector35crossES_,@function
_ZNK7Vector35crossES_:                  # @_ZNK7Vector35crossES_
	.cfi_startproc
# %bb.0:
	movss	(%rdi), %xmm4                   # xmm4 = mem[0],zero,zero,zero
	movss	4(%rdi), %xmm3                  # xmm3 = mem[0],zero,zero,zero
	movss	8(%rdi), %xmm5                  # xmm5 = mem[0],zero,zero,zero
	movaps	%xmm0, %xmm2
	shufps	$85, %xmm0, %xmm2               # xmm2 = xmm2[1,1],xmm0[1,1]
	movaps	%xmm2, %xmm6
	mulss	%xmm5, %xmm6
	mulss	%xmm0, %xmm5
	mulss	%xmm3, %xmm0
	mulss	%xmm1, %xmm3
	subss	%xmm6, %xmm3
	mulss	%xmm4, %xmm1
	subss	%xmm1, %xmm5
	unpcklps	%xmm5, %xmm3                    # xmm3 = xmm3[0],xmm5[0],xmm3[1],xmm5[1]
	mulss	%xmm4, %xmm2
	subss	%xmm0, %xmm2
	movaps	%xmm3, %xmm0
	movaps	%xmm2, %xmm1
	retq
.Lfunc_end5:
	.size	_ZNK7Vector35crossES_, .Lfunc_end5-_ZNK7Vector35crossES_
	.cfi_endproc
                                        # -- End function
	.globl	_ZmlfRK7Vector3                 # -- Begin function _ZmlfRK7Vector3
	.type	_ZmlfRK7Vector3,@function
_ZmlfRK7Vector3:                        # @_ZmlfRK7Vector3
	.cfi_startproc
# %bb.0:
	movaps	%xmm0, %xmm1
	movss	(%rdi), %xmm0                   # xmm0 = mem[0],zero,zero,zero
	mulss	%xmm1, %xmm0
	movss	4(%rdi), %xmm2                  # xmm2 = mem[0],zero,zero,zero
	mulss	%xmm1, %xmm2
	unpcklps	%xmm2, %xmm0                    # xmm0 = xmm0[0],xmm2[0],xmm0[1],xmm2[1]
	mulss	8(%rdi), %xmm1
	retq
.Lfunc_end6:
	.size	_ZmlfRK7Vector3, .Lfunc_end6-_ZmlfRK7Vector3
	.cfi_endproc
                                        # -- End function
	.globl	_ZplRK7Vector3S1_               # -- Begin function _ZplRK7Vector3S1_
	.type	_ZplRK7Vector3S1_,@function
_ZplRK7Vector3S1_:                      # @_ZplRK7Vector3S1_
	.cfi_startproc
# %bb.0:
	movss	(%rdi), %xmm0                   # xmm0 = mem[0],zero,zero,zero
	movss	4(%rdi), %xmm1                  # xmm1 = mem[0],zero,zero,zero
	addss	(%rsi), %xmm0
	addss	4(%rsi), %xmm1
	unpcklps	%xmm1, %xmm0                    # xmm0 = xmm0[0],xmm1[0],xmm0[1],xmm1[1]
	movss	8(%rdi), %xmm1                  # xmm1 = mem[0],zero,zero,zero
	addss	8(%rsi), %xmm1
	retq
.Lfunc_end7:
	.size	_ZplRK7Vector3S1_, .Lfunc_end7-_ZplRK7Vector3S1_
	.cfi_endproc
                                        # -- End function
	.globl	_ZmiRK7Vector3S1_               # -- Begin function _ZmiRK7Vector3S1_
	.type	_ZmiRK7Vector3S1_,@function
_ZmiRK7Vector3S1_:                      # @_ZmiRK7Vector3S1_
	.cfi_startproc
# %bb.0:
	movss	(%rdi), %xmm0                   # xmm0 = mem[0],zero,zero,zero
	movss	4(%rdi), %xmm1                  # xmm1 = mem[0],zero,zero,zero
	subss	(%rsi), %xmm0
	subss	4(%rsi), %xmm1
	unpcklps	%xmm1, %xmm0                    # xmm0 = xmm0[0],xmm1[0],xmm0[1],xmm1[1]
	movss	8(%rdi), %xmm1                  # xmm1 = mem[0],zero,zero,zero
	subss	8(%rsi), %xmm1
	retq
.Lfunc_end8:
	.size	_ZmiRK7Vector3S1_, .Lfunc_end8-_ZmiRK7Vector3S1_
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst16,"aM",@progbits,16
	.p2align	4, 0x0                          # -- Begin function _ZngRK7Vector3
.LCPI9_0:
	.long	0x80000000                      # float -0
	.long	0x80000000                      # float -0
	.long	0x80000000                      # float -0
	.long	0x80000000                      # float -0
	.text
	.globl	_ZngRK7Vector3
	.type	_ZngRK7Vector3,@function
_ZngRK7Vector3:                         # @_ZngRK7Vector3
	.cfi_startproc
# %bb.0:
	movaps	.LCPI9_0(%rip), %xmm2           # xmm2 = [-0.0E+0,-0.0E+0,-0.0E+0,-0.0E+0]
	movss	(%rdi), %xmm0                   # xmm0 = mem[0],zero,zero,zero
	movss	4(%rdi), %xmm3                  # xmm3 = mem[0],zero,zero,zero
	movss	8(%rdi), %xmm1                  # xmm1 = mem[0],zero,zero,zero
	xorps	%xmm2, %xmm0
	xorps	%xmm2, %xmm3
	unpcklps	%xmm3, %xmm0                    # xmm0 = xmm0[0],xmm3[0],xmm0[1],xmm3[1]
	xorps	%xmm2, %xmm1
	retq
.Lfunc_end9:
	.size	_ZngRK7Vector3, .Lfunc_end9-_ZngRK7Vector3
	.cfi_endproc
                                        # -- End function
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.globl	_ZN7Vector3C1Efff
	.type	_ZN7Vector3C1Efff,@function
.set _ZN7Vector3C1Efff, _ZN7Vector3C2Efff
	.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
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.hidden	_ZN7Vector3C1Efff               ; -- Begin function _ZN7Vector3C1Efff
	.globl	_ZN7Vector3C1Efff
	.p2align	2
	.type	_ZN7Vector3C1Efff,@function
_ZN7Vector3C1Efff:                      ; @_ZN7Vector3C1Efff
; %bb.0:
	s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
	flat_store_b96 v[0:1], v[2:4]
	s_waitcnt lgkmcnt(0)
	s_setpc_b64 s[30:31]
.Lfunc_end0:
	.size	_ZN7Vector3C1Efff, .Lfunc_end0-_ZN7Vector3C1Efff
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Function info:
; codeLenInByte = 20
; NumSgprs: 32
; NumVgprs: 5
; ScratchSize: 0
; MemoryBound: 0
	.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_
	.amdgpu_metadata
---
amdhsa.kernels:  []
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
