#include <cuda.h>#include <cuda_runtime.h>#include <cub/cub.cuh>#include <cub/device/device_radix_sort.cuh>
Go to the source code of this file.
Classes | |
| class | arch::buf< T > |
Namespaces | |
| namespace | arch |
Macros | |
| #define | gpuGetLastError cudaGetLastError |
| #define | gpuGetErrorString cudaGetErrorString |
| #define | gpuPeekAtLastError cudaPeekAtLastError |
| #define | gpuSetDevice cudaSetDevice |
| #define | gpuGetDevice cudaGetDevice |
| #define | gpuGetDeviceCount cudaGetDeviceCount |
| #define | gpuGetDeviceProperties cudaGetDeviceProperties |
| #define | gpuDeviceGetAttribute cudaDeviceGetAttribute |
| #define | gpuDeviceSynchronize cudaDeviceSynchronize |
| #define | gpuDeviceReset cudaDeviceReset |
| #define | gpuCpuDeviceId cudaCpuDeviceId |
| #define | gpuMemGetInfo cudaMemGetInfo |
| #define | gpuDevAttrMaxBlocksPerMultiprocessor cudaDevAttrMaxBlocksPerMultiprocessor |
| #define | gpuFree cudaFree |
| #define | gpuFreeHost cudaFreeHost |
| #define | gpuFreeAsync cudaFreeAsync |
| #define | gpuMalloc cudaMalloc |
| #define | gpuMallocHost cudaMallocHost |
| #define | gpuMallocAsync cudaMallocAsync |
| #define | gpuMallocManaged cudaMallocManaged |
| #define | gpuHostAlloc cudaMallocHost |
| #define | gpuHostAllocPortable cudaHostAllocPortable |
| #define | gpuMemcpy cudaMemcpy |
| #define | gpuMemcpyAsync cudaMemcpyAsync |
| #define | gpuMemset cudaMemset |
| #define | gpuMemsetAsync cudaMemsetAsync |
| #define | gpuHostRegister cudaHostRegister |
| #define | gpuHostRegisterPortable cudaHostRegisterPortable |
| #define | gpuMemAdviseSetAccessedBy cudaMemAdviseSetAccessedBy |
| #define | gpuMemAdviseSetPreferredLocation cudaMemAdviseSetPreferredLocation |
| #define | gpuMemAttachSingle cudaMemAttachSingle |
| #define | gpuMemAttachGlobal cudaMemAttachGlobal |
| #define | gpuMemPrefetchAsync cudaMemPrefetchAsync |
| #define | gpuStreamCreate cudaStreamCreate |
| #define | gpuStreamDestroy cudaStreamDestroy |
| #define | gpuStreamWaitEvent cudaStreamWaitEvent |
| #define | gpuStreamSynchronize cudaStreamSynchronize |
| #define | gpuStreamAttachMemAsync cudaStreamAttachMemAsync |
| #define | gpuDeviceGetStreamPriorityRange cudaDeviceGetStreamPriorityRange |
| #define | gpuStreamCreateWithPriority cudaStreamCreateWithPriority |
| #define | gpuStreamDefault cudaStreamDefault |
| #define | gpuEventCreate cudaEventCreate |
| #define | gpuEventCreateWithFlags cudaEventCreateWithFlags |
| #define | gpuEventDestroy cudaEventDestroy |
| #define | gpuEventQuery cudaEventQuery |
| #define | gpuEventRecord cudaEventRecord |
| #define | gpuEventSynchronize cudaEventSynchronize |
| #define | gpuEventElapsedTime cudaEventElapsedTime |
| #define | gpuError_t cudaError_t |
| #define | gpuSuccess cudaSuccess |
| #define | gpuStream_t cudaStream_t |
| #define | gpuDeviceProp cudaDeviceProp |
| #define | gpuEvent_t cudaEvent_t |
| #define | gpuEventDefault cudaEventDefault |
| #define | gpuEventBlockingSync cudaEventBlockingSync |
| #define | gpuEventDisableTiming cudaEventDisableTiming |
| #define | gpuMemcpyKind cudaMemcpyKind |
| #define | gpuMemcpyDeviceToHost cudaMemcpyDeviceToHost |
| #define | gpuMemcpyHostToDevice cudaMemcpyHostToDevice |
| #define | gpuMemcpyDeviceToDevice cudaMemcpyDeviceToDevice |
| #define | gpuMemcpyHostToHost cudaMemcpyHostToHost |
| #define | gpuMemcpyToSymbol cudaMemcpyToSymbol |
| #define | gpuKernelBallot(mask, input) |
| #define | gpuKernelAny(mask, input) |
| #define | gpuKernelShfl(input, source, mask) |
| #define | gpuKernelShflDown(val, offset) |
| #define | gpuWarpSync() |
| #define | ARCH_LOOP_LAMBDA [=] __host__ __device__ |
| #define | ARCH_INNER_BODY2(i, j, aggregate) |
| #define | ARCH_INNER_BODY3(i, j, k, aggregate) |
| #define | ARCH_INNER_BODY4(i, j, k, l, aggregate) |
| #define | CUB_STDERR |
| #define | ARCH_BLOCKSIZE_R 512 |
| #define | ARCH_BLOCKSIZE_R_SMALL 32 |
| #define | GPUTHREADS (32) |
| #define | WARPSPERBLOCK (32) |
| #define | FULL_MASK 0xffffffff |
| #define | CHK_ERR(err) |
Typedefs | |
| typedef cub::BlockReduce< T, Blocksize, cub::BLOCK_REDUCE_RAKING_COMMUTATIVE_ONLY, 1, 1 > | arch::BlockReduce |
Functions | |
| static void | cuda_error (cudaError_t err, const char *file, int line) |
| __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 | arch::device_mempool_check (uint64_t threshold_new) |
| __host__ static __forceinline__ void * | arch::allocate (size_t bytes) |
| __host__ static __forceinline__ void * | arch::allocate (size_t bytes, cudaStream_t stream) |
| template<typename T> | |
| __host__ static __forceinline__ void | arch::free (T *ptr) |
| template<typename T> | |
| __host__ static __forceinline__ void | arch::free (T *ptr, cudaStream_t stream) |
| template<typename T> | |
| static __forceinline__ void | arch::memcpy_h2d (T *dst, T *src, size_t bytes) |
| template<typename T> | |
| static __forceinline__ void | arch::memcpy_h2d (T *dst, T *src, size_t bytes, cudaStream_t stream) |
| template<typename T> | |
| static __forceinline__ void | arch::memcpy_d2h (T *dst, T *src, size_t bytes) |
| template<typename T> | |
| static __forceinline__ void | arch::memcpy_d2h (T *dst, T *src, size_t bytes, cudaStream_t stream) |
| template<typename T> | |
| static __forceinline__ void | arch::host_register (T *ptr, size_t bytes) |
| template<typename T> | |
| static __forceinline__ void | arch::host_unregister (T *ptr) |
| template<typename Lambda, typename T> | |
| __device__ static __forceinline__ void | arch::lambda_eval (const uint(&idx)[1], T *__restrict__ thread_data, Lambda loop_body) |
| template<typename Lambda, typename T> | |
| __device__ static __forceinline__ void | arch::lambda_eval (const uint(&idx)[2], T *__restrict__ thread_data, Lambda loop_body) |
| template<typename Lambda, typename T> | |
| __device__ static __forceinline__ void | arch::lambda_eval (const uint(&idx)[3], T *__restrict__ thread_data, Lambda loop_body) |
| template<typename Lambda, typename T> | |
| __device__ static __forceinline__ void | arch::lambda_eval (const uint(&idx)[4], T *__restrict__ thread_data, Lambda loop_body) |
| 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) |
| template<uint Blocksize, reduce_op Op, uint NDim, uint NReduStatic, typename Lambda, typename T> | |
| static __global__ void | arch::__launch_bounds__ (ARCH_BLOCKSIZE_R) reduction_kernel(Lambda loop_body |
| arch::if (Op==reduce_op::null) | |
| arch::for (uint i=0;i< n_reductions;i++) | |
| arch::if (idx_glob< n_total) | |
| template<reduce_op Op, uint NReduStatic, uint NDim, typename Lambda, typename T> | |
| static __forceinline__ void | arch::parallel_reduce_driver (const uint(&limits)[NDim], Lambda loop_body, T *sum, const uint n_redu_dynamic) |
Variables | |
| cudaStream_t | gpuStreamList [] |
| static __global__ void const T *__restrict__ | arch::init_val |
| static __global__ void const T *__restrict__ T *__restrict__ | arch::rslt |
| static __global__ void const T *__restrict__ T *__restrict__ const uint *__restrict__ | arch::lims |
| static __global__ void const T *__restrict__ T *__restrict__ const uint *__restrict__ const uint | arch::n_total |
| static __global__ void const T *__restrict__ T *__restrict__ const uint *__restrict__ const uint const uint | arch::n_redu_dynamic |
| static __global__ void const T *__restrict__ T *__restrict__ const uint *__restrict__ const uint const uint T * | arch::thread_data_dynamic |
| __shared__ char | arch::temp_storage_dynamic [] |
| constexpr uint | arch::size = NReduStatic ? NReduStatic : 1 |
| __shared__ BlockReduce::TempStorage | arch::temp_storage_static [size] |
| BlockReduce::TempStorage * | arch::temp_storage = NReduStatic ? temp_storage_static : (typename BlockReduce::TempStorage*) temp_storage_dynamic |
| T | arch::thread_data_static [size] |
| T * | arch::thread_data = NReduStatic ? thread_data_static : &thread_data_dynamic[n_redu_dynamic * idx_glob] |
| const uint | arch::n_reductions = NReduStatic ? NReduStatic : n_redu_dynamic |
Macro Definition Documentation
◆ ARCH_BLOCKSIZE_R
| #define ARCH_BLOCKSIZE_R 512 |
Definition at line 108 of file arch_device_cuda.h.
◆ ARCH_BLOCKSIZE_R_SMALL
| #define ARCH_BLOCKSIZE_R_SMALL 32 |
Definition at line 109 of file arch_device_cuda.h.
◆ ARCH_INNER_BODY2
Definition at line 100 of file arch_device_cuda.h.
◆ ARCH_INNER_BODY3
Definition at line 101 of file arch_device_cuda.h.
◆ ARCH_INNER_BODY4
Definition at line 102 of file arch_device_cuda.h.
◆ ARCH_LOOP_LAMBDA
| #define ARCH_LOOP_LAMBDA [=] __host__ __device__ |
Definition at line 99 of file arch_device_cuda.h.
◆ CHK_ERR
| #define CHK_ERR | ( | err | ) |
Definition at line 123 of file arch_device_cuda.h.
◆ CUB_STDERR
| #define CUB_STDERR |
Definition at line 105 of file arch_device_cuda.h.
◆ FULL_MASK
| #define FULL_MASK 0xffffffff |
Definition at line 118 of file arch_device_cuda.h.
◆ gpuCpuDeviceId
| #define gpuCpuDeviceId cudaCpuDeviceId |
Definition at line 27 of file arch_device_cuda.h.
◆ gpuDevAttrMaxBlocksPerMultiprocessor
| #define gpuDevAttrMaxBlocksPerMultiprocessor cudaDevAttrMaxBlocksPerMultiprocessor |
Definition at line 30 of file arch_device_cuda.h.
◆ gpuDeviceGetAttribute
| #define gpuDeviceGetAttribute cudaDeviceGetAttribute |
Definition at line 24 of file arch_device_cuda.h.
◆ gpuDeviceGetStreamPriorityRange
| #define gpuDeviceGetStreamPriorityRange cudaDeviceGetStreamPriorityRange |
Definition at line 61 of file arch_device_cuda.h.
◆ gpuDeviceProp
| #define gpuDeviceProp cudaDeviceProp |
Definition at line 78 of file arch_device_cuda.h.
◆ gpuDeviceReset
| #define gpuDeviceReset cudaDeviceReset |
Definition at line 26 of file arch_device_cuda.h.
◆ gpuDeviceSynchronize
| #define gpuDeviceSynchronize cudaDeviceSynchronize |
Definition at line 25 of file arch_device_cuda.h.
◆ gpuError_t
| #define gpuError_t cudaError_t |
Definition at line 74 of file arch_device_cuda.h.
◆ gpuEvent_t
| #define gpuEvent_t cudaEvent_t |
Definition at line 80 of file arch_device_cuda.h.
◆ gpuEventBlockingSync
| #define gpuEventBlockingSync cudaEventBlockingSync |
Definition at line 82 of file arch_device_cuda.h.
◆ gpuEventCreate
| #define gpuEventCreate cudaEventCreate |
Definition at line 65 of file arch_device_cuda.h.
◆ gpuEventCreateWithFlags
| #define gpuEventCreateWithFlags cudaEventCreateWithFlags |
Definition at line 66 of file arch_device_cuda.h.
◆ gpuEventDefault
| #define gpuEventDefault cudaEventDefault |
Definition at line 81 of file arch_device_cuda.h.
◆ gpuEventDestroy
| #define gpuEventDestroy cudaEventDestroy |
Definition at line 67 of file arch_device_cuda.h.
◆ gpuEventDisableTiming
| #define gpuEventDisableTiming cudaEventDisableTiming |
Definition at line 83 of file arch_device_cuda.h.
◆ gpuEventElapsedTime
| #define gpuEventElapsedTime cudaEventElapsedTime |
Definition at line 71 of file arch_device_cuda.h.
◆ gpuEventQuery
| #define gpuEventQuery cudaEventQuery |
Definition at line 68 of file arch_device_cuda.h.
◆ gpuEventRecord
| #define gpuEventRecord cudaEventRecord |
Definition at line 69 of file arch_device_cuda.h.
◆ gpuEventSynchronize
| #define gpuEventSynchronize cudaEventSynchronize |
Definition at line 70 of file arch_device_cuda.h.
◆ gpuFree
| #define gpuFree cudaFree |
Definition at line 32 of file arch_device_cuda.h.
◆ gpuFreeAsync
| #define gpuFreeAsync cudaFreeAsync |
Definition at line 34 of file arch_device_cuda.h.
◆ gpuFreeHost
| #define gpuFreeHost cudaFreeHost |
Definition at line 33 of file arch_device_cuda.h.
◆ gpuGetDevice
| #define gpuGetDevice cudaGetDevice |
Definition at line 21 of file arch_device_cuda.h.
◆ gpuGetDeviceCount
| #define gpuGetDeviceCount cudaGetDeviceCount |
Definition at line 22 of file arch_device_cuda.h.
◆ gpuGetDeviceProperties
| #define gpuGetDeviceProperties cudaGetDeviceProperties |
Definition at line 23 of file arch_device_cuda.h.
◆ gpuGetErrorString
| #define gpuGetErrorString cudaGetErrorString |
Definition at line 17 of file arch_device_cuda.h.
◆ gpuGetLastError
| #define gpuGetLastError cudaGetLastError |
Definition at line 16 of file arch_device_cuda.h.
◆ gpuHostAlloc
| #define gpuHostAlloc cudaMallocHost |
Definition at line 40 of file arch_device_cuda.h.
◆ gpuHostAllocPortable
| #define gpuHostAllocPortable cudaHostAllocPortable |
Definition at line 41 of file arch_device_cuda.h.
◆ gpuHostRegister
| #define gpuHostRegister cudaHostRegister |
Definition at line 47 of file arch_device_cuda.h.
◆ gpuHostRegisterPortable
| #define gpuHostRegisterPortable cudaHostRegisterPortable |
Definition at line 48 of file arch_device_cuda.h.
◆ gpuKernelAny
| #define gpuKernelAny | ( | mask, | |
| input ) |
Definition at line 93 of file arch_device_cuda.h.
◆ gpuKernelBallot
| #define gpuKernelBallot | ( | mask, | |
| input ) |
Definition at line 92 of file arch_device_cuda.h.
◆ gpuKernelShfl
| #define gpuKernelShfl | ( | input, | |
| source, | |||
| mask ) |
Definition at line 94 of file arch_device_cuda.h.
◆ gpuKernelShflDown
| #define gpuKernelShflDown | ( | val, | |
| offset ) |
Definition at line 95 of file arch_device_cuda.h.
◆ gpuMalloc
| #define gpuMalloc cudaMalloc |
Definition at line 35 of file arch_device_cuda.h.
◆ gpuMallocAsync
| #define gpuMallocAsync cudaMallocAsync |
Definition at line 37 of file arch_device_cuda.h.
◆ gpuMallocHost
| #define gpuMallocHost cudaMallocHost |
Definition at line 36 of file arch_device_cuda.h.
◆ gpuMallocManaged
| #define gpuMallocManaged cudaMallocManaged |
Definition at line 38 of file arch_device_cuda.h.
◆ gpuMemAdviseSetAccessedBy
| #define gpuMemAdviseSetAccessedBy cudaMemAdviseSetAccessedBy |
Definition at line 50 of file arch_device_cuda.h.
◆ gpuMemAdviseSetPreferredLocation
| #define gpuMemAdviseSetPreferredLocation cudaMemAdviseSetPreferredLocation |
Definition at line 51 of file arch_device_cuda.h.
◆ gpuMemAttachGlobal
| #define gpuMemAttachGlobal cudaMemAttachGlobal |
Definition at line 53 of file arch_device_cuda.h.
◆ gpuMemAttachSingle
| #define gpuMemAttachSingle cudaMemAttachSingle |
Definition at line 52 of file arch_device_cuda.h.
◆ gpuMemcpy
| #define gpuMemcpy cudaMemcpy |
Definition at line 42 of file arch_device_cuda.h.
◆ gpuMemcpyAsync
| #define gpuMemcpyAsync cudaMemcpyAsync |
Definition at line 43 of file arch_device_cuda.h.
◆ gpuMemcpyDeviceToDevice
| #define gpuMemcpyDeviceToDevice cudaMemcpyDeviceToDevice |
Definition at line 88 of file arch_device_cuda.h.
◆ gpuMemcpyDeviceToHost
| #define gpuMemcpyDeviceToHost cudaMemcpyDeviceToHost |
Definition at line 86 of file arch_device_cuda.h.
◆ gpuMemcpyHostToDevice
| #define gpuMemcpyHostToDevice cudaMemcpyHostToDevice |
Definition at line 87 of file arch_device_cuda.h.
◆ gpuMemcpyHostToHost
| #define gpuMemcpyHostToHost cudaMemcpyHostToHost |
Definition at line 89 of file arch_device_cuda.h.
◆ gpuMemcpyKind
| #define gpuMemcpyKind cudaMemcpyKind |
Definition at line 85 of file arch_device_cuda.h.
◆ gpuMemcpyToSymbol
| #define gpuMemcpyToSymbol cudaMemcpyToSymbol |
Definition at line 90 of file arch_device_cuda.h.
◆ gpuMemGetInfo
| #define gpuMemGetInfo cudaMemGetInfo |
Definition at line 28 of file arch_device_cuda.h.
◆ gpuMemPrefetchAsync
| #define gpuMemPrefetchAsync cudaMemPrefetchAsync |
Definition at line 54 of file arch_device_cuda.h.
◆ gpuMemset
| #define gpuMemset cudaMemset |
Definition at line 44 of file arch_device_cuda.h.
◆ gpuMemsetAsync
| #define gpuMemsetAsync cudaMemsetAsync |
Definition at line 45 of file arch_device_cuda.h.
◆ gpuPeekAtLastError
| #define gpuPeekAtLastError cudaPeekAtLastError |
Definition at line 18 of file arch_device_cuda.h.
◆ gpuSetDevice
| #define gpuSetDevice cudaSetDevice |
Definition at line 20 of file arch_device_cuda.h.
◆ gpuStream_t
| #define gpuStream_t cudaStream_t |
Definition at line 77 of file arch_device_cuda.h.
◆ gpuStreamAttachMemAsync
| #define gpuStreamAttachMemAsync cudaStreamAttachMemAsync |
Definition at line 60 of file arch_device_cuda.h.
◆ gpuStreamCreate
| #define gpuStreamCreate cudaStreamCreate |
Definition at line 56 of file arch_device_cuda.h.
◆ gpuStreamCreateWithPriority
| #define gpuStreamCreateWithPriority cudaStreamCreateWithPriority |
Definition at line 62 of file arch_device_cuda.h.
◆ gpuStreamDefault
| #define gpuStreamDefault cudaStreamDefault |
Definition at line 63 of file arch_device_cuda.h.
◆ gpuStreamDestroy
| #define gpuStreamDestroy cudaStreamDestroy |
Definition at line 57 of file arch_device_cuda.h.
◆ gpuStreamSynchronize
| #define gpuStreamSynchronize cudaStreamSynchronize |
Definition at line 59 of file arch_device_cuda.h.
◆ gpuStreamWaitEvent
| #define gpuStreamWaitEvent cudaStreamWaitEvent |
Definition at line 58 of file arch_device_cuda.h.
◆ gpuSuccess
| #define gpuSuccess cudaSuccess |
Definition at line 75 of file arch_device_cuda.h.
◆ GPUTHREADS
| #define GPUTHREADS (32) |
Definition at line 113 of file arch_device_cuda.h.
◆ gpuWarpSync
| #define gpuWarpSync | ( | ) |
Definition at line 96 of file arch_device_cuda.h.
◆ WARPSPERBLOCK
| #define WARPSPERBLOCK (32) |
Definition at line 116 of file arch_device_cuda.h.
Function Documentation
◆ atomicMax() [1/2]
|
static |
Definition at line 132 of file arch_device_cuda.h.
◆ atomicMax() [2/2]
|
static |
Definition at line 154 of file arch_device_cuda.h.
◆ atomicMin() [1/2]
|
static |
Definition at line 143 of file arch_device_cuda.h.
◆ atomicMin() [2/2]
|
static |
Definition at line 165 of file arch_device_cuda.h.
◆ cuda_error()
|
inlinestatic |
Definition at line 124 of file arch_device_cuda.h.
Variable Documentation
◆ gpuStreamList
|
extern |
Definition at line 51 of file gpu_base.cpp.
