
#include "planeslicer.cuh"

//#include "stdio.h"

#define flt_eps  1.192092896e-07f


__device__ float dot(float3 a, float3 b)
{
	return (a.x * b.x + a.y * b.y + a.z * b.z);
}

__global__ void computeIntersectionWithPlane_Kernel( int numTriangles , float3 *triangles  , float3 *edges  , int *numEdges , float3 planePosition , float3 planeNormal )
{
   int tid = threadIdx.x + blockIdx.x * blockDim.x;
   
   if( tid > numTriangles - 1 )
      return;

   float3 *t3 = triangles + tid * 3;

   extern __shared__ float sData[];

   int *nIntersections = (int*)sData;
   int *tEid = nIntersections + 1;

   if (threadIdx.x == 0)
   {
	   *nIntersections = 0;

	   *tEid = 0;
   }

   __syncthreads();

   float3 *sIntersections = (float3*)( sData + 2 );
	  
   int p = 0, n = 0;

   float3 d1, d2, d3, l1, l2, l3;

   d1.x = planePosition.x - t3[0].x;
   d1.y = planePosition.y - t3[0].y;
   d1.z = planePosition.z - t3[0].z;

   d2.x = planePosition.x - t3[1].x;
   d2.y = planePosition.y - t3[1].y;
   d2.z = planePosition.z - t3[1].z;

   d3.x = planePosition.x - t3[2].x;
   d3.y = planePosition.y - t3[2].y;
   d3.z = planePosition.z - t3[2].z;

   l1.x = t3[0].x - t3[1].x;
   l1.y = t3[0].y - t3[1].y;
   l1.z = t3[0].z - t3[1].z;

   l2.x = t3[1].x - t3[2].x;
   l2.y = t3[1].y - t3[2].y;
   l2.z = t3[1].z - t3[2].z;

   l3.x = t3[2].x - t3[0].x;
   l3.y = t3[2].y - t3[0].y;
   l3.z = t3[2].z - t3[0].z;

   float f1 =  dot( d1 , planeNormal );
   float f2 =  dot( d2 , planeNormal );
   float f3 =  dot( d3 , planeNormal );

   p += f1 > 0;
   p += f2 > 0;
   p += f3 > 0;

   n += f1 < 0;
   n += f2 < 0;
   n += f3 < 0;

   if (p < 3 && p > 0)
   {

	   int id = atomicAdd(nIntersections, 1);

	   float lambda1 = f2 / dot(l1, planeNormal);
	   float lambda2 = f3 / dot(l2, planeNormal);
	   float lambda3 = f1 / dot(l3, planeNormal);

	   bool p1Found = false, p2Found = false;

	   if (lambda1 >= 0 && lambda1 <= 1)
	   {
		   sIntersections[2 * id].x = lambda1 * t3[0].x + (1 - lambda1) * t3[1].x;
		   sIntersections[2 * id].y = lambda1 * t3[0].y + (1 - lambda1) * t3[1].y;
		   sIntersections[2 * id].z = lambda1 * t3[0].z + (1 - lambda1) * t3[1].z;

		   p1Found = true;
	   }

	   if (lambda2 >= 0 && lambda2 <= 1)
	   {
		   if (!p1Found)
		   {
			   sIntersections[2 * id].x = lambda2 * t3[1].x + (1 - lambda2) * t3[2].x;
			   sIntersections[2 * id].y = lambda2 * t3[1].y + (1 - lambda2) * t3[2].y;
			   sIntersections[2 * id].z = lambda2 * t3[1].z + (1 - lambda2) * t3[2].z;

			   p1Found = true;

		   }
		   else
		   {
			   sIntersections[2 * id + 1].x = lambda2 * t3[1].x + (1 - lambda2) * t3[2].x;
			   sIntersections[2 * id + 1].y = lambda2 * t3[1].y + (1 - lambda2) * t3[2].y;
			   sIntersections[2 * id + 1].z = lambda2 * t3[1].z + (1 - lambda2) * t3[2].z;

			   p2Found = true;
		   }
	   }

	   if (lambda3 >= 0 && lambda3 <= 1)
	   {
		   if (!p2Found)
		   {
			   sIntersections[2 * id + 1].x = lambda3 * t3[2].x + (1 - lambda3) * t3[0].x;
			   sIntersections[2 * id + 1].y = lambda3 * t3[2].y + (1 - lambda3) * t3[0].y;
			   sIntersections[2 * id + 1].z = lambda3 * t3[2].z + (1 - lambda3) * t3[0].z;

			   p2Found = true;
		   }
	   }

	   if (!p1Found )
	   {
		   sIntersections[2 * id].x = sIntersections[2 * id + 1].x;
		   sIntersections[2 * id].y = sIntersections[2 * id + 1].y;
		   sIntersections[2 * id].z = sIntersections[2 * id + 1].z;
	   }

	   if (!p2Found)
	   {
		   sIntersections[2 * id + 1].x = sIntersections[2 * id].x;
		   sIntersections[2 * id + 1].y = sIntersections[2 * id].y;
		   sIntersections[2 * id + 1].z = sIntersections[2 * id].z;
	   }
   }


   __syncthreads();

   if (threadIdx.x == 0)
   {
	  *tEid = atomicAdd( numEdges , *nIntersections ) ;
   }
   
   __syncthreads();

   if ( *nIntersections > 0 && threadIdx.x < *nIntersections )
   {
	   edges[2 * (*tEid + threadIdx.x)]  = sIntersections[2 * threadIdx.x];
	   edges[2 * (*tEid + threadIdx.x) + 1] = sIntersections[2 * threadIdx.x + 1];
   }

}



__global__ void computeIntersectionWithStandardPlane_Kernel( int numTriangles, float3 *triangles , float3* direction , float3 *edges , float3 *edgePointNormals , float3* collectedTris , 
	                                                         int *numEdges, float3 sliceValue , int planeType )
{
	int tid = threadIdx.x + blockIdx.x * blockDim.x;

	if ( tid > numTriangles - 1 )
		return;


	extern __shared__ float sData[];

	int *nIntersections = (int*)sData;
	int *tEid = nIntersections + 1;

	if ( threadIdx.x == 0 )
	{
		*nIntersections = 0;

		*tEid = 0;
	}

	__syncthreads();

	float sv = 0;

	if ( planeType == 0 )
	{
		sv = sliceValue.x;
	}
	else if ( planeType == 1 )
	{
		sv = sliceValue.y;
	}
	else
	{
		sv = sliceValue.z;
	}

	
	float *t = (float*)( triangles + 3 * tid );

	float3 *n3 = direction + 3 * tid;

	float v1 = t[ planeType ] - sv;
	float v2 = t[3 + planeType] - sv;
	float v3 = t[6 + planeType] - sv;

	bool notIntersected = (v1 > 0 && v2 > 0 && v3 > 0) || (v1 < 0 && v2 < 0 && v3 < 0);

	if ( notIntersected )
	{
		return;
	}

	float t1 = v1 / ( t[planeType] - t[3 + planeType] );
	float t2 = v2 / ( t[3 + planeType] - t[6 + planeType]);
	float t3 = v3 / ( t[6 + planeType] - t[planeType] );
	
	float3 end1, end2;
	float interpN1[3], interpN2[3];

	if ( t1 > 0 && t1 < 1 )
	{
		end1.x = t[ 0 ] * t1 + t[ 3 ] * (1 - t1);
		end1.y = t[ 1 ] * t1 + t[ 4 ] * (1 - t1);
		end1.z = t[ 2 ] * t1 + t[ 5 ] * (1 - t1);

		interpN1[0] = n3[0].x * t1 + n3[1].x * (1 - t1);
		interpN1[1] = n3[0].y * t1 + n3[1].y * (1 - t1);
		interpN1[2] = n3[0].z * t1 + n3[1].z * (1 - t1);

		//interpWt1 = wt[0] * t1 + wt[1] * (1 - t1);

		if (t2 > 0 && t2 < 1)
		{
			//end2 = pos2 * t2 + pos3 * (1 - t2);

			end2.x = t[3] * t2 + t[6] * (1 - t2);
			end2.y = t[4] * t2 + t[7] * (1 - t2);
			end2.z = t[5] * t2 + t[8] * (1 - t2);

			interpN2[0] = n3[1].x * t2 + n3[2].x * (1 - t2);
			interpN2[1] = n3[1].y * t2 + n3[2].y * (1 - t2);
			interpN2[2] = n3[1].z * t2 + n3[2].z * (1 - t2);

		}
		else
		{
			//end2 = pos3 * t3 + pos1 * (1 - t3);

			end2.x = t[6] * t3 + t[0] * (1 - t3);
			end2.y = t[7] * t3 + t[1] * (1 - t3);
			end2.z = t[8] * t3 + t[2] * (1 - t3);

			//interpN2 = n3 * t3 + n1 * (1 - t3);

			interpN2[0] = n3[1].x * t3 + n3[2].x * (1 - t3);
			interpN2[1] = n3[1].y * t3 + n3[2].y * (1 - t3);
			interpN2[2] = n3[1].z * t3 + n3[2].z * (1 - t3);

		}

	}
	else if (t2 > 0 && t2 < 1)
	{

		end1.x = t[3] * t2 + t[6] * (1 - t2);
		end1.y = t[4] * t2 + t[7] * (1 - t2);
		end1.z = t[5] * t2 + t[8] * (1 - t2);

		end2.x = t[3] * t2 + t[6] * (1 - t2);
		end2.y = t[4] * t2 + t[7] * (1 - t2);
		end2.z = t[5] * t2 + t[8] * (1 - t2);

		//end1 = pos2 * t2 + pos3 * (1 - t2);
		//end2 = pos3 * t3 + pos1 * (1 - t3);

		//interpN1 = n2 * t2 + n3 * (1 - t2);
		//interpN2 = n3 * t3 + n1 * (1 - t3);

		interpN1[0] = n3[1].x * t2 + n3[2].x * (1 - t2);
		interpN1[1] = n3[1].y * t2 + n3[2].y * (1 - t2);
		interpN1[2] = n3[1].z * t2 + n3[2].z * (1 - t2);

		interpN2[0] = n3[2].x * t3 + n3[0].x * (1 - t3);
		interpN2[1] = n3[2].y * t3 + n3[0].y * (1 - t3);
		interpN2[2] = n3[2].z * t3 + n3[0].z * (1 - t3);


	}
	
	int id = atomicAdd( nIntersections , 2 );

	edges[2 * id] = end1;
	edges[2 * id + 1] = end2;

	edgePointNormals[2 * id].x = interpN1[0];
	edgePointNormals[2 * id].y = interpN1[1];
	edgePointNormals[2 * id].z = interpN1[2];

	edgePointNormals[2 * id + 1].x = interpN2[0];
	edgePointNormals[2 * id + 1].y = interpN2[1];
	edgePointNormals[2 * id + 1].z = interpN2[2];

	__syncthreads();

	atomicAdd( numEdges , *nIntersections );
	

}


cudaError_t computeIntersectionWithStandardPlane(int numTriangles, float3 *triangles, float3 *edges, float3 *oppositeEnds , float3* collectedTris, int *numEdges, float3 sliceValue, int planeType)
{

	dim3 threads(128, 1);

	int wB = (numTriangles + threads.x - 1) / threads.x;
	int hB = 1;

	dim3 blocks(wB, hB);

	int sharedMemorySize = threads.x * 2 * sizeof(float3) + 2 * sizeof(int);

	//computeIntersectionWithPlane_Kernel << < blocks, threads, sharedMemorySize >> >(numTriangles, triangles, edges, numEdges, sliceValue, planeNormal);

	return cudaGetLastError();
}


cudaError_t computeIntersectionWithPlane(int numTriangles, float3 *triangles, float3 planePosition, float3 planeNormal, float3 *edges, int *numEdges)
{
    
	dim3 threads( 128 , 1 );

	int wB = ( numTriangles + threads.x - 1 ) / threads.x;
	int hB = 1;

	dim3 blocks(wB, hB);

	int sharedMemorySize =  threads.x * 2 * sizeof(float3) + 2 * sizeof( int );

	computeIntersectionWithPlane_Kernel << < blocks, threads , sharedMemorySize >> >(numTriangles, triangles, edges, numEdges, planePosition, planeNormal);

	return cudaGetLastError();
}
	.file	"tmpxft_003611b2_00000000-6_planeslicer.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2013:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2013:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z3dot6float3S_
	.type	_Z3dot6float3S_, @function
_Z3dot6float3S_:
.LFB2008:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$56, %rsp
	.cfi_def_cfa_offset 64
	movl	$1, 44(%rsp)
	movl	44(%rsp), %edi
	movq	%xmm0, 16(%rsp)
	movss	%xmm1, 24(%rsp)
	movq	%xmm2, (%rsp)
	movss	%xmm3, 8(%rsp)
	call	exit@PLT
	.cfi_endproc
.LFE2008:
	.size	_Z3dot6float3S_, .-_Z3dot6float3S_
	.globl	_Z36computeIntersectionWithStandardPlaneiP6float3S0_S0_S0_PiS_i
	.type	_Z36computeIntersectionWithStandardPlaneiP6float3S0_S0_S0_PiS_i, @function
_Z36computeIntersectionWithStandardPlaneiP6float3S0_S0_S0_PiS_i:
.LFB2009:
	.cfi_startproc
	endbr64
	jmp	cudaGetLastError@PLT
	.cfi_endproc
.LFE2009:
	.size	_Z36computeIntersectionWithStandardPlaneiP6float3S0_S0_S0_PiS_i, .-_Z36computeIntersectionWithStandardPlaneiP6float3S0_S0_S0_PiS_i
	.globl	_Z71__device_stub__Z35computeIntersectionWithPlane_KerneliP6float3S0_PiS_S_iP6float3S0_PiRS_S2_
	.type	_Z71__device_stub__Z35computeIntersectionWithPlane_KerneliP6float3S0_PiS_S_iP6float3S0_PiRS_S2_, @function
_Z71__device_stub__Z35computeIntersectionWithPlane_KerneliP6float3S0_PiS_S_iP6float3S0_PiRS_S2_:
.LFB2035:
	.cfi_startproc
	endbr64
	subq	$168, %rsp
	.cfi_def_cfa_offset 176
	movl	%edi, 28(%rsp)
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movq	%rcx, (%rsp)
	leaq	48(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 152(%rsp)
	xorl	%eax, %eax
	leaq	28(%rsp), %rax
	movq	%r8, 136(%rsp)
	movq	%rax, 104(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	movq	%rsp, %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%r9, 144(%rsp)
	movq	%rax, 56(%rsp)
	movl	$1, 64(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L5
	pushq	48(%rsp)
	.cfi_def_cfa_offset 184
	leaq	_Z35computeIntersectionWithPlane_KerneliP6float3S0_PiS_S_(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 192
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	120(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 184
	popq	%rdx
	.cfi_def_cfa_offset 176
.L5:
	movq	152(%rsp), %rax
	subq	%fs:40, %rax
	je	.L7
	call	__stack_chk_fail@PLT
.L7:
	addq	$168, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2035:
	.size	_Z71__device_stub__Z35computeIntersectionWithPlane_KerneliP6float3S0_PiS_S_iP6float3S0_PiRS_S2_, .-_Z71__device_stub__Z35computeIntersectionWithPlane_KerneliP6float3S0_PiS_S_iP6float3S0_PiRS_S2_
	.globl	_Z35computeIntersectionWithPlane_KerneliP6float3S0_PiS_S_
	.type	_Z35computeIntersectionWithPlane_KerneliP6float3S0_PiS_S_, @function
_Z35computeIntersectionWithPlane_KerneliP6float3S0_PiS_S_:
.LFB2036:
	.cfi_startproc
	endbr64
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movq	%rsp, %r9
	leaq	16(%rsp), %r8
	movq	%xmm0, 16(%rsp)
	movss	%xmm1, 24(%rsp)
	movq	%xmm2, (%rsp)
	movss	%xmm3, 8(%rsp)
	call	_Z71__device_stub__Z35computeIntersectionWithPlane_KerneliP6float3S0_PiS_S_iP6float3S0_PiRS_S2_
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2036:
	.size	_Z35computeIntersectionWithPlane_KerneliP6float3S0_PiS_S_, .-_Z35computeIntersectionWithPlane_KerneliP6float3S0_PiS_S_
	.globl	_Z28computeIntersectionWithPlaneiP6float3S_S_S0_Pi
	.type	_Z28computeIntersectionWithPlaneiP6float3S_S_S0_Pi, @function
_Z28computeIntersectionWithPlaneiP6float3S_S_S0_Pi:
.LFB2010:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	xorl	%r9d, %r9d
	movq	%rcx, %r13
	movl	$3080, %r8d
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	movq	%rdx, %r12
	movl	$33554433, %edx
	movl	$1, %ecx
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	movq	%rsi, %rbp
	salq	$7, %rdx
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	movl	%edi, %ebx
	subq	$104, %rsp
	.cfi_def_cfa_offset 144
	movq	%xmm0, 16(%rsp)
	movss	%xmm1, 24(%rsp)
	movq	%xmm2, (%rsp)
	movss	%xmm3, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 88(%rsp)
	xorl	%eax, %eax
	leal	127(%rdi), %eax
	shrl	$7, %eax
	movl	%eax, 52(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	52(%rsp), %rdi
	movl	60(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L13
	movq	16(%rsp), %rax
	leaq	76(%rsp), %r9
	movq	%r13, %rcx
	movq	%r12, %rdx
	leaq	64(%rsp), %r8
	movq	%rbp, %rsi
	movl	%ebx, %edi
	movq	%rax, 64(%rsp)
	movl	24(%rsp), %eax
	movl	%eax, 72(%rsp)
	movq	(%rsp), %rax
	movq	%rax, 76(%rsp)
	movl	8(%rsp), %eax
	movl	%eax, 84(%rsp)
	call	_Z71__device_stub__Z35computeIntersectionWithPlane_KerneliP6float3S0_PiS_S_iP6float3S0_PiRS_S2_
.L13:
	movq	88(%rsp), %rax
	subq	%fs:40, %rax
	je	.L14
	call	__stack_chk_fail@PLT
.L14:
	addq	$104, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	jmp	cudaGetLastError@PLT
	.cfi_endproc
.LFE2010:
	.size	_Z28computeIntersectionWithPlaneiP6float3S_S_S0_Pi, .-_Z28computeIntersectionWithPlaneiP6float3S_S_S0_Pi
	.globl	_Z87__device_stub__Z43computeIntersectionWithStandardPlane_KerneliP6float3S0_S0_S0_S0_PiS_iiP6float3S0_S0_S0_S0_PiRS_i
	.type	_Z87__device_stub__Z43computeIntersectionWithStandardPlane_KerneliP6float3S0_S0_S0_S0_PiS_iiP6float3S0_S0_S0_S0_PiRS_i, @function
_Z87__device_stub__Z43computeIntersectionWithStandardPlane_KerneliP6float3S0_S0_S0_S0_PiS_iiP6float3S0_S0_S0_S0_PiRS_i:
.LFB2037:
	.cfi_startproc
	endbr64
	subq	$216, %rsp
	.cfi_def_cfa_offset 224
	movq	224(%rsp), %rax
	movl	%edi, 60(%rsp)
	leaq	80(%rsp), %rdi
	movq	%rsi, 48(%rsp)
	leaq	92(%rsp), %rsi
	movq	%rdx, 40(%rsp)
	leaq	64(%rsp), %rdx
	movq	%rcx, 32(%rsp)
	leaq	72(%rsp), %rcx
	movq	%r8, 24(%rsp)
	movq	%r9, 16(%rsp)
	movq	%rax, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 200(%rsp)
	xorl	%eax, %eax
	leaq	60(%rsp), %rax
	movl	$1, 88(%rsp)
	movq	%rax, 128(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 176(%rsp)
	movq	232(%rsp), %rax
	movl	$1, 100(%rsp)
	movq	%rax, 184(%rsp)
	leaq	240(%rsp), %rax
	movq	%rax, 192(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 80(%rsp)
	movq	%rax, 92(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L17
	pushq	72(%rsp)
	.cfi_def_cfa_offset 232
	leaq	_Z43computeIntersectionWithStandardPlane_KerneliP6float3S0_S0_S0_S0_PiS_i(%rip), %rdi
	pushq	72(%rsp)
	.cfi_def_cfa_offset 240
	movq	108(%rsp), %rcx
	movl	116(%rsp), %r8d
	movq	96(%rsp), %rsi
	movl	104(%rsp), %edx
	leaq	144(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 232
	popq	%rdx
	.cfi_def_cfa_offset 224
.L17:
	movq	200(%rsp), %rax
	subq	%fs:40, %rax
	je	.L19
	call	__stack_chk_fail@PLT
.L19:
	addq	$216, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2037:
	.size	_Z87__device_stub__Z43computeIntersectionWithStandardPlane_KerneliP6float3S0_S0_S0_S0_PiS_iiP6float3S0_S0_S0_S0_PiRS_i, .-_Z87__device_stub__Z43computeIntersectionWithStandardPlane_KerneliP6float3S0_S0_S0_S0_PiS_iiP6float3S0_S0_S0_S0_PiRS_i
	.globl	_Z43computeIntersectionWithStandardPlane_KerneliP6float3S0_S0_S0_S0_PiS_i
	.type	_Z43computeIntersectionWithStandardPlane_KerneliP6float3S0_S0_S0_S0_PiS_i, @function
_Z43computeIntersectionWithStandardPlane_KerneliP6float3S0_S0_S0_S0_PiS_i:
.LFB2038:
	.cfi_startproc
	endbr64
	subq	$32, %rsp
	.cfi_def_cfa_offset 40
	movl	48(%rsp), %eax
	movq	%xmm0, 8(%rsp)
	movss	%xmm1, 16(%rsp)
	pushq	%rax
	.cfi_def_cfa_offset 48
	leaq	16(%rsp), %rax
	pushq	%rax
	.cfi_def_cfa_offset 56
	pushq	56(%rsp)
	.cfi_def_cfa_offset 64
	call	_Z87__device_stub__Z43computeIntersectionWithStandardPlane_KerneliP6float3S0_S0_S0_S0_PiS_iiP6float3S0_S0_S0_S0_PiRS_i
	addq	$56, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2038:
	.size	_Z43computeIntersectionWithStandardPlane_KerneliP6float3S0_S0_S0_S0_PiS_i, .-_Z43computeIntersectionWithStandardPlane_KerneliP6float3S0_S0_S0_S0_PiS_i
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"_Z43computeIntersectionWithStandardPlane_KerneliP6float3S0_S0_S0_S0_PiS_i"
.LC1:
	.string	"_Z35computeIntersectionWithPlane_KerneliP6float3S0_PiS_S_"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2040:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC0(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z43computeIntersectionWithStandardPlane_KerneliP6float3S0_S0_S0_S0_PiS_i(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC1(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z35computeIntersectionWithPlane_KerneliP6float3S0_PiS_S_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2040:
	.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:

#ifndef __PLANESLICER_CUH__
#define __PLANESLICER_CUH__

#include "hip/hip_runtime.h"
#include "hip/hip_runtime_api.h"



hipError_t computeIntersectionWithPlane( int numTriangles, float3 *triangles, float3 planePosition, float3 planeNormal, float3 *edges, int *numEdges );



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