IterateUntilCondition Namespace Reference

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

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

Function Documentation

check_if_participating()

__device__ bool IterateUntilCondition::check_if_participating ( Storage * store )

inline

Definition at line 118 of file IterateUntilCondition.h.

  {
    const bool locked = try_lock_mutex(store);
 
    unsigned int old_count = Storage::NumMaxBlocks;
 
    volatile unsigned int * poll_closed_ptr = static_cast<volatile unsigned int *>(&store->poll_closed);
 
    if (*poll_closed_ptr == 0)
      {
        old_count = atomicAdd(&store->count, 1);
        store->block_indices[blockIdx.x] = old_count;
        unlock_mutex(store);
      }
    else
      {
        if (locked)
          {
            unlock_mutex(store);
          }
        return false;
      }
 
    if (*poll_closed_ptr == 0)
      {
        try_lock_mutex(store);
        *poll_closed_ptr = 1;
        disable_mutex(store);
        unlock_mutex(store);
      }
 
    return (old_count < Storage::NumMaxBlocks);
  }

cooperative_kernel()

template<class HolderLike, class ... Args>

__global__ void IterateUntilCondition::cooperative_kernel

(

Args ...

args

)

Definition at line 65 of file IterateUntilCondition.h.

  {
    cooperative_kernel_impl(HolderLike{}, args...);
  }

cooperative_kernel_impl()

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 )

Definition at line 43 of file IterateUntilCondition.h.

  {
    cooperative_groups::grid_group grid = cooperative_groups::this_grid();
 
    Condition checker;
 
    Before{}(gridDim.x, blockIdx.x, checker, std::forward<Args>(args)...);
 
    while (!checker(gridDim.x, blockIdx.x, std::forward<Args>(args)...))
      {
        auto helper = [&](auto func)
        {
          func(gridDim.x, blockIdx.x, checker, std::forward<Args>(args)...);
          grid.sync();
        };
 
        (helper(Funcs{}), ...);
      }
    After{}(gridDim.x, blockIdx.x, checker, std::forward<Args>(args)...);
  }

disable_mutex()

__device__ void IterateUntilCondition::disable_mutex ( Storage * store )

inline

Definition at line 112 of file IterateUntilCondition.h.

  {
    volatile unsigned int * ptr = static_cast<volatile unsigned int *>(&store->mutex_check);
    *ptr = *ptr | 0x80000000U;
  }

make_holder()

template<class Condition, class Before, class After, class ... Funcs>

auto IterateUntilCondition::make_holder	(	Condition	c,
		Before	b,
		After	a,
		Funcs ...	fs )

Must pass functors!

They will receive two unsigned ints for grid size and block index (for simplicity, we only handle 1D grids), a reference to a mutable Condition (except for Condition) and any arguments you pass to the execute of this return. The functions should not use the actual block indices, but the thread indices inside the block are respected. Condition must return a boolean, with true meaning we have reached the end of the iterations, all others are void. Condition may be (locally) stateful as the same local instance is used throughout the iterations, while the others must be stateless (being constructed every iteration in the case of Funcs).

Definition at line 335 of file IterateUntilCondition.h.

  {
    return Holder<Condition, Before, After, Funcs...> {};
  }

normal_kernel()

template<class HolderLike, class ... Args>

__global__ void IterateUntilCondition::normal_kernel	(	Storage *	store,
		Args ...	args )

Definition at line 264 of file IterateUntilCondition.h.

  {
    normal_kernel_impl(HolderLike{}, store, args...);
  }

normal_kernel_impl()

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 )

Definition at line 153 of file IterateUntilCondition.h.

  {
    __shared__ bool is_participating;
 
    const bool is_reference_thread = (threadIdx.x == 0 && threadIdx.y == 0 && threadIdx.z == 0);
 
    if (is_reference_thread)
      {
        is_participating = check_if_participating(store);
      }
 
    __syncthreads();
 
    const unsigned int this_block_index = store->block_indices[blockIdx.x];
    const unsigned int total_blocks = min(*static_cast<volatile unsigned int *>(&store->count), Storage::NumMaxBlocks);
        
    if (is_participating)
      {
        const bool is_reference_block = (this_block_index == 0);
 
        const unsigned int this_thread_index = threadIdx.z * blockDim.y * blockDim.x +
                                               threadIdx.y * blockDim.x +
                                               threadIdx.x;
 
        const unsigned int num_threads_per_block = blockDim.x * blockDim.y * blockDim.z;
 
        Condition checker;
 
        
        Before{}(total_blocks, this_block_index, checker, std::forward<Args>(args)...);
 
        while (!checker(total_blocks, this_block_index, std::forward<Args>(args)...))
          {
            auto helper = [&](auto func)
            {
 
              func(total_blocks, this_block_index, checker, std::forward<Args>(args)...);
 
              //Technically, for the foreseeable future,
              //this could be simply the if,
              //as the maximum number of concurrent blocks
              //in all devices is smaller than 1024...
              if (is_reference_block)
                {
                  for (unsigned int block_to_check = this_thread_index + 1; block_to_check < total_blocks; block_to_check += num_threads_per_block)
                    {
                      unsigned int last_check = 0;
                      
                      volatile unsigned int * to_check = static_cast<volatile unsigned int*>(&store->wait_flags[block_to_check]);
 
                      do
                        {
                          last_check = *to_check;
                        }
                      while (last_check == 0);
                      //When porting to non-CUDA, this may need to be some form of atomic load.
                    }
 
                  __syncthreads();
 
                  for (unsigned int block_to_check = this_thread_index + 1; block_to_check < total_blocks; block_to_check += num_threads_per_block)
                    {
                      __threadfence();
                      //*static_cast<volatile unsigned int*>(&store->wait_flags[block_to_check]) = 0;
                      atomicAnd(&(store->wait_flags[block_to_check]), 0U);
                    }
 
                  __syncthreads();
                }
              else
                {
                  __syncthreads();
 
                  if (is_reference_thread)
                    {                      
                      atomicOr(&(store->wait_flags[this_block_index]), 1U);
 
                      unsigned int last_check = 1;
                      
                      volatile unsigned int * to_check = static_cast<volatile unsigned int*>(&store->wait_flags[this_block_index]);
 
                      __threadfence();
                      
                      do
                        {
                          last_check = *to_check;
                        }
                      while (last_check != 0);
                      //When porting to non-CUDA, this may need to be some form of atomic load.
                    }
 
                  __syncthreads();
                }
 
            };
 
            (helper(Funcs{}), ...);
          }
 
        After{}(total_blocks, this_block_index, checker, std::forward<Args>(args)...);
      }
    
#if CALORECGPU_ITERATE_UNTIL_CONDITION_DEBUG
    if (is_reference_thread)
      {
        printf("%d | %d | %u %u \n", blockIdx.x, static_cast<int>(is_participating), total_blocks, this_block_index);
      }
#endif
  }

try_lock_mutex()

__device__ bool IterateUntilCondition::try_lock_mutex ( Storage * store )

inline

Definition at line 86 of file IterateUntilCondition.h.

  {
    const unsigned int ticket = atomicAdd(&store->mutex_ticket, 1U);
 
    unsigned int last_check = 0;
 
    bool was_once_valid = false;
 
    volatile unsigned int * to_check = static_cast<volatile unsigned int *>(&store->mutex_check);
    
    do
      {
        last_check = *to_check;
        was_once_valid = !(last_check & 0x80000000U);
      }
    while (last_check < ticket && !(last_check & 0x80000000U));
 
    return was_once_valid;
  }

unlock_mutex()

__device__ void IterateUntilCondition::unlock_mutex ( Storage * store )

inline

Definition at line 106 of file IterateUntilCondition.h.

  {
    volatile unsigned int * ptr = static_cast<volatile unsigned int *>(&store->mutex_check);
    *ptr = *ptr + 1;
  }