#include "base/cs_defs.h"
#include <float.h>
#include "hip/hip_runtime.h"
#include <hip/hip_runtime_api.h>
#include "base/cs_base_hip.h"

Include dependency graph for cs_math_hip.h:

Classes
class	assembled_value< T,... >
	Class for value assembly using warp-based partial summation to reduce conflict occurences in downstream atomic addition. More...

class	assembled_value< T, Head, Tail... >

Functions
template<class V >
__device__ cs_hip_mask_t	_conflict_mask (cs_hip_mask_t mask, V v) noexcept
	Return the bit mask of the threads that have the same value v. More...

template<class T >
__device__ bool	_reduce_add (cs_hip_mask_t mask, cs_hip_mask_t peers, T &v) noexcept
	Reduction step for conflict-free addition. More...

template<typename T >
static __device__ void	cs_math_3_normalize_hip (const T in[3], T out[3])
	Normalise a vector of 3 real values. More...

template<typename T >
static __device__ T	cs_math_3_square_norm_hip (const T in[3])
	Compute the square norm of a vector of 3 real values. More...

Function Documentation

◆ _conflict_mask()

__device__ cs_hip_mask_t _conflict_mask	(	cs_hip_mask_t	mask,
		V	v
	)

noexcept

Return the bit mask of the threads that have the same value v.

This partitionins the incoming threads into groups whose members have the same v value.

This reduced to calling __match_any_sync on recent HIP (>=6.2) architectures.

Parameters

[in]	mask
[in]	v

noexcept

Reduction step for conflict-free addition.

Parameters

[in]	mask	mask for threads in warpp
[in,out]	v	value which should be reduced between threads

static __device__ void cs_math_3_normalize_hip	(	const T	in[3],
		T	out[3]
	)

static

Normalise a vector of 3 real values.

To normalize in-place, vin and vout may point to the same array.

Parameters

[in]	vin	vector
[out]	vout	normalized vector

static __device__ T cs_math_3_square_norm_hip ( const T in[3] )

static

Compute the square norm of a vector of 3 real values.

Parameters

[in] v vector of 3 real values