Vlasiator ebf0dd394 on dev (v5.4.0 + 1054 commits)
Loading...
Searching...
No Matches
arch Namespace Reference

Classes

class  buf

Typedefs

typedef uint32_t uint
typedef cub::BlockReduce< T, Blocksize, cub::BLOCK_REDUCE_RAKING_COMMUTATIVE_ONLY, 1, 1 > BlockReduce

Enumerations

enum  reduce_op {
  max , min , sum , prod ,
  null
}

Functions

template<reduce_op Op, uint NDim, typename Lambda, typename T>
static void parallel_reduce (const uint(&limits)[NDim], Lambda loop_body, T &sum)
template<reduce_op Op, uint NDim, uint NReductions, typename Lambda, typename T>
static void parallel_reduce (const uint(&limits)[NDim], Lambda loop_body, T(&sum)[NReductions])
template<reduce_op Op, uint NDim, typename Lambda, typename T>
static void parallel_reduce (const uint(&limits)[NDim], Lambda loop_body, std::vector< T > &sum)
__host__ static __forceinline__ void device_mempool_check (uint64_t threshold_new)
__host__ static __forceinline__ void * allocate (size_t bytes)
__host__ static __forceinline__ void * allocate (size_t bytes, cudaStream_t stream)
template<typename T>
__host__ static __forceinline__ void free (T *ptr)
template<typename T>
__host__ static __forceinline__ void free (T *ptr, cudaStream_t stream)
template<typename T>
static __forceinline__ void memcpy_h2d (T *dst, T *src, size_t bytes)
template<typename T>
static __forceinline__ void memcpy_h2d (T *dst, T *src, size_t bytes, cudaStream_t stream)
template<typename T>
static __forceinline__ void memcpy_d2h (T *dst, T *src, size_t bytes)
template<typename T>
static __forceinline__ void memcpy_d2h (T *dst, T *src, size_t bytes, cudaStream_t stream)
template<typename T>
static __forceinline__ void host_register (T *ptr, size_t bytes)
template<typename T>
static __forceinline__ void host_unregister (T *ptr)
template<typename Lambda, typename T>
__device__ static __forceinline__ void lambda_eval (const uint(&idx)[1], T *__restrict__ thread_data, Lambda loop_body)
template<typename Lambda, typename T>
__device__ static __forceinline__ void lambda_eval (const uint(&idx)[2], T *__restrict__ thread_data, Lambda loop_body)
template<typename Lambda, typename T>
__device__ static __forceinline__ void lambda_eval (const uint(&idx)[3], T *__restrict__ thread_data, Lambda loop_body)
template<typename Lambda, typename T>
__device__ static __forceinline__ void lambda_eval (const uint(&idx)[4], T *__restrict__ thread_data, Lambda loop_body)
template<uint NDim, typename Lambda, typename T>
__device__ static __forceinline__ void loop_eval (const uint idx_glob, const uint *__restrict__ lims, T *__restrict__ thread_data, Lambda loop_body)
template<uint Blocksize, reduce_op Op, uint NDim, uint NReduStatic, typename Lambda, typename T>
static __global__ void __launch_bounds__ (ARCH_BLOCKSIZE_R) reduction_kernel(Lambda loop_body
 if (Op==reduce_op::null)
 for (uint i=0;i< n_reductions;i++)
 if (idx_glob< n_total)
template<reduce_op Op, uint NReduStatic, uint NDim, typename Lambda, typename T>
static __forceinline__ void parallel_reduce_driver (const uint(&limits)[NDim], Lambda loop_body, T *sum, const uint n_redu_dynamic)
__device__ static __forceinline__ void atomicMax (double *address, double val2)
__device__ static __forceinline__ void atomicMin (double *address, double val2)
__device__ static __forceinline__ void atomicMax (float *address, float val2)
__device__ static __forceinline__ void atomicMin (float *address, float val2)
__host__ static __forceinline__ void device_mempool_check (uint64_t threshold_new)
__host__ static __forceinline__ void * allocate (size_t bytes)
__host__ static __forceinline__ void * allocate (size_t bytes, hipStream_t stream)
template<typename T>
__host__ static __forceinline__ void free (T *ptr)
template<typename T>
__host__ static __forceinline__ void free (T *ptr, hipStream_t stream)
template<typename T>
static __forceinline__ void memcpy_h2d (T *dst, T *src, size_t bytes)
template<typename T>
static __forceinline__ void memcpy_h2d (T *dst, T *src, size_t bytes, hipStream_t stream)
template<typename T>
static __forceinline__ void memcpy_d2h (T *dst, T *src, size_t bytes)
template<typename T>
static __forceinline__ void memcpy_d2h (T *dst, T *src, size_t bytes, hipStream_t stream)
template<typename T>
static __forceinline__ void host_register (T *ptr, size_t bytes)
template<typename T>
static __forceinline__ void host_unregister (T *ptr)
template<typename Lambda, typename T>
__device__ static __forceinline__ void lambda_eval (const uint(&idx)[1], T *__restrict__ thread_data, Lambda loop_body)
template<typename Lambda, typename T>
__device__ static __forceinline__ void lambda_eval (const uint(&idx)[2], T *__restrict__ thread_data, Lambda loop_body)
template<typename Lambda, typename T>
__device__ static __forceinline__ void lambda_eval (const uint(&idx)[3], T *__restrict__ thread_data, Lambda loop_body)
template<typename Lambda, typename T>
__device__ static __forceinline__ void lambda_eval (const uint(&idx)[4], T *__restrict__ thread_data, Lambda loop_body)
template<uint NDim, typename Lambda, typename T>
__device__ static __forceinline__ void loop_eval (const uint idx_glob, const uint *__restrict__ lims, T *__restrict__ thread_data, Lambda loop_body)
template<uint Blocksize, reduce_op Op, uint NDim, uint NReduStatic, typename Lambda, typename T>
static __global__ void __launch_bounds__ (ARCH_BLOCKSIZE_R) reduction_kernel(Lambda loop_body
template<reduce_op Op, uint NReduStatic, uint NDim, typename Lambda, typename T>
static __forceinline__ void parallel_reduce_driver (const uint(&limits)[NDim], Lambda loop_body, T *sum, const uint n_redu_dynamic)
static void * allocate (size_t bytes)
static void free (void *ptr)
template<typename T>
static void memcpy_h2d (T *dst, T *src, size_t bytes)
template<typename T>
static void memcpy_d2h (T *dst, T *src, size_t bytes)
template<typename T>
static void host_register (T *ptr, size_t bytes)
template<typename T>
static void host_unregister (T *ptr)
template<reduce_op Op, uint NReductions, uint NDim, typename Lambda, typename T>
static void parallel_reduce_driver (const uint(&limits)[1], Lambda loop_body, T *sum, const uint n_redu_dynamic)
template<reduce_op Op, uint NReductions, uint NDim, typename Lambda, typename T, typename = typename std::enable_if<std::is_void<typename std::invoke_result<Lambda, uint, uint, T*>::type>::value>::type>
static void parallel_reduce_driver (const uint(&limits)[2], Lambda loop_body, T *sum, const uint n_redu_dynamic)
template<reduce_op Op, uint NReductions, uint NDim, typename Lambda, typename T, typename = typename std::enable_if<!std::is_void<typename std::invoke_result<Lambda, uint, uint, T*>::type>::value>::type, typename = void>
static void parallel_reduce_driver (const uint(&limits)[2], Lambda loop_body, T *sum, const uint n_redu_dynamic)
template<reduce_op Op, uint NReductions, uint NDim, typename Lambda, typename T, typename = typename std::enable_if<std::is_void<typename std::invoke_result<Lambda, uint, uint, uint, T*>::type>::value>::type>
static void parallel_reduce_driver (const uint(&limits)[3], Lambda loop_body, T *sum, const uint n_redu_dynamic)
template<reduce_op Op, uint NReductions, uint NDim, typename Lambda, typename T, typename = typename std::enable_if<!std::is_void<typename std::invoke_result<Lambda, uint, uint, uint, T*>::type>::value>::type, typename = void>
static void parallel_reduce_driver (const uint(&limits)[3], Lambda loop_body, T *sum, const uint n_redu_dynamic)
template<reduce_op Op, uint NReductions, uint NDim, typename Lambda, typename T, typename = typename std::enable_if<std::is_void<typename std::invoke_result<Lambda, uint, uint, uint, uint, T*>::type>::value>::type>
static void parallel_reduce_driver (const uint(&limits)[4], Lambda loop_body, T *sum, const uint n_redu_dynamic)
template<reduce_op Op, uint NReductions, uint NDim, typename Lambda, typename T, typename = typename std::enable_if<!std::is_void<typename std::invoke_result<Lambda, uint, uint, uint, uint, T*>::type>::value>::type, typename = void>
static void parallel_reduce_driver (const uint(&limits)[4], Lambda loop_body, T *sum, const uint n_redu_dynamic)

Variables

static __global__ void const T *__restrict__ init_val
static __global__ void const T *__restrict__ T *__restrict__ rslt
static __global__ void const T *__restrict__ T *__restrict__ const uint *__restrict__ lims
static __global__ void const T *__restrict__ T *__restrict__ const uint *__restrict__ const uint n_total
static __global__ void const T *__restrict__ T *__restrict__ const uint *__restrict__ const uint const uint n_redu_dynamic
static __global__ void const T *__restrict__ T *__restrict__ const uint *__restrict__ const uint const uint T * thread_data_dynamic
__shared__ char temp_storage_dynamic []
constexpr uint size = NReduStatic ? NReduStatic : 1
__shared__ BlockReduce::TempStorage temp_storage_static [size]
BlockReduce::TempStorage * temp_storage = NReduStatic ? temp_storage_static : (typename BlockReduce::TempStorage*) temp_storage_dynamic
T thread_data_static [size]
T * thread_data = NReduStatic ? thread_data_static : &thread_data_dynamic[n_redu_dynamic * idx_glob]
const uint n_reductions = NReduStatic ? NReduStatic : n_redu_dynamic

Typedef Documentation

◆ BlockReduce

typedef hipcub::BlockReduce< T, Blocksize, hipcub::BLOCK_REDUCE_RAKING_COMMUTATIVE_ONLY, 1, 1 > arch::BlockReduce

Definition at line 374 of file arch_device_cuda.h.

◆ uint

typedef uint32_t arch::uint

Definition at line 17 of file arch_device_api.h.

Enumeration Type Documentation

◆ reduce_op

enum arch::reduce_op
Enumerator
max 
min 
sum 
prod 
null 

Definition at line 19 of file arch_device_api.h.

Function Documentation

◆ __launch_bounds__() [1/2]

template<uint Blocksize, reduce_op Op, uint NDim, uint NReduStatic, typename Lambda, typename T>
__global__ void arch::__launch_bounds__ ( ARCH_BLOCKSIZE_R )
static

◆ __launch_bounds__() [2/2]

template<uint Blocksize, reduce_op Op, uint NDim, uint NReduStatic, typename Lambda, typename T>
__global__ void arch::__launch_bounds__ ( ARCH_BLOCKSIZE_R )
static
Here is the call graph for this function:

◆ allocate() [1/5]

__host__ static __forceinline__ void * arch::allocate ( size_t bytes)
static

Definition at line 243 of file arch_device_cuda.h.

Here is the call graph for this function:
Here is the caller graph for this function:

◆ allocate() [2/5]

__host__ static __forceinline__ void * arch::allocate ( size_t bytes)
static

Definition at line 235 of file arch_device_hip.h.

Here is the call graph for this function:

◆ allocate() [3/5]

void * arch::allocate ( size_t bytes)
inlinestatic

Definition at line 39 of file arch_device_host.h.

◆ allocate() [4/5]

__host__ static __forceinline__ void * arch::allocate ( size_t bytes,
cudaStream_t stream )
static

Definition at line 255 of file arch_device_cuda.h.

Here is the call graph for this function:

◆ allocate() [5/5]

__host__ static __forceinline__ void * arch::allocate ( size_t bytes,
hipStream_t stream )
static

Definition at line 243 of file arch_device_hip.h.

Here is the call graph for this function:

◆ atomicMax() [1/2]

__device__ static __forceinline__ void arch::atomicMax ( double * address,
double val2 )
static

Definition at line 131 of file arch_device_hip.h.

◆ atomicMax() [2/2]

__device__ static __forceinline__ void arch::atomicMax ( float * address,
float val2 )
static

Definition at line 153 of file arch_device_hip.h.

◆ atomicMin() [1/2]

__device__ static __forceinline__ void arch::atomicMin ( double * address,
double val2 )
static

Definition at line 142 of file arch_device_hip.h.

◆ atomicMin() [2/2]

__device__ static __forceinline__ void arch::atomicMin ( float * address,
float val2 )
static

Definition at line 164 of file arch_device_hip.h.

◆ device_mempool_check() [1/2]

__host__ static __forceinline__ void arch::device_mempool_check ( uint64_t threshold_new)
static

Definition at line 230 of file arch_device_cuda.h.

Here is the caller graph for this function:

◆ device_mempool_check() [2/2]

__host__ static __forceinline__ void arch::device_mempool_check ( uint64_t threshold_new)
static

Definition at line 222 of file arch_device_hip.h.

◆ for()

arch::for ( )

Definition at line 396 of file arch_device_cuda.h.

◆ free() [1/5]

template<typename T>
__host__ static __forceinline__ void arch::free ( T * ptr)
static

Definition at line 264 of file arch_device_cuda.h.

Here is the caller graph for this function:

◆ free() [2/5]

template<typename T>
__host__ static __forceinline__ void arch::free ( T * ptr)
static

Definition at line 252 of file arch_device_hip.h.

◆ free() [3/5]

template<typename T>
__host__ static __forceinline__ void arch::free ( T * ptr,
cudaStream_t stream )
static

Definition at line 274 of file arch_device_cuda.h.

◆ free() [4/5]

template<typename T>
__host__ static __forceinline__ void arch::free ( T * ptr,
hipStream_t stream )
static

Definition at line 258 of file arch_device_hip.h.

◆ free() [5/5]

void arch::free ( void * ptr)
inlinestatic

Definition at line 44 of file arch_device_host.h.

Here is the call graph for this function:

◆ host_register() [1/3]

template<typename T>
__forceinline__ void arch::host_register ( T * ptr,
size_t bytes )
static

Definition at line 311 of file arch_device_cuda.h.

◆ host_register() [2/3]

template<typename T>
__forceinline__ void arch::host_register ( T * ptr,
size_t bytes )
static

Definition at line 287 of file arch_device_hip.h.

◆ host_register() [3/3]

template<typename T>
void arch::host_register ( T * ptr,
size_t bytes )
inlinestatic

Definition at line 58 of file arch_device_host.h.

◆ host_unregister() [1/3]

template<typename T>
__forceinline__ void arch::host_unregister ( T * ptr)
static

Definition at line 317 of file arch_device_cuda.h.

◆ host_unregister() [2/3]

template<typename T>
__forceinline__ void arch::host_unregister ( T * ptr)
static

Definition at line 293 of file arch_device_hip.h.

◆ host_unregister() [3/3]

template<typename T>
void arch::host_unregister ( T * ptr)
inlinestatic

Definition at line 62 of file arch_device_host.h.

◆ if() [1/2]

arch::if ( )

Definition at line 405 of file arch_device_cuda.h.

Here is the call graph for this function:

◆ if() [2/2]

arch::if ( Op = = reduce_op::null)

Definition at line 364 of file arch_device_cuda.h.

Here is the call graph for this function:

◆ lambda_eval() [1/8]

template<typename Lambda, typename T>
__device__ static __forceinline__ void arch::lambda_eval ( const uint(&) idx[1],
T *__restrict__ thread_data,
Lambda loop_body )
static

Definition at line 323 of file arch_device_cuda.h.

Here is the caller graph for this function:

◆ lambda_eval() [2/8]

template<typename Lambda, typename T>
__device__ static __forceinline__ void arch::lambda_eval ( const uint(&) idx[1],
T *__restrict__ thread_data,
Lambda loop_body )
static

Definition at line 299 of file arch_device_hip.h.

◆ lambda_eval() [3/8]

template<typename Lambda, typename T>
__device__ static __forceinline__ void arch::lambda_eval ( const uint(&) idx[2],
T *__restrict__ thread_data,
Lambda loop_body )
static

Definition at line 326 of file arch_device_cuda.h.

◆ lambda_eval() [4/8]

template<typename Lambda, typename T>
__device__ static __forceinline__ void arch::lambda_eval ( const uint(&) idx[2],
T *__restrict__ thread_data,
Lambda loop_body )
static

Definition at line 302 of file arch_device_hip.h.

◆ lambda_eval() [5/8]

template<typename Lambda, typename T>
__device__ static __forceinline__ void arch::lambda_eval ( const uint(&) idx[3],
T *__restrict__ thread_data,
Lambda loop_body )
static

Definition at line 329 of file arch_device_cuda.h.

◆ lambda_eval() [6/8]

template<typename Lambda, typename T>
__device__ static __forceinline__ void arch::lambda_eval ( const uint(&) idx[3],
T *__restrict__ thread_data,
Lambda loop_body )
static

Definition at line 305 of file arch_device_hip.h.

◆ lambda_eval() [7/8]

template<typename Lambda, typename T>
__device__ static __forceinline__ void arch::lambda_eval ( const uint(&) idx[4],
T *__restrict__ thread_data,
Lambda loop_body )
static

Definition at line 332 of file arch_device_cuda.h.

◆ lambda_eval() [8/8]

template<typename Lambda, typename T>
__device__ static __forceinline__ void arch::lambda_eval ( const uint(&) idx[4],
T *__restrict__ thread_data,
Lambda loop_body )
static

Definition at line 308 of file arch_device_hip.h.

◆ loop_eval() [1/2]

template<uint NDim, typename Lambda, typename T>
__device__ static __forceinline__ void arch::loop_eval ( const uint idx_glob,
const uint *__restrict__ lims,
T *__restrict__ thread_data,
Lambda loop_body )
static

Definition at line 336 of file arch_device_cuda.h.

Here is the call graph for this function:
Here is the caller graph for this function:

◆ loop_eval() [2/2]

template<uint NDim, typename Lambda, typename T>
__device__ static __forceinline__ void arch::loop_eval ( const uint idx_glob,
const uint *__restrict__ lims,
T *__restrict__ thread_data,
Lambda loop_body )
static

Definition at line 312 of file arch_device_hip.h.

Here is the call graph for this function:

◆ memcpy_d2h() [1/5]

template<typename T>
__forceinline__ void arch::memcpy_d2h ( T * dst,
T * src,
size_t bytes )
static

Definition at line 295 of file arch_device_cuda.h.

◆ memcpy_d2h() [2/5]

template<typename T>
__forceinline__ void arch::memcpy_d2h ( T * dst,
T * src,
size_t bytes )
static

Definition at line 275 of file arch_device_hip.h.

◆ memcpy_d2h() [3/5]

template<typename T>
void arch::memcpy_d2h ( T * dst,
T * src,
size_t bytes )
inlinestatic

Definition at line 54 of file arch_device_host.h.

◆ memcpy_d2h() [4/5]

template<typename T>
__forceinline__ void arch::memcpy_d2h ( T * dst,
T * src,
size_t bytes,
cudaStream_t stream )
static

Definition at line 305 of file arch_device_cuda.h.

◆ memcpy_d2h() [5/5]

template<typename T>
__forceinline__ void arch::memcpy_d2h ( T * dst,
T * src,
size_t bytes,
hipStream_t stream )
static

Definition at line 281 of file arch_device_hip.h.

◆ memcpy_h2d() [1/5]

template<typename T>
__forceinline__ void arch::memcpy_h2d ( T * dst,
T * src,
size_t bytes )
static

Definition at line 279 of file arch_device_cuda.h.

◆ memcpy_h2d() [2/5]

template<typename T>
__forceinline__ void arch::memcpy_h2d ( T * dst,
T * src,
size_t bytes )
static

Definition at line 263 of file arch_device_hip.h.

◆ memcpy_h2d() [3/5]

template<typename T>
void arch::memcpy_h2d ( T * dst,
T * src,
size_t bytes )
inlinestatic

Definition at line 50 of file arch_device_host.h.

◆ memcpy_h2d() [4/5]

template<typename T>
__forceinline__ void arch::memcpy_h2d ( T * dst,
T * src,
size_t bytes,
cudaStream_t stream )
static

Definition at line 289 of file arch_device_cuda.h.

◆ memcpy_h2d() [5/5]

template<typename T>
__forceinline__ void arch::memcpy_h2d ( T * dst,
T * src,
size_t bytes,
hipStream_t stream )
static

Definition at line 269 of file arch_device_hip.h.

◆ parallel_reduce() [1/3]

template<reduce_op Op, uint NDim, typename Lambda, typename T>
void arch::parallel_reduce ( const uint(&) limits[NDim],
Lambda loop_body,
std::vector< T > & sum )
inlinestatic

Definition at line 53 of file arch_device_api.h.

Here is the call graph for this function:

◆ parallel_reduce() [2/3]

template<reduce_op Op, uint NDim, typename Lambda, typename T>
void arch::parallel_reduce ( const uint(&) limits[NDim],
Lambda loop_body,
T & sum )
inlinestatic

Definition at line 40 of file arch_device_api.h.

Here is the call graph for this function:
Here is the caller graph for this function:

◆ parallel_reduce() [3/3]

template<reduce_op Op, uint NDim, uint NReductions, typename Lambda, typename T>
void arch::parallel_reduce ( const uint(&) limits[NDim],
Lambda loop_body,
T(&) sum[NReductions] )
inlinestatic

Definition at line 47 of file arch_device_api.h.

Here is the call graph for this function:

◆ parallel_reduce_driver() [1/9]

template<reduce_op Op, uint NReductions, uint NDim, typename Lambda, typename T>
void arch::parallel_reduce_driver ( const uint(&) limits[1],
Lambda loop_body,
T * sum,
const uint n_redu_dynamic )
inlinestatic

Definition at line 66 of file arch_device_host.h.

◆ parallel_reduce_driver() [2/9]

template<reduce_op Op, uint NReductions, uint NDim, typename Lambda, typename T, typename = typename std::enable_if<!std::is_void<typename std::invoke_result<Lambda, uint, uint, T*>::type>::value>::type, typename = void>
void arch::parallel_reduce_driver ( const uint(&) limits[2],
Lambda loop_body,
T * sum,
const uint n_redu_dynamic )
inlinestatic

Definition at line 132 of file arch_device_host.h.

◆ parallel_reduce_driver() [3/9]

template<reduce_op Op, uint NReductions, uint NDim, typename Lambda, typename T, typename = typename std::enable_if<std::is_void<typename std::invoke_result<Lambda, uint, uint, T*>::type>::value>::type>
void arch::parallel_reduce_driver ( const uint(&) limits[2],
Lambda loop_body,
T * sum,
const uint n_redu_dynamic )
inlinestatic

Definition at line 95 of file arch_device_host.h.

◆ parallel_reduce_driver() [4/9]

template<reduce_op Op, uint NReductions, uint NDim, typename Lambda, typename T, typename = typename std::enable_if<!std::is_void<typename std::invoke_result<Lambda, uint, uint, uint, T*>::type>::value>::type, typename = void>
void arch::parallel_reduce_driver ( const uint(&) limits[3],
Lambda loop_body,
T * sum,
const uint n_redu_dynamic )
inlinestatic

Definition at line 218 of file arch_device_host.h.

◆ parallel_reduce_driver() [5/9]

template<reduce_op Op, uint NReductions, uint NDim, typename Lambda, typename T, typename = typename std::enable_if<std::is_void<typename std::invoke_result<Lambda, uint, uint, uint, T*>::type>::value>::type>
void arch::parallel_reduce_driver ( const uint(&) limits[3],
Lambda loop_body,
T * sum,
const uint n_redu_dynamic )
inlinestatic

Definition at line 173 of file arch_device_host.h.

◆ parallel_reduce_driver() [6/9]

template<reduce_op Op, uint NReductions, uint NDim, typename Lambda, typename T, typename = typename std::enable_if<!std::is_void<typename std::invoke_result<Lambda, uint, uint, uint, uint, T*>::type>::value>::type, typename = void>
void arch::parallel_reduce_driver ( const uint(&) limits[4],
Lambda loop_body,
T * sum,
const uint n_redu_dynamic )
inlinestatic

Definition at line 320 of file arch_device_host.h.

◆ parallel_reduce_driver() [7/9]

template<reduce_op Op, uint NReductions, uint NDim, typename Lambda, typename T, typename = typename std::enable_if<std::is_void<typename std::invoke_result<Lambda, uint, uint, uint, uint, T*>::type>::value>::type>
void arch::parallel_reduce_driver ( const uint(&) limits[4],
Lambda loop_body,
T * sum,
const uint n_redu_dynamic )
inlinestatic

Definition at line 267 of file arch_device_host.h.

◆ parallel_reduce_driver() [8/9]

template<reduce_op Op, uint NReduStatic, uint NDim, typename Lambda, typename T>
__forceinline__ void arch::parallel_reduce_driver ( const uint(&) limits[NDim],
Lambda loop_body,
T * sum,
const uint n_redu_dynamic )
static

Definition at line 442 of file arch_device_cuda.h.

Here is the call graph for this function:
Here is the caller graph for this function:

◆ parallel_reduce_driver() [9/9]

template<reduce_op Op, uint NReduStatic, uint NDim, typename Lambda, typename T>
__forceinline__ void arch::parallel_reduce_driver ( const uint(&) limits[NDim],
Lambda loop_body,
T * sum,
const uint n_redu_dynamic )
static

Definition at line 419 of file arch_device_hip.h.

Here is the call graph for this function:

Variable Documentation

◆ init_val

__global__ void const T *__restrict__ arch::init_val

Definition at line 359 of file arch_device_cuda.h.

◆ lims

__global__ void const T *__restrict__ T *__restrict__ const uint *__restrict__ arch::lims

Definition at line 359 of file arch_device_cuda.h.

◆ n_redu_dynamic

__global__ void const T *__restrict__ T *__restrict__ const uint *__restrict__ const uint const uint arch::n_redu_dynamic

Definition at line 359 of file arch_device_cuda.h.

◆ n_reductions

const uint arch::n_reductions = NReduStatic ? NReduStatic : n_redu_dynamic

Definition at line 393 of file arch_device_cuda.h.

◆ n_total

__global__ void const T *__restrict__ T *__restrict__ const uint *__restrict__ const uint arch::n_total

Definition at line 359 of file arch_device_cuda.h.

◆ rslt

__global__ void const T *__restrict__ T *__restrict__ arch::rslt

Definition at line 359 of file arch_device_cuda.h.

◆ size

uint arch::size = NReduStatic ? NReduStatic : 1
constexpr

Definition at line 380 of file arch_device_cuda.h.

◆ temp_storage

BlockReduce::TempStorage * arch::temp_storage = NReduStatic ? temp_storage_static : (typename BlockReduce::TempStorage*) temp_storage_dynamic

Definition at line 384 of file arch_device_cuda.h.

◆ temp_storage_dynamic

__shared__ char arch::temp_storage_dynamic
extern

Definition at line 353 of file arch_device_hip.h.

◆ temp_storage_static

__shared__ BlockReduce::TempStorage arch::temp_storage_static

Definition at line 381 of file arch_device_cuda.h.

◆ thread_data

T * arch::thread_data = NReduStatic ? thread_data_static : &thread_data_dynamic[n_redu_dynamic * idx_glob]

Definition at line 390 of file arch_device_cuda.h.

◆ thread_data_dynamic

__global__ void const T *__restrict__ T *__restrict__ const uint *__restrict__ const uint const uint T * arch::thread_data_dynamic
Initial value:
{
const uint idx_glob = blockIdx.x * blockDim.x + threadIdx.x
arch::uint
uint32_t uint
Definition arch_device_api.h:17

Definition at line 359 of file arch_device_cuda.h.

◆ thread_data_static

T arch::thread_data_static

Definition at line 387 of file arch_device_cuda.h.