Description

Collaboration diagram for CUDA functions:

Functions
__device__ __host__ int64_t3	convertPosLocalToGlobal (unsigned int ownerSD, const int3 &local_pos, ChSystemGpu_impl::GranParamsPtr gran_params)
	Convert position from its owner subdomain local frame to the global big domain frame.

__device__ void	figureOutTouchedSD (int sphCenter_X_local, int sphCenter_Y_local, int sphCenter_Z_local, int3 ownerSD, unsigned int SDs[MAX_SDs_TOUCHED_BY_SPHERE], ChSystemGpu_impl::GranParamsPtr gran_params)
	Takes in a sphere's position and inserts into the given int array[8] which subdomains, if any, are touched The array is indexed with the ones bit equal to +/- x, twos bit equal to +/- y, and the fours bit equal to +/- z A bit set to 0 means the lower index, whereas 1 means the higher index (lower + 1) The kernel computes global x, y, and z indices for the bottom-left subdomain and then uses those to figure out which subdomains described in the corresponding 8-SD cube are touched by the sphere. More...

template<typename T >
__global__ void	elementalArray3Squared (T sqSum, const T arrX, const T arrY, const T arrZ, size_t nSpheres)
	Compute the elementwise squared sum of array XYZ components.

template<unsigned int CUB_THREADS>
__global__ void	getNumberOfSpheresTouchingEachSD (ChSystemGpu_impl::GranSphereDataPtr sphere_data, unsigned int nSpheres, ChSystemGpu_impl::GranParamsPtr gran_params)
	Template arguments: More...

__device__ int3	getOffsetFromSDs (unsigned int thisSD, unsigned int otherSD, ChSystemGpu_impl::GranParamsPtr gran_params)
	Get position offset between two SDs NOTE this assumes they are close together.

__device__ void	findNewLocalCoords (ChSystemGpu_impl::GranSphereDataPtr sphere_data, unsigned int mySphereID, int64_t global_pos_X, int64_t global_pos_Y, int64_t global_pos_Z, ChSystemGpu_impl::GranParamsPtr gran_params)
	update local positions and SD based on global position

__device__ void	applyGravity (float3 &sphere_force, ChSystemGpu_impl::GranParamsPtr gran_params)
	Apply gravity to a sphere.

__device__ void	applyExternalForces_frictionless (unsigned int ownerSD, const int3 &sphPos_local, const float3 &sphVel, float3 &sphere_force, ChSystemGpu_impl::GranParamsPtr gran_params, ChSystemGpu_impl::GranSphereDataPtr sphere_data, BC_type bc_type_list, BC_params_t< int64_t, int64_t3 > bc_params_list, unsigned int nBCs)
	Compute forces on a sphere from walls, BCs, and gravity.

__device__ void	applyExternalForces (unsigned int currSphereID, unsigned int ownerSD, const int3 &sphPos_local, const float3 &sphVel, const float3 &sphOmega, float3 &sphere_force, float3 &sphere_ang_acc, ChSystemGpu_impl::GranParamsPtr gran_params, ChSystemGpu_impl::GranSphereDataPtr sphere_data, BC_type bc_type_list, BC_params_t< int64_t, int64_t3 > bc_params_list, unsigned int nBCs)
	Compute forces on a sphere from walls, BCs, and gravity.

__device__ float3	computeSphereNormalForces (float &reciplength, float3 &vrel_t, float3 &delta_r, const int3 &sphereA_pos, const int3 &sphereB_pos, const float3 &sphereA_vel, const float3 &sphereB_vel, ChSystemGpu_impl::GranParamsPtr gran_params)
	Compute normal forces for a contacting pair.

__device__ float3	computeSphereNormalForces_matBased (float3 &vrel_t, float3 &contact_normal, float &sqrt_Rd, float &beta, const int3 &sphereA_pos, const int3 &sphereB_pos, const float3 &sphereA_vel, const float3 &sphereB_vel, ChSystemGpu_impl::GranParamsPtr gran_params)
	Compute normal forces for a contacting pair.

__device__ bool	evaluateRollingFriction (ChSystemGpu_impl::GranParamsPtr gran_params, const float &E_eff, const float &R_eff, const float &beta, const float &m_eff, const float &time_contact, float &t_collision)

__device__ float	integrateForwardEuler (float stepsize_SU, float val_dt)
	Compute update for a quantity using Forward Euler integrator.

__device__ float	integrateChung_vel (float stepsize_SU, float acc, float acc_old)
	Compute update for a velocity using Chung integrator.

__device__ float	integrateChung_pos (float stepsize_SU, float vel_old, float acc, float acc_old)
	Compute update for a position using Chung integrator.

template<class IN_T , class IN_T3 , class OUT_T3 = IN_T3>
__device__ OUT_T3	apply_frame_transform (const IN_T3 &point, const IN_T pos, const IN_T rot_mat)
	Point is in the LRF, rot_mat rotates LRF to GRF, pos translates LRF to GRF LRF: local reference frame GRF: global reference frame.

template<class T3 >
__device__ void	convert_pos_UU2SU (T3 &pos, ChSystemGpu_impl::GranParamsPtr gran_params)
	Convert position vector from user units to scaled units.

__inline__ __device__ void	triangle_figureOutSDBox (const float3 &vA, const float3 &vB, const float3 &vC, int L, int U, ChSystemGpu_impl::GranParamsPtr gran_params)
	Takes in a triangle ID and figures out an SD AABB for broadphase use.

__device__ unsigned int	triangle_countTouchedSDs (unsigned int triangleID, const ChSystemGpuMesh_impl::TriangleSoupPtr triangleSoup, ChSystemGpu_impl::GranParamsPtr gran_params, ChSystemGpuMesh_impl::MeshParamsPtr tri_params)
	Takes in a triangle's position in UU and finds out how many SDs it touches. More...

__device__ void	triangle_figureOutTouchedSDs (unsigned int triangleID, const ChSystemGpuMesh_impl::TriangleSoupPtr triangleSoup, unsigned int *touchedSDs, const ChSystemGpu_impl::GranParamsPtr &gran_params, const ChSystemGpuMesh_impl::MeshParamsPtr &tri_params)
	Takes in a triangle's position in UU and finds out what SDs it touches. More...

__global__ void	determineCountOfSDsTouchedByEachTriangle (const ChSystemGpuMesh_impl::TriangleSoupPtr d_triangleSoup, unsigned int *Triangle_NumSDsTouching, ChSystemGpu_impl::GranParamsPtr gran_params, ChSystemGpuMesh_impl::MeshParamsPtr mesh_params)

__global__ void	storeSDsTouchedByEachTriangle (const ChSystemGpuMesh_impl::TriangleSoupPtr d_triangleSoup, const unsigned int Triangle_NumSDsTouching, const unsigned int TriangleSDCompositeOffsets, unsigned int Triangle_SDsComposite, unsigned int Triangle_TriIDsComposite, ChSystemGpu_impl::GranParamsPtr gran_params, ChSystemGpuMesh_impl::MeshParamsPtr mesh_params)
	kernel is called to populate two arrays: one holds on to what SDs touch a triangle, and then a companion array that keeps repeating the triangle ID; in the end, these two arrays will be handled by a CUB sort-by-key op More...

__global__ void	finalizeSD_numTrianglesTouching (const unsigned int d_SDs_touched, const unsigned int d_howManyTrianglesTouchTheTouchedSDs, const unsigned int nSDs_touchedByTriangles, unsigned int pSD_numTrianglesTouching)
	Upon input, pSD_numTrianglesTouching contains a bunch of zeros. More...

__device__ bool	snap_to_face (const double3 &A, const double3 &B, const double3 &C, const double3 &P, double3 &res)
	This utility function takes the location 'P' and snaps it to the closest point on the triangular face with given vertices (A, B, and C). More...

__device__ bool	face_sphere_cd (const double3 &A, const double3 &B, const double3 &C, const double3 &sphere_pos, const int radius, float3 &normal, float &depth, double3 &pt1)
	/brief TRIANGLE FACE - SPHERE NARROW-PHASE COLLISION DETECTION More...

Variables
const int	SAFETY_PARAM = 1000

Function Documentation

◆ determineCountOfSDsTouchedByEachTriangle()

__global__ void determineCountOfSDsTouchedByEachTriangle	(	const ChSystemGpuMesh_impl::TriangleSoupPtr	d_triangleSoup,
		unsigned int *	Triangle_NumSDsTouching,
		ChSystemGpu_impl::GranParamsPtr	gran_params,
		ChSystemGpuMesh_impl::MeshParamsPtr	mesh_params
	)

Parameters

Triangle_NumSDsTouching number of SDs touching this Triangle

◆ face_sphere_cd()

__device__ bool face_sphere_cd	(	const double3 &	A,
		const double3 &	B,
		const double3 &	C,
		const double3 &	sphere_pos,
		const int	radius,
		float3 &	normal,
		float &	depth,
		double3 &	pt1
	)

/brief TRIANGLE FACE - SPHERE NARROW-PHASE COLLISION DETECTION

The triangular face is defined by points A, B, C. The sequence is important as it defines the positive face via a right-hand rule. The sphere is centered at sphere_pos and has radius. The index "1" is associated with the triangle. The index "2" is associated with the sphere. The coordinates of the face and sphere are assumed to be provided in the same reference frame.

Output:

pt1: contact point on triangle
depth: penetration distance (a negative value means that overlap exists)
normal: contact normal, from pt2 to pt1 A return value of "true" signals collision.

Parameters

A	First vertex of the triangle
B	Second vertex of the triangle
C	Third vertex of the triangle
sphere_pos	Location of the center of the sphere
radius	Sphere radius
normal	contact normal
depth	penetration
pt1	contact point on triangle

◆ figureOutTouchedSD()

__device__ void figureOutTouchedSD	(	int	sphCenter_X_local,
		int	sphCenter_Y_local,
		int	sphCenter_Z_local,
		int3	ownerSD,
		unsigned int	SDs[MAX_SDs_TOUCHED_BY_SPHERE],
		ChSystemGpu_impl::GranParamsPtr	gran_params
	)

inline

Takes in a sphere's position and inserts into the given int array[8] which subdomains, if any, are touched The array is indexed with the ones bit equal to +/- x, twos bit equal to +/- y, and the fours bit equal to +/- z A bit set to 0 means the lower index, whereas 1 means the higher index (lower + 1) The kernel computes global x, y, and z indices for the bottom-left subdomain and then uses those to figure out which subdomains described in the corresponding 8-SD cube are touched by the sphere.

The kernel then converts these indices to indices into the global SD list via the (currently local) conv[3] data structure Should be mostly bug-free, especially away from boundaries

◆ finalizeSD_numTrianglesTouching()

__global__ void finalizeSD_numTrianglesTouching	(	const unsigned int *	d_SDs_touched,
		const unsigned int *	d_howManyTrianglesTouchTheTouchedSDs,
		const unsigned int *	nSDs_touchedByTriangles,
		unsigned int *	pSD_numTrianglesTouching
	)

Upon input, pSD_numTrianglesTouching contains a bunch of zeros.

Upon return, the array will reflect the fact that some SDs are touched by one or more triangles

Parameters

d_SDs_touched	- list of SDs that happen to be touched by at least one triangle
d_howManyTrianglesTouchTheTouchedSDs	- if SD is in the list above, it says how many triangles touch it
nSDs_touchedByTriangles	- how many SDs are actually touched by at least one triangle
pSD_numTrianglesTouching	- [in/out] array of size SDs, populated with numbre of triangles touched by each SD

Returns

◆ getNumberOfSpheresTouchingEachSD()

template<unsigned int CUB_THREADS>

__global__ void getNumberOfSpheresTouchingEachSD	(	ChSystemGpu_impl::GranSphereDataPtr	sphere_data,
		unsigned int	nSpheres,
		ChSystemGpu_impl::GranParamsPtr	gran_params
	)

Template arguments:

CUB_THREADS: the number of threads used in this kernel, comes into play when invoking CUB block collectives

Assumptions:

Granular material is made up of monodisperse spheres.
The function below assumes the spheres are in a box
The box has dimensions L x D x H.
The reference frame associated with the box:
- The x-axis is along the length L of the box
- The y-axis is along the width D of the box
- The z-axis is along the height H of the box
A sphere cannot touch more than eight SDs

Basic idea: use domain decomposition on the rectangular box and figure out how many sphere touche each SD. The subdomains are axis-aligned relative to the reference frame associated with the box. The origin of the box is at the center of the box. The orientation of the box is defined relative to a world inertial reference frame.

Nomenclature:

SD: subdomain.
BD: the big-domain, which is the union of all SDs
NULL_CHGPU_ID: the equivalent of a non-sphere SD ID, or a non-sphere ID

◆ snap_to_face()

__device__ bool snap_to_face	(	const double3 &	A,
		const double3 &	B,
		const double3 &	C,
		const double3 &	P,
		double3 &	res
	)

This utility function takes the location 'P' and snaps it to the closest point on the triangular face with given vertices (A, B, and C).

The result is returned in 'res'. Both 'P' and 'res' are assumed to be specified in the same frame as the face vertices. This function returns 'true' if the result is on an edge of this face and 'false' if the result is inside the triangle. Code from Ericson, "real-time collision detection", 2005, pp. 141

◆ storeSDsTouchedByEachTriangle()

__global__ void storeSDsTouchedByEachTriangle	(	const ChSystemGpuMesh_impl::TriangleSoupPtr	d_triangleSoup,
		const unsigned int *	Triangle_NumSDsTouching,
		const unsigned int *	TriangleSDCompositeOffsets,
		unsigned int *	Triangle_SDsComposite,
		unsigned int *	Triangle_TriIDsComposite,
		ChSystemGpu_impl::GranParamsPtr	gran_params,
		ChSystemGpuMesh_impl::MeshParamsPtr	mesh_params
	)

kernel is called to populate two arrays: one holds on to what SDs touch a triangle, and then a companion array that keeps repeating the triangle ID; in the end, these two arrays will be handled by a CUB sort-by-key op

Parameters

d_triangleSoup	- the collection of triangles in this mesh soup
Triangle_NumSDsTouching	- number of SDs touched by each triangle
TriangleSDCompositeOffsets	- offsets in the array that say each SD what triangles touch it
Triangle_SDsComposite	- the list of SDs touched by each triangle; triangle by triangle
Triangle_TriIDsComposite	- array that goes hand in hand with Triangle_SDsComposite; it repeats the triangle ID for a subsequent sort by key op that is performed elsewhere
gran_params
mesh_params

Returns

◆ triangle_countTouchedSDs()

__device__ unsigned int triangle_countTouchedSDs	(	unsigned int	triangleID,
		const ChSystemGpuMesh_impl::TriangleSoupPtr	triangleSoup,
		ChSystemGpu_impl::GranParamsPtr	gran_params,
		ChSystemGpuMesh_impl::MeshParamsPtr	tri_params
	)

inline

Takes in a triangle's position in UU and finds out how many SDs it touches.

Triangle broadphase is done in float by applying the frame transform and then converting the GRF position to SU

◆ triangle_figureOutTouchedSDs()

__device__ void triangle_figureOutTouchedSDs	(	unsigned int	triangleID,
		const ChSystemGpuMesh_impl::TriangleSoupPtr	triangleSoup,
		unsigned int *	touchedSDs,
		const ChSystemGpu_impl::GranParamsPtr &	gran_params,
		const ChSystemGpuMesh_impl::MeshParamsPtr &	tri_params
	)

inline

Takes in a triangle's position in UU and finds out what SDs it touches.

Triangle broadphase is done in float by applying the frame transform and then converting the GRF position to SU

Description

Functions

Variables

Function Documentation

◆ determineCountOfSDsTouchedByEachTriangle()

◆ face_sphere_cd()

◆ figureOutTouchedSD()

◆ finalizeSD_numTrianglesTouching()

◆ getNumberOfSpheresTouchingEachSD()

◆ snap_to_face()

◆ storeSDsTouchedByEachTriangle()

◆ triangle_countTouchedSDs()

◆ triangle_figureOutTouchedSDs()