MatrixFreeDevice.hip.cpp File Reference

#include <electrostatics/MatrixFreeDevice.h>
#include <utils/DeviceKernelLauncherHelpers.h>
#include <utils/DeviceAPICalls.h>

Include dependency graph for MatrixFreeDevice.hip.cpp:

Functions
__device__ dftefe::uInt	getMultiVectorIndex (const dftefe::uInt node, const dftefe::uInt batch, const dftefe::uInt nOwnedDofs, const dftefe::uInt nGhostDofs, const dftefe::uInt *__restrict__ ghostMap)
template<typename T , std::uint32_t nDofsPerDim, std::uint32_t batchSize>
__global__ void	constraintsDistributeKernel (T *__restrict__ x, const dftefe::uInt *__restrict__ constrainingNodeBuckets, const dftefe::uInt *__restrict__ constrainingNodeOffset, const dftefe::uInt *__restrict__ constrainedNodeBuckets, const dftefe::uInt *__restrict__ constrainedNodeOffset, const T *__restrict__ weightMatrixList, const dftefe::uInt *__restrict__ weightMatrixOffset, const T *__restrict__ inhomogenityList, const dftefe::uInt *__restrict__ ghostMap, const dftefe::uInt nOwnedDofs, const dftefe::uInt nGhostDofs)
template<typename T , std::uint32_t nDofsPerDim, std::uint32_t batchSize>
__global__ void	constraintsDistributeTransposeKernel (T *__restrict__ Ax, T *__restrict__ x, const dftefe::uInt *__restrict__ constrainingNodeBuckets, const dftefe::uInt *__restrict__ constrainingNodeOffset, const dftefe::uInt *__restrict__ constrainedNodeBuckets, const dftefe::uInt *__restrict__ constrainedNodeOffset, const T *__restrict__ weightMatrixList, const dftefe::uInt *__restrict__ weightMatrixOffset, const dftefe::uInt *__restrict__ ghostMap, const dftefe::uInt nOwnedDofs, const dftefe::uInt nGhostDofs)
template<typename T , std::uint32_t nDofsPerDim, std::uint32_t nQuadPointsPerDim, std::uint32_t batchSize, std::uint32_t dim>
__global__ void	LaplaceKernel (T *__restrict__ dst, const T *__restrict__ src, const T *__restrict__ J, const dftefe::uInt *__restrict__ map)
template<typename T , std::uint32_t nDofsPerDim, std::uint32_t nQuadPointsPerDim, std::uint32_t batchSize, std::uint32_t dim>
__global__ void	HelmholtzKernel (T *__restrict__ dst, const T *__restrict__ src, const T *__restrict__ J, const dftefe::uInt *__restrict__ map, const T coeffHelmholtz)

Variables
constexpr std::uint32_t	maxDofsPerDim = 17
__constant__ double	constMem [(maxDofsPerDim *maxDofsPerDim *5+maxDofsPerDim) *static_cast< std::uint32_t >(operatorList::Count)]

Function Documentation

constraintsDistributeKernel()

template<typename T , std::uint32_t nDofsPerDim, std::uint32_t batchSize>

__global__ void constraintsDistributeKernel	(	T *__restrict__	x,
		const dftefe::uInt *__restrict__	constrainingNodeBuckets,
		const dftefe::uInt *__restrict__	constrainingNodeOffset,
		const dftefe::uInt *__restrict__	constrainedNodeBuckets,
		const dftefe::uInt *__restrict__	constrainedNodeOffset,
		const T *__restrict__	weightMatrixList,
		const dftefe::uInt *__restrict__	weightMatrixOffset,
		const T *__restrict__	inhomogenityList,
		const dftefe::uInt *__restrict__	ghostMap,
		const dftefe::uInt	nOwnedDofs,
		const dftefe::uInt	nGhostDofs
	)

Here is the call graph for this function:

constraintsDistributeTransposeKernel()

template<typename T , std::uint32_t nDofsPerDim, std::uint32_t batchSize>

__global__ void constraintsDistributeTransposeKernel	(	T *__restrict__	Ax,
		T *__restrict__	x,
		const dftefe::uInt *__restrict__	constrainingNodeBuckets,
		const dftefe::uInt *__restrict__	constrainingNodeOffset,
		const dftefe::uInt *__restrict__	constrainedNodeBuckets,
		const dftefe::uInt *__restrict__	constrainedNodeOffset,
		const T *__restrict__	weightMatrixList,
		const dftefe::uInt *__restrict__	weightMatrixOffset,
		const dftefe::uInt *__restrict__	ghostMap,
		const dftefe::uInt	nOwnedDofs,
		const dftefe::uInt	nGhostDofs
	)

Here is the call graph for this function:

getMultiVectorIndex()

__device__ dftefe::uInt getMultiVectorIndex ( const dftefe::uInt  node, const dftefe::uInt  batch, const dftefe::uInt  nOwnedDofs, const dftefe::uInt  nGhostDofs, const dftefe::uInt *__restrict__  ghostMap  )

inline

Here is the caller graph for this function:

HelmholtzKernel()

template<typename T , std::uint32_t nDofsPerDim, std::uint32_t nQuadPointsPerDim, std::uint32_t batchSize, std::uint32_t dim>

__global__ void HelmholtzKernel	(	T *__restrict__	dst,
		const T *__restrict__	src,
		const T *__restrict__	J,
		const dftefe::uInt *__restrict__	map,
		const T	coeffHelmholtz
	)

LaplaceKernel()

template<typename T , std::uint32_t nDofsPerDim, std::uint32_t nQuadPointsPerDim, std::uint32_t batchSize, std::uint32_t dim>

__global__ void LaplaceKernel	(	T *__restrict__	dst,
		const T *__restrict__	src,
		const T *__restrict__	J,
		const dftefe::uInt *__restrict__	map
	)

Variable Documentation

constMem

__constant__ double constMem[(maxDofsPerDim *maxDofsPerDim *5+maxDofsPerDim) * static_cast< std::uint32_t >(operatorList::Count)]

maxDofsPerDim

constexpr std::uint32_t maxDofsPerDim = 17

constexpr