NVBIO
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Groups Pages
types.h
Go to the documentation of this file.
1 /*
2  * nvbio
3  * Copyright (c) 2011-2014, 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 #pragma once
29 
30 #include <nvbio/basic/version.h>
31 #include <assert.h>
32 #include <vector_types.h>
33 #include <vector_functions.h>
34 #include <thrust/iterator/detail/host_system_tag.h>
35 #include <thrust/iterator/detail/device_system_tag.h>
36 #include <thrust/iterator/detail/any_system_tag.h>
37 
38 #ifdef __CUDACC__
39  #define NVBIO_HOST_DEVICE __host__ __device__
40  #define NVBIO_HOST __host__
41  #define NVBIO_DEVICE __device__
42 #else
43  #define NVBIO_HOST_DEVICE
44  #define NVBIO_HOST
45  #define NVBIO_DEVICE
46 #endif
47 
48 #ifdef __CUDA_ARCH__
49 #define NVBIO_RESTRICT __restrict__
50 #define NVBIO_SHARED __shared__
51 #else
52 #define NVBIO_RESTRICT
53 #define NVBIO_SHARED
54 #endif
55 
56 #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ > 0
57 #define NVBIO_DEVICE_COMPILATION
58 #endif
59 
60 #define NVBIO_API_CS
61 #define NVBIO_API_SS
62 
63 #ifdef WIN32
64 #define NVBIO_FORCEINLINE __forceinline
65 #else
66 #define NVBIO_FORCEINLINE __inline__
67 #endif
68 
69 #if defined(NVBIO_CUDA_DEBUG)
70 #define NVBIO_CUDA_DEBUG_STATEMENT(x) x
71 #else
72 #define NVBIO_CUDA_DEBUG_STATEMENT(x)
73 #endif
74 
75 #if defined(NVBIO_CUDA_ASSERTS)
76  // the following trickery eliminates the "controlling expression is constant" warning from nvcc when doing assert(!"some string")
77  #define NVBIO_CUDA_ASSERT(x) { const bool __yes = true; assert(x && __yes); }
78  #define NVBIO_CUDA_ASSERT_IF(cond, x, ...) if ((cond) && !(x)) {printf(__VA_ARGS__); NVBIO_CUDA_ASSERT(x); }
79  #define NVBIO_CUDA_DEBUG_ASSERT(x,...) if (!(x)) { printf(__VA_ARGS__); NVBIO_CUDA_ASSERT(x); }
80 #elif defined(NVBIO_CUDA_NON_BLOCKING_ASSERTS) // !defined(NVBIO_CUDA_ASSERTS)
81  #define NVBIO_CUDA_ASSERT(x)
82  #define NVBIO_CUDA_ASSERT_IF(cond, x, ...) if ((cond) && !(x)) { printf(__VA_ARGS__); }
83  #define NVBIO_CUDA_DEBUG_ASSERT(x,...) if (!(x)) { printf(__VA_ARGS__); }
84 #else // !defined(NVBIO_NON_BLOCKING_ASSERTS) && !defined(NVBIO_CUDA_ASSERTS)
85  #define NVBIO_CUDA_ASSERT(x)
86  #define NVBIO_CUDA_ASSERT_IF(cond, x, ...)
87  #define NVBIO_CUDA_DEBUG_ASSERT(x,...)
88 #endif
89 
90 #if defined(NVBIO_CUDA_DEBUG)
91  #define NVBIO_CUDA_DEBUG_PRINT(...) printf(__VA_ARGS__)
92  #define NVBIO_CUDA_DEBUG_PRINT_IF(cond,...) if (cond) printf(__VA_ARGS__)
93  #define NVBIO_CUDA_DEBUG_SELECT(debug_val,normal_val) (debug_val)
94 #else // !defined(NVBIO_CUDA_DEBUG)
95  #define NVBIO_CUDA_DEBUG_PRINT(...)
96  #define NVBIO_CUDA_DEBUG_PRINT_IF(cond,...)
97  #define NVBIO_CUDA_DEBUG_SELECT(debug_val,normal_val) (normal_val)
98 #endif
99 
100 #if defined(NVBIO_CUDA_DEBUG)
101  #if defined(NVBIO_CUDA_ASSERTS)
102  #define NVBIO_CUDA_DEBUG_CHECK_IF(cond, check,...) if ((cond) && (!(check))) { printf(__VA_ARGS__); assert(check); }
103  #else // !defined(NVBIO_CUDA_ASSERTS)
104  #define NVBIO_CUDA_DEBUG_CHECK_IF(cond, check,...) if ((cond) && (!(check))) printf(__VA_ARGS__)
105  #endif
106 #else // !defined(NVBIO_CUDA_DEBUG)
107 #define NVBIO_CUDA_DEBUG_CHECK_IF(cond, check,...)
108 #endif
109 
110 #if defined(__CUDACC__)
111 #define NVBIO_HOST_DEVICE_TEMPLATE \
112 #pragma hd_warning_disable
113 #else
114 #define NVBIO_HOST_DEVICE_TEMPLATE
115 #endif
116 
117 #ifdef WIN32
118 #define WINONLY(x) x
119 #else
120 #define WINONLY(x)
121 #endif
122 
123 // NVBIO_VAR_UNUSED can be prepended to a variable to turn off unused variable warnings
124 // this should only be used when the variable actually is used and the warning is wrong
125 // (e.g., variables which are used only as template parameters for kernel launches)
126 #if defined(__GNUC__)
127 #define NVBIO_VAR_UNUSED __attribute__((unused))
128 #else
129 #define NVBIO_VAR_UNUSED
130 #endif
131 
132 namespace nvbio {
133 
136 
137 typedef unsigned long long uint64;
138 typedef unsigned int uint32;
139 typedef unsigned short uint16;
140 typedef unsigned char uint8;
141 typedef long long int64;
142 typedef int int32;
143 typedef short int16;
144 typedef signed char int8;
145 
150 
153 struct host_tag : public thrust::host_system_tag {};
154 
157 struct device_tag : public thrust::device_system_tag {};
158 
160 
163 struct null_type {};
164 
167 
172 
175 
178 template <typename T> struct to_const { typedef T type; };
179 template <typename T> struct to_const<T&> { typedef const T& type; };
180 template <typename T> struct to_const<T*> { typedef const T* type; };
181 template <typename T> struct to_const<const T&> { typedef const T& type; };
182 template <typename T> struct to_const<const T*> { typedef const T* type; };
183 
186 template <typename T> struct reference_subtype { typedef typename T::reference type; };
187 template <typename T> struct reference_subtype<T*> { typedef T& type; };
188 template <typename T> struct reference_subtype<const T*> { typedef const T& type; };
189 template <> struct reference_subtype<null_type> { typedef null_type type; };
190 
193 template <typename T> struct device_view_subtype { typedef typename T::device_view_type type; };
194 template <> struct device_view_subtype<null_type> { typedef null_type type; };
195 template <typename T> struct device_view_subtype<const T*> { typedef const T* type; };
196 template <typename T> struct device_view_subtype<T*> { typedef T* type; };
197 
200 template <typename T> struct plain_view_subtype { typedef typename T::plain_view_type type; };
201 template <typename T> struct plain_view_subtype<const T> { typedef typename T::const_plain_view_type type; };
202 template <> struct plain_view_subtype<null_type> { typedef null_type type; };
203 template <typename T> struct plain_view_subtype<const T*> { typedef const T* type; };
204 template <typename T> struct plain_view_subtype<T*> { typedef T* type; };
205 
208 template <typename T> struct signed_type {};
209 template <> struct signed_type<uint32> { typedef int32 type; };
210 template <> struct signed_type<uint64> { typedef int64 type; };
211 template <> struct signed_type<int32> { typedef int32 type; };
212 template <> struct signed_type<int64> { typedef int64 type; };
213 
216 template <typename T> struct unsigned_type {};
217 template <> struct unsigned_type<uint32> { typedef uint32 type; };
218 template <> struct unsigned_type<uint64> { typedef uint64 type; };
219 template <> struct unsigned_type<int32> { typedef uint32 type; };
220 template <> struct unsigned_type<int64> { typedef uint64 type; };
221 
224 template <typename T1, typename T2> struct same_type { static const bool pred = false; };
225 template <typename T> struct same_type<T,T> { static const bool pred = true; };
226 
229 template <typename A, typename B>
230 NVBIO_FORCEINLINE NVBIO_HOST_DEVICE bool equal() { return same_type<A,B>::pred; }
231 
234 template <bool predicate, typename T, typename F> struct if_true {};
235 template <typename T, typename F> struct if_true<true,T,F> { typedef T type; };
236 template <typename T, typename F> struct if_true<false,T,F> { typedef F type; };
237 
240 template <typename A, typename B, typename T, typename F> struct if_equal
241 {
242  typedef typename if_true< same_type<A,B>::pred, T, F >::type type;
243 };
244 
247 template <typename A, typename B, uint32 N> struct binary_switch { typedef B type; };
248 
251 template <typename A, typename B> struct binary_switch<A,B,0> { typedef A type; };
252 
254 
257 template <typename Out, typename In>
258 union BinaryCast
259 {
260  In in;
261  Out out;
262 };
263 
266 template <typename Out, typename In>
267 NVBIO_FORCEINLINE NVBIO_HOST_DEVICE Out binary_cast(const In in)
268 {
269 #if defined(__CUDA_ARCH__)
270  return reinterpret_cast<const Out&>(in);
271 #else
272  BinaryCast<Out,In> inout;
273  inout.in = in;
274  return inout.out;
275 #endif
276 }
277 
278 template <uint32 C>
279 NVBIO_FORCEINLINE NVBIO_HOST_DEVICE bool is_pow2() { return (C & (C-1)) == 0u; }
280 
281 // round up to next multiple of N, where N is a power of 2.
282 template <uint32 N, typename I> NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
283 I align(const I a) { return (N > 1) ? I(a + N-1) & I(~(N-1)) : a; }
284 
285 // round down to previous multiple of N, where N is a power of 2.
286 template <uint32 N, typename I> NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
287 I align_down(const I a) { return (N > 1) ? I((a / N) * N) : a; }
288 
291 
292 } // namespace nvbio
293 
294 //
295 // add basic C++ operators to CUDA vector types - it's not nice, but we need it to apply
296 // basic algorithms requiring their presence
297 //
298 
301 NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
302 bool operator==(const uint2 op1, const uint2 op2) { return op1.x == op2.x && op1.y == op2.y; }
303 
306 NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
307 bool operator!=(const uint2 op1, const uint2 op2) { return op1.x != op2.x || op1.y != op2.y; }
308 
311 NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
312 bool operator==(const uint3 op1, const uint3 op2) { return op1.x == op2.x && op1.y == op2.y && op1.z == op2.z; }
313 
316 NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
317 bool operator!=(const uint3 op1, const uint3 op2) { return op1.x != op2.x || op1.y != op2.y || op1.z != op2.z; }
318 
321 NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
322 bool operator==(const uint4 op1, const uint4 op2) { return op1.x == op2.x && op1.y == op2.y && op1.z == op2.z && op1.w == op2.w; }
323 
326 NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
327 bool operator!=(const uint4 op1, const uint4 op2) { return op1.x != op2.x || op1.y != op2.y || op1.z != op2.z || op1.w != op2.w; }
328 
331 NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
332 bool operator==(const int2 op1, const int2 op2) { return op1.x == op2.x && op1.y == op2.y; }
333 
336 NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
337 bool operator!=(const int2 op1, const int2 op2) { return op1.x != op2.x || op1.y != op2.y; }
338 
341 NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
342 bool operator==(const int3 op1, const int3 op2) { return op1.x == op2.x && op1.y == op2.y && op1.z == op2.z; }
343 
346 NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
347 bool operator!=(const int3 op1, const int3 op2) { return op1.x != op2.x || op1.y != op2.y || op1.z != op2.z; }
348 
351 NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
352 bool operator==(const int4 op1, const int4 op2) { return op1.x == op2.x && op1.y == op2.y && op1.z == op2.z && op1.w == op2.w; }
353 
356 NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
357 bool operator!=(const int4 op1, const int4 op2) { return op1.x != op2.x || op1.y != op2.y || op1.z != op2.z || op1.w != op2.w; }
358 
361 NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
362 bool operator==(const float2 op1, const float2 op2) { return op1.x == op2.x && op1.y == op2.y; }
363 
366 NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
367 bool operator!=(const float2 op1, const float2 op2) { return op1.x != op2.x || op1.y != op2.y; }
368 
371 NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
372 bool operator==(const float3 op1, const float3 op2) { return op1.x == op2.x && op1.y == op2.y && op1.z == op2.z; }
373 
376 NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
377 bool operator!=(const float3 op1, const float3 op2) { return op1.x != op2.x || op1.y != op2.y || op1.z != op2.z; }
378 
381 NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
382 bool operator==(const float4 op1, const float4 op2) { return op1.x == op2.x && op1.y == op2.y && op1.z == op2.z && op1.w == op2.w; }
383 
386 NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
387 bool operator!=(const float4 op1, const float4 op2) { return op1.x != op2.x || op1.y != op2.y || op1.z != op2.z || op1.w != op2.w; }
Generated on Wed Feb 25 2015 08:32:57 for NVBIO by   doxygen 1.8.4