Namespaces
namespace	Launch

Functions
__global__ void	set (Tree tree, integer count, integer start, integer child, integer sorted, integer index, integer toDeleteLeaf, integer toDeleteNode, real minX, real maxX)
	Kernel call to setter. More...

__global__ void	info (Tree tree, Particles particles, integer n, integer m)
	Info Kernel for tree class (for debugging purposes) More...

__global__ void	info (Tree tree, Particles particles)

__global__ void	testTree (Tree tree, Particles particles, integer n, integer m)

__global__ void	set (Tree tree, integer count, integer start, integer child, integer sorted, integer index, integer toDeleteLeaf, integer toDeleteNode, real minX, real maxX, real minY, real maxY)
	Kernel call to setter. More...

__global__ void	set (Tree tree, integer count, integer start, integer child, integer sorted, integer index, integer toDeleteLeaf, integer toDeleteNode, real minX, real maxX, real minY, real maxY, real minZ, real maxZ)

__global__ void	sumParticles (Tree *tree)
	Kernel call to sum particles within tree. More...

__global__ void	buildTree (Tree tree, Particles particles, integer n, integer m)
	Kernel to construct the tree using the particles within `particles` More...

__global__ void	prepareSorting (Tree tree, Particles particles, integer n, integer m)

__global__ void	calculateCentersOfMass (Tree tree, Particles particles, integer n, integer level)
	Compute center of masses (level wise). More...

__global__ void	computeBoundingBox (Tree tree, Particles particles, integer *mutex, integer n, integer blockSize)
	Kernel to compute the bounding box/borders of the tree or rather the particles within the tree. More...

__global__ void	centerOfMass (Tree tree, Particles particles, integer n)

__global__ void	sort (Tree *tree, integer n, integer m)
	Kernel to sort tree/child indices to optimize cache efficiency. More...

__global__ void	getParticleKeys (Tree tree, Particles particles, keyType *keys, integer maxLevel, integer n, Curve::Type curveType=Curve::lebesgue)
	Kernel to get all particle's keys. More...

__global__ void	globalCOM (Tree tree, Particles particles, real com[DIM])
	Compute center of mass for all particles. More...

Function Documentation

◆ buildTree()

__global__ void TreeNS::Kernel::buildTree	(	Tree *	tree,
		Particles *	particles,
		integer	n,
		integer	m
	)

Kernel to construct the tree using the particles within particles

‍Corresponding wrapper function: TreeNS::Kernel::Launch::buildTree()

The algorithm shows an iterative tree construction algorithm utilizing lightweight locks for leaf nodes to avoid race conditions. Particles are assigned to threads, those threads insert its body by traversing the tree from root to the correct node trying to lock the corresponding child pointer via the array index "-2" using atomic operations. If the locking is successful the particle is inserted into the tree, releasing the lock and executing a memory fence to ensure visibility for the other threads. If a particle is already stored at the desired node, a new cell is automatically generated or rather requested and the old and new body are inserted correspondingly. However, if the locking is not successful, the thread need to try again until accomplishing the task. This approach is very similar to the one presented in An Efficient CUDA Implementation of the Tree-Based Barnes Hut n-Body Algorithm.

Parameters

tree	Tree class target instance
particles	Particles class instance/particles to be inserted in tree
n	number of particles
m	number of potential particles to be inserted (needed for start index of nodes)

Definition at line 768 of file tree.cu.

Here is the call graph for this function:

◆ calculateCentersOfMass()

__global__ void TreeNS::Kernel::calculateCentersOfMass	(	Tree *	tree,
		Particles *	particles,
		integer	n,
		integer	level
	)

Compute center of masses (level wise).

The local COM computation kernel is called iteratively starting with the deepest tree level and finishing at the root. The pseudo-particles children entries are weighted to derive the COM, their masses are accumulated and finally the position is normalized by dividing through the summed mass.

Parameters

tree	Tree class instance
particles	Particle class instance
n	Number of particles
level	Current (relevant) tree level

Definition at line 1416 of file tree.cu.

Here is the call graph for this function:

◆ centerOfMass()

__global__ void TreeNS::Kernel::centerOfMass	(	Tree *	tree,
		Particles *	particles,
		integer	n
	)

Kernel to compute center of mass for pseudo-particles/nodes within tree

‍Corresponding wrapper function: TreeNS::Kernel::Launch::centerOfMass()

Parameters

tree	Tree class instance
particles	Particles class instance
n

Definition at line 1501 of file tree.cu.

◆ computeBoundingBox()

__global__ void TreeNS::Kernel::computeBoundingBox	(	Tree *	tree,
		Particles *	particles,
		integer *	mutex,
		integer	n,
		integer	blockSize
	)

Kernel to compute the bounding box/borders of the tree or rather the particles within the tree.

‍Corresponding wrapper function: TreeNS::Kernel::Launch::computeBoundingBox()

The algorithm to compute the bounding box is basically a parallel reduction primitive to derive the minimum and maximum for each coordinate axis.

Parameters

tree	Tree class target instance
particles	Particles class instance/particles within the tree
mutex	mutex/lock
n	number of particles
blockSize	device block size

Definition at line 302 of file tree.cu.

Here is the call graph for this function:

◆ getParticleKeys()

__global__ void TreeNS::Kernel::getParticleKeys	(	Tree *	tree,
		Particles *	particles,
		keyType *	keys,
		integer	maxLevel,
		integer	n,
		Curve::Type	curveType = `Curve::lebesgue`
	)

Kernel to get all particle's keys.

‍Corresponding wrapper function: TreeNS::Kernel::Launch::getParticleKeys()

Parameters

[in]	tree	Tree class instance
[in]	particles	Particles class instance
[out]	keys	input particle's keys
[in]	maxLevel	Tree maximum level
[in]	n
[in]	curveType	SFC curve type (Lebesgue/Hilbert)

Definition at line 1618 of file tree.cu.

◆ globalCOM()

__global__ void TreeNS::Kernel::globalCOM	(	Tree *	tree,
		Particles *	particles,
		real	com[DIM]
	)

Compute center of mass for all particles.

‍Corresponding wrapper function: TreeNS::Kernel::Launch::globalCOM()

Parameters

[in]	tree	Tree class instance
[in]	particles	Particle class instance
[out]	com	Center of mass

Definition at line 1645 of file tree.cu.

Here is the call graph for this function:

◆ info() [1/2]

__global__ void TreeNS::Kernel::info	(	Tree *	tree,
		Particles *	particles
	)

Definition at line 1788 of file tree.cu.

◆ info() [2/2]

__global__ void TreeNS::Kernel::info	(	Tree *	tree,
		Particles *	particles,
		integer	n,
		integer	m
	)

Info Kernel for tree class (for debugging purposes)

‍Corresponding wrapper function: TreeNS::Kernel::Launch::info()

Parameters

tree
n
m

Definition at line 1684 of file tree.cu.

◆ prepareSorting()

__global__ void TreeNS::Kernel::prepareSorting	(	Tree *	tree,
		Particles *	particles,
		integer	n,
		integer	m
	)

Definition at line 1404 of file tree.cu.

◆ set() [1/3]

__global__ void TreeNS::Kernel::set	(	Tree *	tree,
		integer *	count,
		integer *	start,
		integer *	child,
		integer *	sorted,
		integer *	index,
		integer *	toDeleteLeaf,
		integer *	toDeleteNode,
		real *	minX,
		real *	maxX
	)

Kernel call to setter.

‍Corresponding wrapper function: TreeNS::Kernel::Launch::set()

Parameters

tree	Tree class instance (to be constructed)
count
start
child
sorted
index
toDeleteLeaf
toDeleteNode
minX
maxX

Definition at line 1679 of file tree.cu.

◆ set() [2/3]

__global__ void TreeNS::Kernel::set	(	Tree *	tree,
		integer *	count,
		integer *	start,
		integer *	child,
		integer *	sorted,
		integer *	index,
		integer *	toDeleteLeaf,
		integer *	toDeleteNode,
		real *	minX,
		real *	maxX,
		real *	minY,
		real *	maxY
	)

Kernel call to setter.

Parameters

tree
count
start
child
sorted
index
toDeleteLeaf
toDeleteNode
minX
maxX
minY
maxY

Definition at line 1893 of file tree.cu.

◆ set() [3/3]

__global__ void TreeNS::Kernel::set	(	Tree *	tree,
		integer *	count,
		integer *	start,
		integer *	child,
		integer *	sorted,
		integer *	index,
		integer *	toDeleteLeaf,
		integer *	toDeleteNode,
		real *	minX,
		real *	maxX,
		real *	minY,
		real *	maxY,
		real *	minZ,
		real *	maxZ
	)

Kernel call to setter

Parameters

tree
count
start
child
sorted
index
toDeleteLeaf
toDeleteNode
minX
maxX
minY
maxY
minZ
maxZ

Definition at line 1909 of file tree.cu.

◆ sort()

__global__ void TreeNS::Kernel::sort	(	Tree *	tree,
		integer	n,
		integer	m
	)

Kernel to sort tree/child indices to optimize cache efficiency.

‍Corresponding wrapper function: TreeNS::Kernel::Launch::sort()

Parameters

tree	Tree class instance
n
m

Definition at line 1527 of file tree.cu.

◆ sumParticles()

__global__ void TreeNS::Kernel::sumParticles ( Tree * tree )

Kernel call to sum particles within tree.

‍Corresponding wrapper function: TreeNS::Kernel::Launch::sumParticles()

Parameters

tree	target tree instance

Definition at line 440 of file tree.cu.

Here is the call graph for this function:

◆ testTree()

__global__ void TreeNS::Kernel::testTree	(	Tree *	tree,
		Particles *	particles,
		integer	n,
		integer	m
	)

Definition at line 1815 of file tree.cu.

Namespaces

Functions

Function Documentation

◆ buildTree()

◆ calculateCentersOfMass()

◆ centerOfMass()

◆ computeBoundingBox()

◆ getParticleKeys()

◆ globalCOM()

◆ info() [1/2]

◆ info() [2/2]

◆ prepareSorting()

◆ set() [1/3]

◆ set() [2/3]

◆ set() [3/3]

◆ sort()

◆ sumParticles()

◆ testTree()