Add clang tidy

.clang-format

@@ -6,15 +6,21 @@ Language: Cpp
 AccessModifierOffset: -1
 AlignAfterOpenBracket: Align
 AlignConsecutiveAssignments: true
+AlignConsecutiveBitFields: true
 AlignConsecutiveDeclarations: false
+AlignConsecutiveMacros: true
 AlignEscapedNewlines: Left
 AlignOperands: true
 AlignTrailingComments: true
+AllowAllArgumentsOnNextLine: true
+AllowAllConstructorInitializersOnNextLine: true
 AllowAllParametersOfDeclarationOnNextLine: true
 AllowShortBlocksOnASingleLine: true 
 AllowShortCaseLabelsOnASingleLine: true
+AllowShortEnumsOnASingleLine: true
 AllowShortFunctionsOnASingleLine: All
 AllowShortIfStatementsOnASingleLine: true
+AllowShortLambdasOnASingleLine: true
 AllowShortLoopsOnASingleLine: false
 # This is deprecated
 AlwaysBreakAfterDefinitionReturnType: None
@@ -40,14 +46,14 @@ BraceWrapping:
   SplitEmptyFunction: false
   SplitEmptyRecord: false
   SplitEmptyNamespace: false
+BreakAfterJavaFieldAnnotations: false
 BreakBeforeBinaryOperators: None
 BreakBeforeBraces: WebKit
 BreakBeforeInheritanceComma: false
-BreakInheritanceList: BeforeColon
 BreakBeforeTernaryOperators: true
 BreakConstructorInitializersBeforeComma: false
 BreakConstructorInitializers: BeforeColon
-BreakAfterJavaFieldAnnotations: false
+BreakInheritanceList: BeforeColon
 BreakStringLiterals: true
 ColumnLimit: 100
 CommentPragmas: '^ IWYU pragma:'
@@ -57,7 +63,7 @@ ConstructorInitializerAllOnOneLineOrOnePerLine: true
 ConstructorInitializerIndentWidth: 2
 ContinuationIndentWidth: 2
 Cpp11BracedListStyle: true
-DerivePointerAlignment: true
+DerivePointerAlignment: false
 DisableFormat: false
 ExperimentalAutoDetectBinPacking: false
 FixNamespaceComments: true
@@ -66,15 +72,6 @@ ForEachMacros:
   - Q_FOREACH
   - BOOST_FOREACH
 IncludeBlocks: Preserve
-IncludeCategories:
-  - Regex:           '^<ext/.*\.h>'
-    Priority:        2
-  - Regex:           '^<.*\.h>'
-    Priority:        1
-  - Regex:           '^<.*'
-    Priority:        2
-  - Regex:           '.*'
-    Priority:        3
 IncludeIsMainRegex: '([-_](test|unittest))?$'
 IndentCaseLabels: true
 IndentPPDirectives: None
@@ -139,14 +136,17 @@ SpaceBeforeCtorInitializerColon: true
 SpaceBeforeInheritanceColon: true
 SpaceBeforeParens: ControlStatements
 SpaceBeforeRangeBasedForLoopColon: true
+SpaceBeforeSquareBrackets: false
+SpaceInEmptyBlock: false
 SpaceInEmptyParentheses: false
 SpacesBeforeTrailingComments: 2
 SpacesInAngles: false
+SpacesInConditionalStatement: false
 SpacesInContainerLiterals: true
 SpacesInCStyleCastParentheses: false
 SpacesInParentheses: false
 SpacesInSquareBrackets: false
-Standard: Cpp11
+Standard: c++17
 StatementMacros:
   - Q_UNUSED
   - QT_REQUIRE_VERSION

.clang-tidy

@@ -0,0 +1,68 @@
+---
+Checks:          'clang-diagnostic-*,
+                  clang-analyzer-*,
+                  cppcoreguidelines-*,
+                  modernize-*,
+                  bugprone-*,
+                  performance-*,
+                  readability-*,
+                  llvm-*,
+                  -cppcoreguidelines-macro-usage,
+                  -llvm-header-guard,
+                  -modernize-use-trailing-return-type,
+                  -readability-static-accessed-through-instance,
+                  -readability-named-parameter'
+WarningsAsErrors: ''
+HeaderFilterRegex: ''
+AnalyzeTemporaryDtors: false
+FormatStyle:     none
+CheckOptions:
+  - key:             cert-dcl16-c.NewSuffixes
+    value:           'L;LL;LU;LLU'
+  - key:             cert-oop54-cpp.WarnOnlyIfThisHasSuspiciousField
+    value:           '0'
+  - key:             cert-str34-c.DiagnoseSignedUnsignedCharComparisons
+    value:           '0'
+  - key:             cppcoreguidelines-explicit-virtual-functions.IgnoreDestructors
+    value:           '1'
+  - key:             cppcoreguidelines-non-private-member-variables-in-classes.IgnoreClassesWithAllMemberVariablesBeingPublic
+    value:           '1'
+  - key:             google-readability-braces-around-statements.ShortStatementLines
+    value:           '1'
+  - key:             google-readability-function-size.StatementThreshold
+    value:           '800'
+  - key:             google-readability-namespace-comments.ShortNamespaceLines
+    value:           '10'
+  - key:             google-readability-namespace-comments.SpacesBeforeComments
+    value:           '2'
+  - key:             llvm-else-after-return.WarnOnConditionVariables
+    value:           '0'
+  - key:             llvm-else-after-return.WarnOnUnfixable
+    value:           '0'
+  - key:             llvm-qualified-auto.AddConstToQualified
+    value:           '0'
+  - key:             modernize-loop-convert.MaxCopySize
+    value:           '16'
+  - key:             modernize-loop-convert.MinConfidence
+    value:           reasonable
+  - key:             modernize-loop-convert.NamingStyle
+    value:           CamelCase
+  - key:             modernize-pass-by-value.IncludeStyle
+    value:           llvm
+  - key:             modernize-replace-auto-ptr.IncludeStyle
+    value:           llvm
+  - key:             modernize-use-nullptr.NullMacros
+    value:           'NULL'
+  - key:             readability-identifier-length.IgnoredParameterNames
+    value:           'mr|os'
+  - key:             readability-identifier-length.IgnoredVariableNames
+    value:           'mr|_'
+  #- key:             readability-function-cognitive-complexity.IgnoreMacros
+  #  value:           '1'
+  - key:             bugprone-easily-swappable-parameters.IgnoredParameterNames
+    value:           'alignment'
+  - key:             cppcoreguidelines-avoid-magic-numbers.IgnorePowersOf2IntegerValues
+    value:           '1'
+  - key:             readability-magic-numbers.IgnorePowersOf2IntegerValues
+    value:           '1'
+...

CMakeLists.txt

@@ -29,6 +29,9 @@ include(rapids-find)
 # * Enable the CMake CUDA language
 rapids_cuda_init_architectures(CUCO)
 
+# this is needed for clang-tidy runs
+set(CMAKE_EXPORT_COMPILE_COMMANDS ON)
+
 project(CUCO VERSION 0.0.1 LANGUAGES CXX CUDA)
 
 # Write the version header

benchmarks/hash_table/static_map_bench.cu

@@ -14,13 +14,13 @@
  * limitations under the License.
  */
 
+#include "cuco/static_map.cuh"
 #include <benchmark/benchmark.h>
-#include <thrust/device_vector.h>
-#include <thrust/for_each.h>
 #include <fstream>
 #include <iostream>
 #include <random>
-#include "cuco/static_map.cuh"
+#include <thrust/device_vector.h>
+#include <thrust/for_each.h>
 
 enum class dist_type { UNIQUE, UNIFORM, GAUSSIAN };
 

benchmarks/hash_table/static_multimap/retrieve_bench.cu

@@ -16,8 +16,8 @@
 
 #include <nvbench/nvbench.cuh>
 
-#include <thrust/device_vector.h>
 #include <random>
+#include <thrust/device_vector.h>
 
 #include "cuco/static_multimap.cuh"
 

benchmarks/hash_table/static_multimap/static_multimap_bench.cu

@@ -16,9 +16,9 @@
 
 #include <random>
 
+#include <nvbench/nvbench.cuh>
 #include <thrust/device_vector.h>
 #include <thrust/iterator/discard_iterator.h>
-#include <nvbench/nvbench.cuh>
 
 #include <cuco/static_multimap.cuh>
 #include <key_generator.hpp>

benchmarks/reduce_by_key/reduce_by_key.cu

@@ -16,13 +16,13 @@
 
 #include <benchmark/benchmark.h>
 
+#include <thrust/device_vector.h>
 #include <thrust/iterator/iterator_traits.h>
 #include <thrust/iterator/zip_iterator.h>
 #include <thrust/pair.h>
 #include <thrust/reduce.h>
 #include <thrust/sort.h>
 #include <thrust/transform.h>
-#include <thrust/device_vector.h>
 
 /**
  * @brief Generates input sizes and number of unique keys
@@ -76,13 +76,11 @@ static void BM_thrust(::benchmark::State& state)
   }
 }
 BENCHMARK_TEMPLATE(BM_thrust, int32_t, int32_t)
-    ->Unit(benchmark::kMillisecond)
-    ->Apply(generate_size_and_num_unique);
+  ->Unit(benchmark::kMillisecond)
+  ->Apply(generate_size_and_num_unique);
 
 BENCHMARK_TEMPLATE(BM_thrust, int64_t, int64_t)
-    ->Unit(benchmark::kMillisecond)
-    ->Apply(generate_size_and_num_unique);
+  ->Unit(benchmark::kMillisecond)
+  ->Apply(generate_size_and_num_unique);
 
 // TODO: Hash based reduce by key benchmark
-
-

benchmarks/synchronization.hpp

@@ -81,7 +81,9 @@ class cuda_event_timer {
    *                            every iteration.
    * @param[in] stream_ The CUDA stream we are measuring time on.
    */
-  cuda_event_timer(benchmark::State &state, bool flush_l2_cache = false, cudaStream_t stream = 0)
+  cuda_event_timer(benchmark::State& state,
+                   bool flush_l2_cache = false,
+                   cudaStream_t stream = nullptr)
     : p_state(&state), stream_(stream)
   {
     // flush all of L2$
@@ -95,7 +97,7 @@ class cuda_event_timer {
 
       if (l2_cache_bytes > 0) {
         const int memset_value = 0;
-        int *l2_cache_buffer   = nullptr;
+        int* l2_cache_buffer   = nullptr;
         BENCH_CUDA_TRY(cudaMalloc(&l2_cache_buffer, l2_cache_bytes));
         BENCH_CUDA_TRY(cudaMemsetAsync(l2_cache_buffer, memset_value, l2_cache_bytes, stream_));
         BENCH_CUDA_TRY(cudaFree(l2_cache_buffer));
@@ -128,5 +130,5 @@ class cuda_event_timer {
   cudaEvent_t start_;
   cudaEvent_t stop_;
   cudaStream_t stream_;
-  benchmark::State *p_state;
-};
+  benchmark::State* p_state;
+};

ci/checks/style.sh

@@ -0,0 +1,24 @@
+#!/bin/bash
+# Copyright (c) 2018-2021, NVIDIA CORPORATION.
+
+# Ignore errors and set path
+set +e
+PATH=/conda/bin:$PATH
+LC_ALL=C.UTF-8
+LANG=C.UTF-8
+
+# Activate common conda env
+. /opt/conda/etc/profile.d/conda.sh
+conda activate rapids
+
+# Run clang-format and check for a consistent code format
+CLANG_FORMAT=`python scripts/run-clang-format.py 2>&1`
+CLANG_FORMAT_RETVAL=$?
+
+if [ "$CLANG_FORMAT_RETVAL" != "0" ]; then
+  echo -e "\n\n>>>> FAILED: clang format check; begin output\n\n"
+  echo -e "$CLANG_FORMAT"
+  echo -e "\n\n>>>> FAILED: clang format check; end output\n\n"
+else
+  echo -e "\n\n>>>> PASSED: clang format check\n\n"
+fi

include/cuco/detail/error.hpp

@@ -32,7 +32,7 @@ struct cuda_error : public std::runtime_error {
 }  // namespace cuco
 
 #define STRINGIFY_DETAIL(x) #x
-#define CUCO_STRINGIFY(x) STRINGIFY_DETAIL(x)
+#define CUCO_STRINGIFY(x)   STRINGIFY_DETAIL(x)
 
 /**
  * @brief Error checking macro for CUDA runtime API functions.

include/cuco/detail/hash_functions.cuh

@@ -81,12 +81,12 @@ struct MurmurHash3_32 {
   }
 
  private:
-  constexpr __host__ __device__ uint32_t rotl32(uint32_t x, int8_t r) const noexcept
+  [[nodiscard]] constexpr __host__ __device__ uint32_t rotl32(uint32_t x, int8_t r) const noexcept
   {
     return (x << r) | (x >> (32 - r));
   }
 
-  constexpr __host__ __device__ uint32_t fmix32(uint32_t h) const noexcept
+  [[nodiscard]] constexpr __host__ __device__ uint32_t fmix32(uint32_t h) const noexcept
   {
     h ^= h >> 16;
     h *= 0x85ebca6b;

include/cuco/detail/probe_sequence_impl.cuh

@@ -112,7 +112,7 @@ class probe_sequence_impl_base {
   /**
    * @brief Returns the capacity of the hash map.
    */
-  __host__ __device__ __forceinline__ std::size_t get_capacity() const noexcept
+  [[nodiscard]] __host__ __device__ __forceinline__ std::size_t get_capacity() const noexcept
   {
     return capacity_;
   }
@@ -125,7 +125,10 @@ class probe_sequence_impl_base {
   /**
    * @brief Returns slots array.
    */
-  __device__ __forceinline__ const_iterator get_slots() const noexcept { return slots_; }
+  [[nodiscard]] __device__ __forceinline__ const_iterator get_slots() const noexcept
+  {
+    return slots_;
+  }
 
  protected:
   iterator slots_;              ///< Pointer to beginning of the hash map slots

include/cuco/detail/static_map.inl

@@ -383,10 +383,11 @@ static_map<Key, Value, Scope, Allocator>::device_view::find(Key const& k,
 
 template <typename Key, typename Value, cuda::thread_scope Scope, typename Allocator>
 template <typename Hash, typename KeyEqual>
-__device__ typename static_map<Key, Value, Scope, Allocator>::device_view::const_iterator
-static_map<Key, Value, Scope, Allocator>::device_view::find(Key const& k,
-                                                            Hash hash,
-                                                            KeyEqual key_equal) const noexcept
+[[nodiscard]] __device__
+  typename static_map<Key, Value, Scope, Allocator>::device_view::const_iterator
+  static_map<Key, Value, Scope, Allocator>::device_view::find(Key const& k,
+                                                              Hash hash,
+                                                              KeyEqual key_equal) const noexcept
 {
   auto current_slot = initial_slot(k, hash);
 
@@ -502,7 +503,7 @@ __device__ bool static_map<Key, Value, Scope, Allocator>::device_view::contains(
 
 template <typename Key, typename Value, cuda::thread_scope Scope, typename Allocator>
 template <typename CG, typename Hash, typename KeyEqual>
-__device__ bool static_map<Key, Value, Scope, Allocator>::device_view::contains(
+[[nodiscard]] __device__ bool static_map<Key, Value, Scope, Allocator>::device_view::contains(
   CG g, Key const& k, Hash hash, KeyEqual key_equal) const noexcept
 {
   auto current_slot = initial_slot(g, k, hash);

include/cuco/detail/static_multimap/device_view_impl.inl

@@ -145,7 +145,7 @@ class static_multimap<Key, Value, Scope, Allocator, ProbeSequence>::device_view_
    *
    * @return The sentinel value used to represent an empty key slot
    */
-  __host__ __device__ __forceinline__ Key get_empty_key_sentinel() const noexcept
+  [[nodiscard]] __host__ __device__ __forceinline__ Key get_empty_key_sentinel() const noexcept
   {
     return empty_key_sentinel_;
   }
@@ -155,7 +155,7 @@ class static_multimap<Key, Value, Scope, Allocator, ProbeSequence>::device_view_
    *
    * @return The sentinel value used to represent an empty value slot
    */
-  __host__ __device__ __forceinline__ Value get_empty_value_sentinel() const noexcept
+  [[nodiscard]] __host__ __device__ __forceinline__ Value get_empty_value_sentinel() const noexcept
   {
     return empty_value_sentinel_;
   }
@@ -175,7 +175,7 @@ class static_multimap<Key, Value, Scope, Allocator, ProbeSequence>::device_view_
    *
    * @return Slots array
    */
-  __device__ __forceinline__ pair_atomic_type const* get_slots() const noexcept
+  [[nodiscard]] __device__ __forceinline__ pair_atomic_type const* get_slots() const noexcept
   {
     return probe_sequence_.get_slots();
   }
@@ -185,7 +185,7 @@ class static_multimap<Key, Value, Scope, Allocator, ProbeSequence>::device_view_
    *
    * @return The maximum number of elements the hash map can hold
    */
-  __host__ __device__ __forceinline__ std::size_t get_capacity() const noexcept
+  [[nodiscard]] __host__ __device__ __forceinline__ std::size_t get_capacity() const noexcept
   {
     return probe_sequence_.get_capacity();
   }