#include <stdio.h>
#include <math.h>
#include <stdlib.h>

__host__ __device__ double2 d2add(double2 a, double2 b) {
	/*
	 * Arguments: two 2d vectors
	 * Returns: the vector addition of the two vectors
	 */
	double2 ret;
	ret.x=a.x+b.x;
	ret.y=a.y+b.y;
	return ret;
}

__host__ __device__ double2 d2sub(double2 a, double2 b) {
	/*
	 * Arguments: two 2d vectors
	 * Returns: the vector subtraction of the two vectors
	 */
	double2 ret;
	ret.x=a.x-b.x;
	ret.y=a.y-b.y;
	return ret;
}

__host__ __device__ double2 d2multscal(double2 a, double b) {
	/*
	 * Arguments: a 2d vector and a scalar
	 * Returns: the vector with each component multiplied by the scalar
	 */
	double2 ret;
	ret.x=b*a.x;
	ret.y=b*a.y;
	return ret;
}

__host__ __device__ double2 d2multscal2(double2 a, double b) {
	/*
	 * Arguments: a 2d vector and a scalar
	 * Returns: the vector with each component multiplied by the scalar squared
	 */
	double2 ret;
	ret.x=b*b*a.x;
	ret.y=b*b*a.y;
	return ret;
}

__host__ __device__ double2 d2divscal(double2 a, double b) {
	/*
	 * Arguments: a 2d vector and a scalar
	 * Returns: the vector with each component divided by the scalar
	 */
	double2 ret;
	ret.x=a.x/b;
	ret.y=a.y/b;
	return ret;
}

__host__ __device__ double d2dotp(double2 a, double2 b) {
	/*
	 * Arguments: two 2d vectors
	 * Returns: the dot product between the two vectors
	 */
	return (a.x*b.x)+(a.y*b.y);
}

__host__ __device__ double d2mag(double2 a) {
	/*
	 * Arguments: one 2d vector
	 * Returns: the magnitude of the argument vector
	 */
	return sqrt((a.x)*(a.x)+(a.y)*(a.y));
}

__host__ __device__ double d2dist(double2 a, double2 b) {
	/*
	 * Arguments: two 2d vectors
	 * Returns: the magnitude of the first vector minus the second vector
	 */
	return d2mag(d2sub(a,b));
}

__host__ __device__ double2 d2unit(double2 a) {
	/*
	 * Arguments: one 2d vector
	 * Returns: a unit vector pointing in the same direction as the argument vector
	 */ 
	return d2divscal(a, d2mag(a));
}

__host__ __device__ double2 d2null() {
	/*
	 * Arguments: none
	 * Returns: the null vector (2d)
	 */
	double2 a;
	a.x=0.0;
	a.y=0.0;
	return a;
}


	.file	"tmpxft_00217b4f_00000000-6_d2func.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2039:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2039:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z5d2add7double2S_
	.type	_Z5d2add7double2S_, @function
_Z5d2add7double2S_:
.LFB2027:
	.cfi_startproc
	endbr64
	addsd	%xmm3, %xmm1
	addsd	%xmm2, %xmm0
	ret
	.cfi_endproc
.LFE2027:
	.size	_Z5d2add7double2S_, .-_Z5d2add7double2S_
	.globl	_Z5d2sub7double2S_
	.type	_Z5d2sub7double2S_, @function
_Z5d2sub7double2S_:
.LFB2028:
	.cfi_startproc
	endbr64
	subsd	%xmm3, %xmm1
	subsd	%xmm2, %xmm0
	ret
	.cfi_endproc
.LFE2028:
	.size	_Z5d2sub7double2S_, .-_Z5d2sub7double2S_
	.globl	_Z10d2multscal7double2d
	.type	_Z10d2multscal7double2d, @function
_Z10d2multscal7double2d:
.LFB2029:
	.cfi_startproc
	endbr64
	mulsd	%xmm2, %xmm1
	mulsd	%xmm2, %xmm0
	ret
	.cfi_endproc
.LFE2029:
	.size	_Z10d2multscal7double2d, .-_Z10d2multscal7double2d
	.globl	_Z11d2multscal27double2d
	.type	_Z11d2multscal27double2d, @function
_Z11d2multscal27double2d:
.LFB2030:
	.cfi_startproc
	endbr64
	mulsd	%xmm2, %xmm2
	mulsd	%xmm2, %xmm1
	mulsd	%xmm2, %xmm0
	ret
	.cfi_endproc
.LFE2030:
	.size	_Z11d2multscal27double2d, .-_Z11d2multscal27double2d
	.globl	_Z9d2divscal7double2d
	.type	_Z9d2divscal7double2d, @function
_Z9d2divscal7double2d:
.LFB2031:
	.cfi_startproc
	endbr64
	divsd	%xmm2, %xmm1
	divsd	%xmm2, %xmm0
	ret
	.cfi_endproc
.LFE2031:
	.size	_Z9d2divscal7double2d, .-_Z9d2divscal7double2d
	.globl	_Z6d2dotp7double2S_
	.type	_Z6d2dotp7double2S_, @function
_Z6d2dotp7double2S_:
.LFB2032:
	.cfi_startproc
	endbr64
	mulsd	%xmm2, %xmm0
	mulsd	%xmm3, %xmm1
	addsd	%xmm1, %xmm0
	ret
	.cfi_endproc
.LFE2032:
	.size	_Z6d2dotp7double2S_, .-_Z6d2dotp7double2S_
	.globl	_Z5d2mag7double2
	.type	_Z5d2mag7double2, @function
_Z5d2mag7double2:
.LFB2033:
	.cfi_startproc
	endbr64
	mulsd	%xmm0, %xmm0
	mulsd	%xmm1, %xmm1
	addsd	%xmm1, %xmm0
	jmp	sqrt@PLT
	.cfi_endproc
.LFE2033:
	.size	_Z5d2mag7double2, .-_Z5d2mag7double2
	.globl	_Z6d2dist7double2S_
	.type	_Z6d2dist7double2S_, @function
_Z6d2dist7double2S_:
.LFB2034:
	.cfi_startproc
	endbr64
	subsd	%xmm2, %xmm0
	subsd	%xmm3, %xmm1
	jmp	_Z5d2mag7double2
	.cfi_endproc
.LFE2034:
	.size	_Z6d2dist7double2S_, .-_Z6d2dist7double2S_
	.globl	_Z6d2unit7double2
	.type	_Z6d2unit7double2, @function
_Z6d2unit7double2:
.LFB2035:
	.cfi_startproc
	endbr64
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movsd	%xmm1, (%rsp)
	movsd	%xmm0, 8(%rsp)
	call	_Z5d2mag7double2
	movsd	8(%rsp), %xmm2
	movsd	(%rsp), %xmm1
	addq	$24, %rsp
	.cfi_def_cfa_offset 8
	movaps	%xmm0, %xmm3
	divsd	%xmm0, %xmm2
	divsd	%xmm3, %xmm1
	movaps	%xmm2, %xmm0
	ret
	.cfi_endproc
.LFE2035:
	.size	_Z6d2unit7double2, .-_Z6d2unit7double2
	.globl	_Z6d2nullv
	.type	_Z6d2nullv, @function
_Z6d2nullv:
.LFB2036:
	.cfi_startproc
	endbr64
	xorps	%xmm0, %xmm0
	movaps	%xmm0, %xmm1
	ret
	.cfi_endproc
.LFE2036:
	.size	_Z6d2nullv, .-_Z6d2nullv
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2062:
	.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
.LFE2062:
	.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
	.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 <stdio.h>
#include <math.h>
#include <stdlib.h>

__host__ __device__ double2 d2add(double2 a, double2 b) {
	/*
	 * Arguments: two 2d vectors
	 * Returns: the vector addition of the two vectors
	 */
	double2 ret;
	ret.x=a.x+b.x;
	ret.y=a.y+b.y;
	return ret;
}

__host__ __device__ double2 d2sub(double2 a, double2 b) {
	/*
	 * Arguments: two 2d vectors
	 * Returns: the vector subtraction of the two vectors
	 */
	double2 ret;
	ret.x=a.x-b.x;
	ret.y=a.y-b.y;
	return ret;
}

__host__ __device__ double2 d2multscal(double2 a, double b) {
	/*
	 * Arguments: a 2d vector and a scalar
	 * Returns: the vector with each component multiplied by the scalar
	 */
	double2 ret;
	ret.x=b*a.x;
	ret.y=b*a.y;
	return ret;
}

__host__ __device__ double2 d2multscal2(double2 a, double b) {
	/*
	 * Arguments: a 2d vector and a scalar
	 * Returns: the vector with each component multiplied by the scalar squared
	 */
	double2 ret;
	ret.x=b*b*a.x;
	ret.y=b*b*a.y;
	return ret;
}

__host__ __device__ double2 d2divscal(double2 a, double b) {
	/*
	 * Arguments: a 2d vector and a scalar
	 * Returns: the vector with each component divided by the scalar
	 */
	double2 ret;
	ret.x=a.x/b;
	ret.y=a.y/b;
	return ret;
}

__host__ __device__ double d2dotp(double2 a, double2 b) {
	/*
	 * Arguments: two 2d vectors
	 * Returns: the dot product between the two vectors
	 */
	return (a.x*b.x)+(a.y*b.y);
}

__host__ __device__ double d2mag(double2 a) {
	/*
	 * Arguments: one 2d vector
	 * Returns: the magnitude of the argument vector
	 */
	return sqrt((a.x)*(a.x)+(a.y)*(a.y));
}

__host__ __device__ double d2dist(double2 a, double2 b) {
	/*
	 * Arguments: two 2d vectors
	 * Returns: the magnitude of the first vector minus the second vector
	 */
	return d2mag(d2sub(a,b));
}

__host__ __device__ double2 d2unit(double2 a) {
	/*
	 * Arguments: one 2d vector
	 * Returns: a unit vector pointing in the same direction as the argument vector
	 */ 
	return d2divscal(a, d2mag(a));
}

__host__ __device__ double2 d2null() {
	/*
	 * Arguments: none
	 * Returns: the null vector (2d)
	 */
	double2 a;
	a.x=0.0;
	a.y=0.0;
	return a;
}


	.text
	.file	"d2func.hip"
	.globl	_Z5d2add15HIP_vector_typeIdLj2EES0_ # -- Begin function _Z5d2add15HIP_vector_typeIdLj2EES0_
	.type	_Z5d2add15HIP_vector_typeIdLj2EES0_,@function
_Z5d2add15HIP_vector_typeIdLj2EES0_:    # @_Z5d2add15HIP_vector_typeIdLj2EES0_
	.cfi_startproc
# %bb.0:
	addsd	%xmm2, %xmm0
	addsd	%xmm3, %xmm1
	retq
.Lfunc_end0:
	.size	_Z5d2add15HIP_vector_typeIdLj2EES0_, .Lfunc_end0-_Z5d2add15HIP_vector_typeIdLj2EES0_
	.cfi_endproc
                                        # -- End function
	.globl	_Z5d2sub15HIP_vector_typeIdLj2EES0_ # -- Begin function _Z5d2sub15HIP_vector_typeIdLj2EES0_
	.type	_Z5d2sub15HIP_vector_typeIdLj2EES0_,@function
_Z5d2sub15HIP_vector_typeIdLj2EES0_:    # @_Z5d2sub15HIP_vector_typeIdLj2EES0_
	.cfi_startproc
# %bb.0:
	subsd	%xmm2, %xmm0
	subsd	%xmm3, %xmm1
	retq
.Lfunc_end1:
	.size	_Z5d2sub15HIP_vector_typeIdLj2EES0_, .Lfunc_end1-_Z5d2sub15HIP_vector_typeIdLj2EES0_
	.cfi_endproc
                                        # -- End function
	.globl	_Z10d2multscal15HIP_vector_typeIdLj2EEd # -- Begin function _Z10d2multscal15HIP_vector_typeIdLj2EEd
	.type	_Z10d2multscal15HIP_vector_typeIdLj2EEd,@function
_Z10d2multscal15HIP_vector_typeIdLj2EEd: # @_Z10d2multscal15HIP_vector_typeIdLj2EEd
	.cfi_startproc
# %bb.0:
	mulsd	%xmm2, %xmm0
	mulsd	%xmm2, %xmm1
	retq
.Lfunc_end2:
	.size	_Z10d2multscal15HIP_vector_typeIdLj2EEd, .Lfunc_end2-_Z10d2multscal15HIP_vector_typeIdLj2EEd
	.cfi_endproc
                                        # -- End function
	.globl	_Z11d2multscal215HIP_vector_typeIdLj2EEd # -- Begin function _Z11d2multscal215HIP_vector_typeIdLj2EEd
	.type	_Z11d2multscal215HIP_vector_typeIdLj2EEd,@function
_Z11d2multscal215HIP_vector_typeIdLj2EEd: # @_Z11d2multscal215HIP_vector_typeIdLj2EEd
	.cfi_startproc
# %bb.0:
	mulsd	%xmm2, %xmm2
	mulsd	%xmm2, %xmm0
	mulsd	%xmm2, %xmm1
	retq
.Lfunc_end3:
	.size	_Z11d2multscal215HIP_vector_typeIdLj2EEd, .Lfunc_end3-_Z11d2multscal215HIP_vector_typeIdLj2EEd
	.cfi_endproc
                                        # -- End function
	.globl	_Z9d2divscal15HIP_vector_typeIdLj2EEd # -- Begin function _Z9d2divscal15HIP_vector_typeIdLj2EEd
	.type	_Z9d2divscal15HIP_vector_typeIdLj2EEd,@function
_Z9d2divscal15HIP_vector_typeIdLj2EEd:  # @_Z9d2divscal15HIP_vector_typeIdLj2EEd
	.cfi_startproc
# %bb.0:
	divsd	%xmm2, %xmm0
	divsd	%xmm2, %xmm1
	retq
.Lfunc_end4:
	.size	_Z9d2divscal15HIP_vector_typeIdLj2EEd, .Lfunc_end4-_Z9d2divscal15HIP_vector_typeIdLj2EEd
	.cfi_endproc
                                        # -- End function
	.globl	_Z6d2dotp15HIP_vector_typeIdLj2EES0_ # -- Begin function _Z6d2dotp15HIP_vector_typeIdLj2EES0_
	.type	_Z6d2dotp15HIP_vector_typeIdLj2EES0_,@function
_Z6d2dotp15HIP_vector_typeIdLj2EES0_:   # @_Z6d2dotp15HIP_vector_typeIdLj2EES0_
	.cfi_startproc
# %bb.0:
	mulsd	%xmm2, %xmm0
	mulsd	%xmm3, %xmm1
	addsd	%xmm1, %xmm0
	retq
.Lfunc_end5:
	.size	_Z6d2dotp15HIP_vector_typeIdLj2EES0_, .Lfunc_end5-_Z6d2dotp15HIP_vector_typeIdLj2EES0_
	.cfi_endproc
                                        # -- End function
	.globl	_Z5d2mag15HIP_vector_typeIdLj2EE # -- Begin function _Z5d2mag15HIP_vector_typeIdLj2EE
	.type	_Z5d2mag15HIP_vector_typeIdLj2EE,@function
_Z5d2mag15HIP_vector_typeIdLj2EE:       # @_Z5d2mag15HIP_vector_typeIdLj2EE
	.cfi_startproc
# %bb.0:
	mulsd	%xmm0, %xmm0
	mulsd	%xmm1, %xmm1
	addsd	%xmm0, %xmm1
	xorps	%xmm0, %xmm0
	sqrtsd	%xmm1, %xmm0
	retq
.Lfunc_end6:
	.size	_Z5d2mag15HIP_vector_typeIdLj2EE, .Lfunc_end6-_Z5d2mag15HIP_vector_typeIdLj2EE
	.cfi_endproc
                                        # -- End function
	.globl	_Z6d2dist15HIP_vector_typeIdLj2EES0_ # -- Begin function _Z6d2dist15HIP_vector_typeIdLj2EES0_
	.type	_Z6d2dist15HIP_vector_typeIdLj2EES0_,@function
_Z6d2dist15HIP_vector_typeIdLj2EES0_:   # @_Z6d2dist15HIP_vector_typeIdLj2EES0_
	.cfi_startproc
# %bb.0:
	subsd	%xmm2, %xmm0
	subsd	%xmm3, %xmm1
	mulsd	%xmm0, %xmm0
	mulsd	%xmm1, %xmm1
	addsd	%xmm0, %xmm1
	xorps	%xmm0, %xmm0
	sqrtsd	%xmm1, %xmm0
	retq
.Lfunc_end7:
	.size	_Z6d2dist15HIP_vector_typeIdLj2EES0_, .Lfunc_end7-_Z6d2dist15HIP_vector_typeIdLj2EES0_
	.cfi_endproc
                                        # -- End function
	.globl	_Z6d2unit15HIP_vector_typeIdLj2EE # -- Begin function _Z6d2unit15HIP_vector_typeIdLj2EE
	.type	_Z6d2unit15HIP_vector_typeIdLj2EE,@function
_Z6d2unit15HIP_vector_typeIdLj2EE:      # @_Z6d2unit15HIP_vector_typeIdLj2EE
	.cfi_startproc
# %bb.0:
	movapd	%xmm0, %xmm2
	mulsd	%xmm0, %xmm2
	movapd	%xmm1, %xmm3
	mulsd	%xmm1, %xmm3
	addsd	%xmm2, %xmm3
	xorps	%xmm2, %xmm2
	sqrtsd	%xmm3, %xmm2
	divsd	%xmm2, %xmm0
	divsd	%xmm2, %xmm1
	retq
.Lfunc_end8:
	.size	_Z6d2unit15HIP_vector_typeIdLj2EE, .Lfunc_end8-_Z6d2unit15HIP_vector_typeIdLj2EE
	.cfi_endproc
                                        # -- End function
	.globl	_Z6d2nullv                      # -- Begin function _Z6d2nullv
	.type	_Z6d2nullv,@function
_Z6d2nullv:                             # @_Z6d2nullv
	.cfi_startproc
# %bb.0:
	xorps	%xmm0, %xmm0
	xorps	%xmm1, %xmm1
	retq
.Lfunc_end9:
	.size	_Z6d2nullv, .Lfunc_end9-_Z6d2nullv
	.cfi_endproc
                                        # -- End function
	.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
