CUB
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Groups
util_arch.cuh
Go to the documentation of this file.
1 /******************************************************************************
2  * Copyright (c) 2011, Duane Merrill. All rights reserved.
3  * Copyright (c) 2011-2015, NVIDIA CORPORATION. All rights reserved.
4  *
5  * Redistribution and use in source and binary forms, with or without
6  * modification, are permitted provided that the following conditions are met:
7  * * Redistributions of source code must retain the above copyright
8  * notice, this list of conditions and the following disclaimer.
9  * * Redistributions in binary form must reproduce the above copyright
10  * notice, this list of conditions and the following disclaimer in the
11  * documentation and/or other materials provided with the distribution.
12  * * Neither the name of the NVIDIA CORPORATION nor the
13  * names of its contributors may be used to endorse or promote products
14  * derived from this software without specific prior written permission.
15  *
16  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
17  * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
18  * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
19  * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
20  * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
21  * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
22  * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
23  * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
24  * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
25  * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
26  *
27  ******************************************************************************/
28 
34 #pragma once
35 
36 #include "util_namespace.cuh"
37 
39 CUB_NS_PREFIX
40 
42 namespace cub {
43 
44 #ifndef DOXYGEN_SHOULD_SKIP_THIS // Do not document
45 
46 
48 #ifndef CUB_PTX_ARCH
49  #ifndef __CUDA_ARCH__
50  #define CUB_PTX_ARCH 0
51  #else
52  #define CUB_PTX_ARCH __CUDA_ARCH__
53  #endif
54 #endif
55 
56 
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__
62  #else
63  #define CUB_RUNTIME_FUNCTION __host__
64  #endif
65 #endif
66 
67 
69 #ifndef CUB_LOG_WARP_THREADS
70  #define CUB_LOG_WARP_THREADS(arch) \
71  (5)
72  #define CUB_WARP_THREADS(arch) \
73  (1 << CUB_LOG_WARP_THREADS(arch))
74 
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)
77 #endif
78 
79 
81 #ifndef CUB_LOG_SMEM_BANKS
82  #define CUB_LOG_SMEM_BANKS(arch) \
83  ((arch >= 200) ? \
84  (5) : \
85  (4))
86  #define CUB_SMEM_BANKS(arch) \
87  (1 << CUB_LOG_SMEM_BANKS(arch))
88 
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)
91 #endif
92 
93 
95 #ifndef CUB_SUBSCRIPTION_FACTOR
96  #define CUB_SUBSCRIPTION_FACTOR(arch) \
97  ((arch >= 300) ? \
98  (5) : \
99  ((arch >= 200) ? \
100  (3) : \
101  (10)))
102  #define CUB_PTX_SUBSCRIPTION_FACTOR CUB_SUBSCRIPTION_FACTOR(CUB_PTX_ARCH)
103 #endif
104 
105 
107 #ifndef CUB_PREFER_CONFLICT_OVER_PADDING
108  #define CUB_PREFER_CONFLICT_OVER_PADDING(arch) \
109  ((arch >= 300) ? \
110  (1) : \
111  (4))
112  #define CUB_PTX_PREFER_CONFLICT_OVER_PADDING CUB_PREFER_CONFLICT_OVER_PADDING(CUB_PTX_ARCH)
113 #endif
114 
115 
116 #endif // Do not document
117 
118 } // CUB namespace
119 CUB_NS_POSTFIX // Optional outer namespace(s)