Go to the documentation of this file.
36 #include "util_namespace.cuh"
44 #ifndef DOXYGEN_SHOULD_SKIP_THIS // Do not document
50 #define CUB_PTX_ARCH 0
52 #define CUB_PTX_ARCH __CUDA_ARCH__
58 #ifndef CUB_RUNTIME_FUNCTION
59 #if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__>= 350 && defined(__CUDACC_RDC__))
60 #define CUB_RUNTIME_ENABLED
61 #define CUB_RUNTIME_FUNCTION __host__ __device__
63 #define CUB_RUNTIME_FUNCTION __host__
69 #ifndef CUB_LOG_WARP_THREADS
70 #define CUB_LOG_WARP_THREADS(arch) \
72 #define CUB_WARP_THREADS(arch) \
73 (1 << CUB_LOG_WARP_THREADS(arch))
75 #define CUB_PTX_WARP_THREADS CUB_WARP_THREADS(CUB_PTX_ARCH)
76 #define CUB_PTX_LOG_WARP_THREADS CUB_LOG_WARP_THREADS(CUB_PTX_ARCH)
81 #ifndef CUB_LOG_SMEM_BANKS
82 #define CUB_LOG_SMEM_BANKS(arch) \
86 #define CUB_SMEM_BANKS(arch) \
87 (1 << CUB_LOG_SMEM_BANKS(arch))
89 #define CUB_PTX_LOG_SMEM_BANKS CUB_LOG_SMEM_BANKS(CUB_PTX_ARCH)
90 #define CUB_PTX_SMEM_BANKS CUB_SMEM_BANKS(CUB_PTX_ARCH)
95 #ifndef CUB_SUBSCRIPTION_FACTOR
96 #define CUB_SUBSCRIPTION_FACTOR(arch) \
102 #define CUB_PTX_SUBSCRIPTION_FACTOR CUB_SUBSCRIPTION_FACTOR(CUB_PTX_ARCH)
107 #ifndef CUB_PREFER_CONFLICT_OVER_PADDING
108 #define CUB_PREFER_CONFLICT_OVER_PADDING(arch) \
112 #define CUB_PTX_PREFER_CONFLICT_OVER_PADDING CUB_PREFER_CONFLICT_OVER_PADDING(CUB_PTX_ARCH)
117 #define CUB_BLOCK_THREADS(NOMINAL_4B_BLOCK_THREADS, T, PTX_ARCH) \
118 (CUB_MIN(NOMINAL_4B_BLOCK_THREADS, CUB_MAX(3, ((NOMINAL_4B_BLOCK_THREADS / CUB_WARP_THREADS(PTX_ARCH)) * 4) / sizeof(T)) * CUB_WARP_THREADS(PTX_ARCH)))
121 #define CUB_ITEMS_PER_THREAD(NOMINAL_4B_ITEMS_PER_THREAD, NOMINAL_4B_BLOCK_THREADS, T, PTX_ARCH) \
122 (CUB_MIN(NOMINAL_4B_ITEMS_PER_THREAD, CUB_MAX(1, (NOMINAL_4B_ITEMS_PER_THREAD * NOMINAL_4B_BLOCK_THREADS * 4 / sizeof(T)) / CUB_BLOCK_THREADS(NOMINAL_4B_BLOCK_THREADS, T, PTX_ARCH))))
126 #endif // Do not document