IterateUntilCondition.h File Reference

This provides functionality to iterate until a condition is reached. More...

#include <cooperative_groups.h>

Include dependency graph for IterateUntilCondition.h:

Go to the source code of this file.

Classes
struct	IterateUntilCondition::BasicStorage
struct	IterateUntilCondition::Storage
struct	IterateUntilCondition::Holder< Condition, Before, After, Funcs >
	Condition, Before, After and Funcs must all be functor classes. More...

Namespaces
namespace	IterateUntilCondition

Macros
#define	CALORECGPU_ITERATE_UNTIL_CONDITION_DEBUG   0
#define	CALORECGPU_ITERATE_UNTIL_CONDITION_INCLUDE_ASSERTS   0

Functions
template<class Condition, class Before, class After, class ... Funcs, class ... Args>
__device__ void	IterateUntilCondition::cooperative_kernel_impl (const Holder< Condition, Before, After, Funcs... > &, Args &&... args)
template<class HolderLike, class ... Args>
__global__ void	IterateUntilCondition::cooperative_kernel (Args ... args)
__device__ bool	IterateUntilCondition::try_lock_mutex (Storage *store)
__device__ void	IterateUntilCondition::unlock_mutex (Storage *store)
__device__ void	IterateUntilCondition::disable_mutex (Storage *store)
__device__ bool	IterateUntilCondition::check_if_participating (Storage *store)
template<class Condition, class Before, class After, class ... Funcs, class ... Args>
__device__ void	IterateUntilCondition::normal_kernel_impl (const Holder< Condition, Before, After, Funcs... > &, Storage *store, Args &&... args)
template<class HolderLike, class ... Args>
__global__ void	IterateUntilCondition::normal_kernel (Storage *store, Args ... args)
template<class Condition, class Before, class After, class ... Funcs>
auto	IterateUntilCondition::make_holder (Condition c, Before b, After a, Funcs ... fs)
	Must pass functors!

Detailed Description

This provides functionality to iterate until a condition is reached.

The "right" way to do this would be using CUDA's cooperative kernel grid-wide synchronisation. However, following from https://doi.org/10.1145/2983990.2984032, one can achieve this with reasonable portability by employing an appropriate discovery and synchronisation strategy.

Definition in file IterateUntilCondition.h.

Macro Definition Documentation

CALORECGPU_ITERATE_UNTIL_CONDITION_DEBUG

#define CALORECGPU_ITERATE_UNTIL_CONDITION_DEBUG   0

Definition at line 25 of file IterateUntilCondition.h.

CALORECGPU_ITERATE_UNTIL_CONDITION_INCLUDE_ASSERTS

#define CALORECGPU_ITERATE_UNTIL_CONDITION_INCLUDE_ASSERTS   0

Definition at line 32 of file IterateUntilCondition.h.