OpenVDB  12.0.0
Util.h
Go to the documentation of this file.
1 // Copyright Contributors to the OpenVDB Project
2 // SPDX-License-Identifier: Apache-2.0
3 
4 /*!
5  \file nanovdb/util/Util.h
6 
7  \author Ken Museth
8 
9  \date January 8, 2020
10 
11  \brief Utility functions
12 */
13 
14 #ifndef NANOVDB_UTIL_UTIL_H_HAS_BEEN_INCLUDED
15 #define NANOVDB_UTIL_UTIL_H_HAS_BEEN_INCLUDED
16 
17 #ifdef __CUDACC_RTC__
18 
19 typedef signed char int8_t;
20 typedef short int16_t;
21 typedef int int32_t;
22 typedef long long int64_t;
23 typedef unsigned char uint8_t;
24 typedef unsigned int uint32_t;
25 typedef unsigned short uint16_t;
26 typedef unsigned long long uint64_t;
27 
28 #define NANOVDB_ASSERT(x)
29 
30 #ifndef UINT64_C
31 #define UINT64_C(x) (x ## ULL)
32 #endif
33 
34 #else // !__CUDACC_RTC__
35 
36 #include <stdlib.h> // for abs in clang7
37 #include <stdint.h> // for types like int32_t etc
38 #include <stddef.h> // for size_t type
39 #include <cassert> // for assert
40 #include <cstdio> // for stderr and snprintf
41 #include <cmath> // for sqrt and fma
42 #include <limits> // for numeric_limits
43 #include <utility>// for std::move
44 #ifdef NANOVDB_USE_IOSTREAMS
45 #include <fstream>// for read/writeUncompressedGrids
46 #endif// ifdef NANOVDB_USE_IOSTREAMS
47 
48 // All asserts can be disabled here, even for debug builds
49 #if 1
50 #define NANOVDB_ASSERT(x) assert(x)
51 #else
52 #define NANOVDB_ASSERT(x)
53 #endif
54 
55 #if defined(NANOVDB_USE_INTRINSICS) && defined(_MSC_VER)
56 #include <intrin.h>
57 #pragma intrinsic(_BitScanReverse)
58 #pragma intrinsic(_BitScanForward)
59 #pragma intrinsic(_BitScanReverse64)
60 #pragma intrinsic(_BitScanForward64)
61 #endif
62 
63 #endif // __CUDACC_RTC__
64 
65 #if defined(__CUDACC__) || defined(__HIP__)
66 // Only define __hostdev__ qualifier when using NVIDIA CUDA or HIP compilers
67 #ifndef __hostdev__
68 #define __hostdev__ __host__ __device__ // Runs on the CPU and GPU, called from the CPU or the GPU
69 #endif
70 #else
71 // Dummy definitions of macros only defined by CUDA and HIP compilers
72 #ifndef __hostdev__
73 #define __hostdev__ // Runs on the CPU and GPU, called from the CPU or the GPU
74 #endif
75 #ifndef __global__
76 #define __global__ // Runs on the GPU, called from the CPU or the GPU
77 #endif
78 #ifndef __device__
79 #define __device__ // Runs on the GPU, called from the GPU
80 #endif
81 #ifndef __host__
82 #define __host__ // Runs on the CPU, called from the CPU
83 #endif
84 
85 #endif // if defined(__CUDACC__) || defined(__HIP__)
86 
87 // The following macro will suppress annoying warnings when nvcc
88 // compiles functions that call (host) intrinsics (which is perfectly valid)
89 #if defined(_MSC_VER) && defined(__CUDACC__)
90 #define NANOVDB_HOSTDEV_DISABLE_WARNING __pragma("hd_warning_disable")
91 #elif defined(__GNUC__) && defined(__CUDACC__)
92 #define NANOVDB_HOSTDEV_DISABLE_WARNING _Pragma("hd_warning_disable")
93 #else
94 #define NANOVDB_HOSTDEV_DISABLE_WARNING
95 #endif
96 
97 // Define compiler warnings that work with all compilers
98 //#if defined(_MSC_VER)
99 //#define NANO_WARNING(msg) _pragma("message" #msg)
100 //#else
101 //#define NANO_WARNING(msg) _Pragma("message" #msg)
102 //#endif
103 
104 //==============================================
105 /// @brief Defines macros that issues warnings for deprecated header files
106 /// @details Example:
107 /// @code
108 /// #include <nanovdb/util/Util.h> // for NANOVDB_DEPRECATED_HEADER
109 /// #include <nanovdb/path/Alternative.h>
110 /// NANOVDB_DEPRECATED_HEADER("This header file is deprecated, please use <nanovdb/path/Alternative.h> instead")
111 /// @endcode
112 #ifdef __GNUC__
113 #define NANOVDB_PRAGMA(X) _Pragma(#X)
114 #define NANOVDB_DEPRECATED_HEADER(MSG) NANOVDB_PRAGMA(GCC warning MSG)
115 #elif defined(_MSC_VER)
116 #define NANOVDB_STRINGIZE_(MSG) #MSG
117 #define NANOVDB_STRINGIZE(MSG) NANOVDB_STRINGIZE_(MSG)
118 #define NANOVDB_DEPRECATED_HEADER(MSG) \
119  __pragma(message(__FILE__ "(" NANOVDB_STRINGIZE(__LINE__) ") : Warning: " MSG))
120 #endif
121 
122 // A portable implementation of offsetof - unfortunately it doesn't work with static_assert
123 #define NANOVDB_OFFSETOF(CLASS, MEMBER) ((int)(size_t)((char*)&((CLASS*)0)->MEMBER - (char*)0))
124 
125 namespace nanovdb {// =================================================================
126 
127 namespace util {// ====================================================================
128 
129 /// @brief Minimal implementation of std::declval, which converts any type @c T to
130 //// a reference type, making it possible to use member functions in the operand
131 /// of the decltype specifier without the need to go through constructors.
132 /// @tparam T Template type to be converted to T&&
133 /// @return T&&
134 /// @warning Unlike std::declval, this version does not work when T = void! However,
135 /// NVRTC does not like std::declval, so we provide our own implementation.
136 template<typename T>
137 T&& declval() noexcept;
138 
139 // --------------------------> string utility functions <------------------------------------
140 
141 /// @brief tests if a c-string @c str is empty, that is its first value is '\0'
142 /// @param str c-string to be tested for null termination
143 /// @return true if str[0] = '\0'
144 __hostdev__ inline bool empty(const char* str)
145 {
146  NANOVDB_ASSERT(str != nullptr);
147  return *str == '\0';
148 }// util::empty
149 
150 /// @brief length of a c-sting, excluding '\0'.
151 /// @param str c-string
152 /// @return the number of characters that precede the terminating null character.
153 __hostdev__ inline size_t strlen(const char *str)
154 {
155  NANOVDB_ASSERT(str != nullptr);
156  const char *s = str;
157  while(*s) ++s; ;
158  return (s - str);
159 }// util::strlen
160 
161 /// @brief Copy characters from @c src to @c dst.
162 /// @param dst pointer to the destination string.
163 /// @param src pointer to the null-terminated source string.
164 /// @return destination string @c dst.
165 /// @note Emulates the behaviour of std::strcpy, except this version also runs on the GPU.
166 __hostdev__ inline char* strcpy(char *dst, const char *src)
167 {
168  NANOVDB_ASSERT(dst != nullptr && src != nullptr);
169  for (char *p = dst; (*p++ = *src) != '\0'; ++src);
170  return dst;
171 }// util::strcpy(char*, const char*)
172 
173 /// @brief Copies the first num characters of @c src to @c dst.
174 /// If the end of the source C string (which is signaled by a
175 /// null-character) is found before @c max characters have been
176 /// copied, @c dst is padded with zeros until a total of @c max
177 /// characters have been written to it.
178 /// @param dst destination string
179 /// @param src source string
180 /// @param max maximum number of character in destination string
181 /// @return destination string @c dst
182 /// @warning if strncpy(dst, src, max)[max-1]!='\0' then @c src has more
183 /// characters than @c max and the return string needs to be
184 /// manually null-terminated, i.e. strncpy(dst, src, max)[max-1]='\0'
185 __hostdev__ inline char* strncpy(char *dst, const char *src, size_t max)
186 {
187  NANOVDB_ASSERT(dst != nullptr && src != nullptr);
188  size_t i = 0;
189  for (; i < max && src[i] != '\0'; ++i) dst[i] = src[i];
190  for (; i < max; ++i) dst[i] = '\0';
191  return dst;
192 }// util::strncpy(char *dst, const char *src, size_t max)
193 
194 /// @brief converts a number to a string using a specific base
195 /// @param dst destination string
196 /// @param num signed number to be concatenated after @c dst
197 /// @param bas base used when converting @c num to a string
198 /// @return destination string @c dst
199 /// @note Emulates the behaviour of itoa, except this verion also works on the GPU.
200 __hostdev__ inline char* strcpy(char* dst, int num, int bas = 10)
201 {
202  NANOVDB_ASSERT(dst != nullptr && bas > 0);
203  int len = 0;// length of number once converted to a string
204  if (num == 0) dst[len++] = '0';
205  for (int abs = num < 0 && bas == 10 ? -num : num; abs; abs /= bas) {
206  const int rem = abs % bas;
207  dst[len++] = rem > 9 ? rem - 10 + 'a' : rem + '0';
208  }
209  if (num < 0) dst[len++] = '-';// append '-' if negative
210  for (char *a = dst, *b = a + len - 1; a < b; ++a, --b) {// reverse dst
211  dst[len] = *a;// use end of string as temp
212  *a = *b;
213  *b = dst[len];
214  }
215  dst[len] = '\0';// explicitly terminate end of string
216  return dst;
217 }// util::strcpy(char*, int, int)
218 
219 /// @brief Appends a copy of the character string pointed to by @c src to
220 /// the end of the character string pointed to by @c dst on the device.
221 /// @param dst pointer to the null-terminated byte string to append to.
222 /// @param src pointer to the null-terminated byte string to copy from.
223 /// @return pointer to the character array being appended to.
224 /// @note Emulates the behaviour of std::strcat, except this version also runs on the GPU.
225 __hostdev__ inline char* strcat(char *dst, const char *src)
226 {
227  NANOVDB_ASSERT(dst != nullptr && src != nullptr);
228  char *p = dst;
229  while (*p != '\0') ++p;// advance till end of dst
230  strcpy(p, src);// append src
231  return dst;
232 }// util::strcat(char*, const char*)
233 
234 /// @brief concatenates a number after a string using a specific base
235 /// @param dst null terminated destination string
236 /// @param num signed number to be concatenated after @c dst
237 /// @param bas base used when converting @c num to a string
238 /// @return destination string @c dst
239 __hostdev__ inline char* strcat(char* dst, int num, int bas = 10)
240 {
241  NANOVDB_ASSERT(dst != nullptr);
242  char *p = dst;
243  while (*p != '\0') ++p;
244  strcpy(p, num, bas);
245  return dst;
246 }// util::strcat(char*, int, int)
247 
248 /// @brief Compares two null-terminated byte strings lexicographically.
249 /// @param lhs pointer to the null-terminated byte strings to compare
250 /// @param rhs pointer to the null-terminated byte strings to compare
251 /// @return Negative value if @c lhs appears before @c rhs in lexicographical order.
252 /// Zero if @c lhs and @c rhs compare equal. Positive value if @c lhs appears
253 /// after @c rhs in lexicographical order.
254 /// @note Emulates the behaviour of std::strcmp, except this version also runs on the GPU.
255 __hostdev__ inline int strcmp(const char *lhs, const char *rhs)
256 {
257  while(*lhs != '\0' && (*lhs == *rhs)){
258  lhs++;
259  rhs++;
260  }
261  return *(const unsigned char*)lhs - *(const unsigned char*)rhs;// zero if lhs == rhs
262 }// util::strcmp(const char*, const char*)
263 
264 /// @brief Test if two null-terminated byte strings are the same
265 /// @param lhs pointer to the null-terminated byte strings to compare
266 /// @param rhs pointer to the null-terminated byte strings to compare
267 /// @return true if the two c-strings are identical
268 __hostdev__ inline bool streq(const char *lhs, const char *rhs)
269 {
270  return strcmp(lhs, rhs) == 0;
271 }// util::streq
272 
273 namespace impl {// =======================================================
274 // Base-case implementation of Variadic Template function impl::sprint
275 __hostdev__ inline char* sprint(char *dst){return dst;}
276 // Variadic Template function impl::sprint
277 template <typename T, typename... Types>
278 __hostdev__ inline char* sprint(char *dst, T var1, Types... var2)
279 {
280  return impl::sprint(strcat(dst, var1), var2...);
281 }
282 }// namespace impl =========================================================
283 
284 /// @brief prints a variable number of string and/or numbers to a destination string
285 template <typename T, typename... Types>
286 __hostdev__ inline char* sprint(char *dst, T var1, Types... var2)
287 {
288  return impl::sprint(strcpy(dst, var1), var2...);
289 }// util::sprint
290 
291 // --------------------------> memzero <------------------------------------
292 
293 /// @brief Zero initialization of memory
294 /// @param dst pointer to destination
295 /// @param byteCount number of bytes to be initialized to zero
296 /// @return destination pointer @c dst
297 __hostdev__ inline static void* memzero(void *dst, size_t byteCount)
298 {
299  NANOVDB_ASSERT(dst);
300  const size_t wordCount = byteCount >> 3;
301  if (wordCount << 3 == byteCount) {
302  for (auto *d = (uint64_t*)dst, *e = d + wordCount; d != e; ++d) *d = 0ULL;
303  } else {
304  for (auto *d = (char*)dst, *e = d + byteCount; d != e; ++d) *d = '\0';
305  }
306  return dst;
307 }// util::memzero
308 
309 // --------------------------> util::is_same <------------------------------------
310 
311 /// @brief C++11 implementation of std::is_same
312 /// @note When more than two arguments are provided value = T0==T1 || T0==T2 || ...
313 template<typename T0, typename T1, typename ...T>
314 struct is_same
315 {
316  static constexpr bool value = is_same<T0, T1>::value || is_same<T0, T...>::value;
317 };
318 
319 template<typename T0, typename T1>
320 struct is_same<T0, T1> {static constexpr bool value = false;};
321 
322 template<typename T>
323 struct is_same<T, T> {static constexpr bool value = true;};
324 
325 // --------------------------> util::is_floating_point <------------------------------------
326 
327 /// @brief C++11 implementation of std::is_floating_point
328 template<typename T>
329 struct is_floating_point {static constexpr bool value = is_same<T, float, double>::value;};
330 
331 // --------------------------> util::enable_if <------------------------------------
332 
333 /// @brief C++11 implementation of std::enable_if
334 template <bool, typename T = void>
335 struct enable_if {};
336 
337 template <typename T>
338 struct enable_if<true, T> {using type = T;};
339 
340 // --------------------------> util::disable_if <------------------------------------
341 
342 template<bool, typename T = void>
343 struct disable_if {using type = T;};
344 
345 template<typename T>
346 struct disable_if<true, T> {};
347 
348 // --------------------------> util::is_const <------------------------------------
349 
350 template<typename T>
351 struct is_const {static constexpr bool value = false;};
352 
353 template<typename T>
354 struct is_const<const T> {static constexpr bool value = true;};
355 
356 // --------------------------> util::is_pointer <------------------------------------
357 
358 /// @brief Trait used to identify template parameter that are pointers
359 /// @tparam T Template parameter to be tested
360 template<class T>
361 struct is_pointer {static constexpr bool value = false;};
362 
363 /// @brief Template specialization of pointers
364 /// @tparam T Template parameter to be tested
365 /// @note T can be both a non-const and const type
366 template<class T>
367 struct is_pointer<T*> {static constexpr bool value = true;};
368 
369 // --------------------------> util::conditional <------------------------------------
370 
371 /// @brief C++11 implementation of std::conditional
372 template<bool, class TrueT, class FalseT>
373 struct conditional { using type = TrueT; };
374 
375 /// @brief Template specialization of conditional
376 /// @tparam FalseT Type used when boolean is false
377 /// @tparam TrueT Type used when boolean is true
378 template<class TrueT, class FalseT>
379 struct conditional<false, TrueT, FalseT> { using type = FalseT; };
380 
381 // --------------------------> util::remove_const <------------------------------------
382 
383 /// @brief Trait use to const from type. Default implementation is just a pass-through
384 /// @tparam T Type
385 /// @details remove_pointer<float>::type = float
386 template<typename T>
387 struct remove_const {using type = T;};
388 
389 /// @brief Template specialization of trait class use to remove const qualifier type from a type
390 /// @tparam T Type of the const type
391 /// @details remove_pointer<const float>::type = float
392 template<typename T>
393 struct remove_const<const T> {using type = T;};
394 
395 // --------------------------> util::remove_reference <------------------------------------
396 
397 /// @brief Trait use to remove reference, i.e. "&", qualifier from a type. Default implementation is just a pass-through
398 /// @tparam T Type
399 /// @details remove_pointer<float>::type = float
400 template <typename T>
401 struct remove_reference {using type = T;};
402 
403 /// @brief Template specialization of trait class use to remove reference, i.e. "&", qualifier from a type
404 /// @tparam T Type of the reference
405 /// @details remove_pointer<float&>::type = float
406 template <typename T>
407 struct remove_reference<T&> {using type = T;};
408 
409 // --------------------------> util::remove_pointer <------------------------------------
410 
411 /// @brief Trait use to remove pointer, i.e. "*", qualifier from a type. Default implementation is just a pass-through
412 /// @tparam T Type
413 /// @details remove_pointer<float>::type = float
414 template <typename T>
415 struct remove_pointer {using type = T;};
416 
417 /// @brief Template specialization of trait class use to to remove pointer, i.e. "*", qualifier from a type
418 /// @tparam T Type of the pointer
419 /// @details remove_pointer<float*>::type = float
420 template <typename T>
421 struct remove_pointer<T*> {using type = T;};
422 
423 // --------------------------> util::match_const <------------------------------------
424 
425 /// @brief Trait used to transfer the const-ness of a reference type to another type
426 /// @tparam T Type whose const-ness needs to match the reference type
427 /// @tparam ReferenceT Reference type that is not const
428 /// @details match_const<const int, float>::type = int
429 /// match_const<int, float>::type = int
430 template<typename T, typename ReferenceT>
431 struct match_const {using type = typename remove_const<T>::type;};
432 
433 /// @brief Template specialization used to transfer the const-ness of a reference type to another type
434 /// @tparam T Type that will adopt the const-ness of the reference type
435 /// @tparam ReferenceT Reference type that is const
436 /// @details match_const<const int, const float>::type = const int
437 /// match_const<int, const float>::type = const int
438 template<typename T, typename ReferenceT>
439 struct match_const<T, const ReferenceT> {using type = const typename remove_const<T>::type;};
440 
441 // --------------------------> util::is_specialization <------------------------------------
442 
443 /// @brief Metafunction used to determine if the first template
444 /// parameter is a specialization of the class template
445 /// given in the second template parameter.
446 ///
447 /// @details is_specialization<Vec3<float>, Vec3>::value == true;
448 /// is_specialization<Vec3f, Vec3>::value == true;
449 /// is_specialization<std::vector<float>, std::vector>::value == true;
450 template<typename AnyType, template<typename...> class TemplateType>
451 struct is_specialization {static const bool value = false;};
452 template<typename... Args, template<typename...> class TemplateType>
453 struct is_specialization<TemplateType<Args...>, TemplateType>
454 {
455  static const bool value = true;
456 };// util::is_specialization
457 
458 // --------------------------> util::PtrDiff <------------------------------------
459 
460 /// @brief Compute the distance, in bytes, between two pointers, dist = p - q
461 /// @param p fist pointer, assumed to NOT be NULL
462 /// @param q second pointer, assumed to NOT be NULL
463 /// @return signed distance between pointer, p - q, addresses in units of bytes
464 __hostdev__ inline static int64_t PtrDiff(const void* p, const void* q)
465 {
466  NANOVDB_ASSERT(p && q);
467  return reinterpret_cast<const char*>(p) - reinterpret_cast<const char*>(q);
468 }// util::PtrDiff
469 
470 // --------------------------> util::PtrAdd <------------------------------------
471 
472 /// @brief Adds a byte offset to a non-const pointer to produce another non-const pointer
473 /// @tparam DstT Type of the return pointer (defaults to void)
474 /// @param p non-const input pointer, assumed to NOT be NULL
475 /// @param offset signed byte offset
476 /// @return a non-const pointer defined as the offset of an input pointer
477 template<typename DstT = void>
478 __hostdev__ inline static DstT* PtrAdd(void* p, int64_t offset)
479 {
480  NANOVDB_ASSERT(p);
481  return reinterpret_cast<DstT*>(reinterpret_cast<char*>(p) + offset);
482 }// util::PtrAdd
483 
484 /// @brief Adds a byte offset to a const pointer to produce another const pointer
485 /// @tparam DstT Type of the return pointer (defaults to void)
486 /// @param p const input pointer, assumed to NOT be NULL
487 /// @param offset signed byte offset
488 /// @return a const pointer defined as the offset of a const input pointer
489 template<typename DstT = void>
490 __hostdev__ inline static const DstT* PtrAdd(const void* p, int64_t offset)
491 {
492  NANOVDB_ASSERT(p);
493  return reinterpret_cast<const DstT*>(reinterpret_cast<const char*>(p) + offset);
494 }// util::PtrAdd
495 
496 // -------------------> findLowestOn <----------------------------
497 
498 /// @brief Returns the index of the lowest, i.e. least significant, on bit in the specified 32 bit word
499 ///
500 /// @warning Assumes that at least one bit is set in the word, i.e. @a v != uint32_t(0)!
502 __hostdev__ inline uint32_t findLowestOn(uint32_t v)
503 {
504  NANOVDB_ASSERT(v);
505 #if (defined(__CUDA_ARCH__) || defined(__HIP__)) && defined(NANOVDB_USE_INTRINSICS)
506  return __ffs(v) - 1; // one based indexing
507 #elif defined(_MSC_VER) && defined(NANOVDB_USE_INTRINSICS)
508  unsigned long index;
509  _BitScanForward(&index, v);
510  return static_cast<uint32_t>(index);
511 #elif (defined(__GNUC__) || defined(__clang__)) && defined(NANOVDB_USE_INTRINSICS)
512  return static_cast<uint32_t>(__builtin_ctzl(v));
513 #else
514  //NANO_WARNING("Using software implementation for findLowestOn(uint32_t v)")
515  static const unsigned char DeBruijn[32] = {
516  0, 1, 28, 2, 29, 14, 24, 3, 30, 22, 20, 15, 25, 17, 4, 8, 31, 27, 13, 23, 21, 19, 16, 7, 26, 12, 18, 6, 11, 5, 10, 9};
517 // disable unary minus on unsigned warning
518 #if defined(_MSC_VER) && !defined(__NVCC__)
519 #pragma warning(push)
520 #pragma warning(disable : 4146)
521 #endif
522  return DeBruijn[uint32_t((v & -v) * 0x077CB531U) >> 27];
523 #if defined(_MSC_VER) && !defined(__NVCC__)
524 #pragma warning(pop)
525 #endif
526 
527 #endif
528 }// util::findLowestOn(uint32_t)
529 
530 /// @brief Returns the index of the lowest, i.e. least significant, on bit in the specified 64 bit word
531 ///
532 /// @warning Assumes that at least one bit is set in the word, i.e. @a v != uint32_t(0)!
534 __hostdev__ inline uint32_t findLowestOn(uint64_t v)
535 {
536  NANOVDB_ASSERT(v);
537 #if (defined(__CUDA_ARCH__) || defined(__HIP__)) && defined(NANOVDB_USE_INTRINSICS)
538  return __ffsll(static_cast<unsigned long long int>(v)) - 1; // one based indexing
539 #elif defined(_MSC_VER) && defined(NANOVDB_USE_INTRINSICS)
540  unsigned long index;
541  _BitScanForward64(&index, v);
542  return static_cast<uint32_t>(index);
543 #elif (defined(__GNUC__) || defined(__clang__)) && defined(NANOVDB_USE_INTRINSICS)
544  return static_cast<uint32_t>(__builtin_ctzll(v));
545 #else
546  //NANO_WARNING("Using software implementation for util::findLowestOn(uint64_t)")
547  static const unsigned char DeBruijn[64] = {
548  0, 1, 2, 53, 3, 7, 54, 27, 4, 38, 41, 8, 34, 55, 48, 28,
549  62, 5, 39, 46, 44, 42, 22, 9, 24, 35, 59, 56, 49, 18, 29, 11,
550  63, 52, 6, 26, 37, 40, 33, 47, 61, 45, 43, 21, 23, 58, 17, 10,
551  51, 25, 36, 32, 60, 20, 57, 16, 50, 31, 19, 15, 30, 14, 13, 12,
552  };
553 // disable unary minus on unsigned warning
554 #if defined(_MSC_VER) && !defined(__NVCC__)
555 #pragma warning(push)
556 #pragma warning(disable : 4146)
557 #endif
558  return DeBruijn[uint64_t((v & -v) * UINT64_C(0x022FDD63CC95386D)) >> 58];
559 #if defined(_MSC_VER) && !defined(__NVCC__)
560 #pragma warning(pop)
561 #endif
562 
563 #endif
564 }// util::findLowestOn(uint64_t)
565 
566 // -------------------> findHighestOn <----------------------------
567 
568 /// @brief Returns the index of the highest, i.e. most significant, on bit in the specified 32 bit word
569 ///
570 /// @warning Assumes that at least one bit is set in the word, i.e. @a v != uint32_t(0)!
572 __hostdev__ inline uint32_t findHighestOn(uint32_t v)
573 {
574  NANOVDB_ASSERT(v);
575 #if (defined(__CUDA_ARCH__) || defined(__HIP__)) && defined(NANOVDB_USE_INTRINSICS)
576  return sizeof(uint32_t) * 8 - 1 - __clz(v); // Return the number of consecutive high-order zero bits in a 32-bit integer.
577 #elif defined(_MSC_VER) && defined(NANOVDB_USE_INTRINSICS)
578  unsigned long index;
579  _BitScanReverse(&index, v);
580  return static_cast<uint32_t>(index);
581 #elif (defined(__GNUC__) || defined(__clang__)) && defined(NANOVDB_USE_INTRINSICS)
582  return sizeof(unsigned long) * 8 - 1 - __builtin_clzl(v);
583 #else
584  //NANO_WARNING("Using software implementation for util::findHighestOn(uint32_t)")
585  static const unsigned char DeBruijn[32] = {
586  0, 9, 1, 10, 13, 21, 2, 29, 11, 14, 16, 18, 22, 25, 3, 30,
587  8, 12, 20, 28, 15, 17, 24, 7, 19, 27, 23, 6, 26, 5, 4, 31};
588  v |= v >> 1; // first round down to one less than a power of 2
589  v |= v >> 2;
590  v |= v >> 4;
591  v |= v >> 8;
592  v |= v >> 16;
593  return DeBruijn[uint32_t(v * 0x07C4ACDDU) >> 27];
594 #endif
595 }// util::findHighestOn
596 
597 /// @brief Returns the index of the highest, i.e. most significant, on bit in the specified 64 bit word
598 ///
599 /// @warning Assumes that at least one bit is set in the word, i.e. @a v != uint32_t(0)!
601 __hostdev__ inline uint32_t findHighestOn(uint64_t v)
602 {
603  NANOVDB_ASSERT(v);
604 #if (defined(__CUDA_ARCH__) || defined(__HIP__)) && defined(NANOVDB_USE_INTRINSICS)
605  return sizeof(unsigned long) * 8 - 1 - __clzll(static_cast<unsigned long long int>(v));
606 #elif defined(_MSC_VER) && defined(NANOVDB_USE_INTRINSICS)
607  unsigned long index;
608  _BitScanReverse64(&index, v);
609  return static_cast<uint32_t>(index);
610 #elif (defined(__GNUC__) || defined(__clang__)) && defined(NANOVDB_USE_INTRINSICS)
611  return sizeof(unsigned long) * 8 - 1 - __builtin_clzll(v);
612 #else
613  const uint32_t* p = reinterpret_cast<const uint32_t*>(&v);
614  return p[1] ? 32u + findHighestOn(p[1]) : findHighestOn(p[0]);
615 #endif
616 }// util::findHighestOn
617 
618 // ----------------------------> util::countOn <--------------------------------------
619 
620 /// @return Number of bits that are on in the specified 64-bit word
622 __hostdev__ inline uint32_t countOn(uint64_t v)
623 {
624 #if (defined(__CUDA_ARCH__) || defined(__HIP__)) && defined(NANOVDB_USE_INTRINSICS)
625  //#warning Using popcll for util::countOn
626  return __popcll(v);
627 // __popcnt64 intrinsic support was added in VS 2019 16.8
628 #elif defined(_MSC_VER) && defined(_M_X64) && (_MSC_VER >= 1928) && defined(NANOVDB_USE_INTRINSICS)
629  //#warning Using popcnt64 for util::countOn
630  return uint32_t(__popcnt64(v));
631 #elif (defined(__GNUC__) || defined(__clang__)) && defined(NANOVDB_USE_INTRINSICS)
632  //#warning Using builtin_popcountll for util::countOn
633  return __builtin_popcountll(v);
634 #else // use software implementation
635  //NANO_WARNING("Using software implementation for util::countOn")
636  v = v - ((v >> 1) & uint64_t(0x5555555555555555));
637  v = (v & uint64_t(0x3333333333333333)) + ((v >> 2) & uint64_t(0x3333333333333333));
638  return (((v + (v >> 4)) & uint64_t(0xF0F0F0F0F0F0F0F)) * uint64_t(0x101010101010101)) >> 56;
639 #endif
640 }// util::countOn(uint64_t)
641 
642 }// namespace util ==================================================================
643 
644 [[deprecated("Use nanovdb::util::findLowestOn instead")]]
645 __hostdev__ inline uint32_t FindLowestOn(uint32_t v){return util::findLowestOn(v);}
646 [[deprecated("Use nanovdb::util::findLowestOn instead")]]
647 __hostdev__ inline uint32_t FindLowestOn(uint64_t v){return util::findLowestOn(v);}
648 [[deprecated("Use nanovdb::util::findHighestOn instead")]]
649 __hostdev__ inline uint32_t FindHighestOn(uint32_t v){return util::findHighestOn(v);}
650 [[deprecated("Use nanovdb::util::findHighestOn instead")]]
651 __hostdev__ inline uint32_t FindHighestOn(uint64_t v){return util::findHighestOn(v);}
652 [[deprecated("Use nanovdb::util::countOn instead")]]
653 __hostdev__ inline uint32_t CountOn(uint64_t v){return util::countOn(v);}
654 
655 } // namespace nanovdb ===================================================================
656 
657 #endif // end of NANOVDB_UTIL_UTIL_H_HAS_BEEN_INCLUDED
Definition: Util.h:343
C++11 implementation of std::enable_if.
Definition: Util.h:335
Trait used to identify template parameter that are pointers.
Definition: Util.h:361
T type
Definition: Util.h:343
Definition: Util.h:351
static int64_t PtrDiff(const void *p, const void *q)
Compute the distance, in bytes, between two pointers, dist = p - q.
Definition: Util.h:464
Trait use to const from type. Default implementation is just a pass-through.
Definition: Util.h:387
char * strcpy(char *dst, const char *src)
Copy characters from src to dst.
Definition: Util.h:166
Trait use to remove reference, i.e. "&", qualifier from a type. Default implementation is just a pass...
Definition: Util.h:401
char * strncpy(char *dst, const char *src, size_t max)
Copies the first num characters of src to dst. If the end of the source C string (which is signaled b...
Definition: Util.h:185
uint32_t FindLowestOn(uint32_t v)
Definition: Util.h:645
char * strcat(char *dst, const char *src)
Appends a copy of the character string pointed to by src to the end of the character string pointed t...
Definition: Util.h:225
#define __hostdev__
Definition: Util.h:73
char * sprint(char *dst)
Definition: Util.h:275
const std::enable_if<!VecTraits< T >::IsVec, T >::type & max(const T &a, const T &b)
Definition: Composite.h:110
Metafunction used to determine if the first template parameter is a specialization of the class templ...
Definition: Util.h:451
Definition: GridHandle.h:27
C++11 implementation of std::conditional.
Definition: Util.h:373
size_t strlen(const char *str)
length of a c-sting, excluding &#39;\0&#39;.
Definition: Util.h:153
T type
Definition: Util.h:401
uint32_t countOn(uint64_t v)
Definition: Util.h:622
T && declval() noexcept
Minimal implementation of std::declval, which converts any type T to.
int strcmp(const char *lhs, const char *rhs)
Compares two null-terminated byte strings lexicographically.
Definition: Util.h:255
uint32_t CountOn(uint64_t v)
Definition: Util.h:653
Trait use to remove pointer, i.e. "*", qualifier from a type. Default implementation is just a pass-t...
Definition: Util.h:415
uint32_t FindHighestOn(uint32_t v)
Definition: Util.h:649
bool empty(const char *str)
tests if a c-string str is empty, that is its first value is &#39;\0&#39;
Definition: Util.h:144
#define NANOVDB_HOSTDEV_DISABLE_WARNING
Definition: Util.h:94
uint32_t findLowestOn(uint32_t v)
Returns the index of the lowest, i.e. least significant, on bit in the specified 32 bit word...
Definition: Util.h:502
bool streq(const char *lhs, const char *rhs)
Test if two null-terminated byte strings are the same.
Definition: Util.h:268
T type
Definition: Util.h:415
#define NANOVDB_ASSERT(x)
Definition: Util.h:50
C++11 implementation of std::is_floating_point.
Definition: Util.h:329
static void * memzero(void *dst, size_t byteCount)
Zero initialization of memory.
Definition: Util.h:297
static DstT * PtrAdd(void *p, int64_t offset)
Adds a byte offset to a non-const pointer to produce another non-const pointer.
Definition: Util.h:478
C++11 implementation of std::is_same.
Definition: Util.h:314
T type
Definition: Util.h:387
T type
Definition: Util.h:338
Trait used to transfer the const-ness of a reference type to another type.
Definition: Util.h:431
uint32_t findHighestOn(uint32_t v)
Returns the index of the highest, i.e. most significant, on bit in the specified 32 bit word...
Definition: Util.h:572
T type
Definition: Util.h:421
TrueT type
Definition: Util.h:373
typename remove_const< T >::type type
Definition: Util.h:431
const typename remove_const< T >::type type
Definition: Util.h:439