Replace typedef by alias declarations

Using:
* clang-tidy modernize-use-using
* regex replace: "typedef ([\w<>,:\s*&+:\[\]\.]+)\s+([\w_]+);" by "using $2 = $1;"
* manual search and edits

Fixes: NVIDIA#1747

cub/README.md

@@ -34,9 +34,9 @@ __global__ void BlockSortKernel(int *d_in, int *d_out)
 
      // Specialize BlockRadixSort, BlockLoad, and BlockStore for 128 threads
      // owning 16 integer items each
-     typedef BlockRadixSort<int, 128, 16>                     BlockRadixSort;
-     typedef BlockLoad<int, 128, 16, BLOCK_LOAD_TRANSPOSE>   BlockLoad;
-     typedef BlockStore<int, 128, 16, BLOCK_STORE_TRANSPOSE> BlockStore;
+     using BlockRadixSort = BlockRadixSort<int, 128, 16>;
+     using BlockLoad = BlockLoad<int, 128, 16, BLOCK_LOAD_TRANSPOSE>;
+     using BlockStore = BlockStore<int, 128, 16, BLOCK_STORE_TRANSPOSE> BlockStore;
 
      // Allocate shared memory
      __shared__ union {

cub/cub/agent/agent_histogram.cuh

@@ -230,19 +230,19 @@ struct AgentHistogram
                                SampleIteratorT>;
 
   /// Pixel input iterator type (for applying cache modifier)
-  typedef CacheModifiedInputIterator<LOAD_MODIFIER, PixelT, OffsetT> WrappedPixelIteratorT;
+  using WrappedPixelIteratorT = CacheModifiedInputIterator<LOAD_MODIFIER, PixelT, OffsetT>;
 
   /// Qaud input iterator type (for applying cache modifier)
-  typedef CacheModifiedInputIterator<LOAD_MODIFIER, VecT, OffsetT> WrappedVecsIteratorT;
+  using WrappedVecsIteratorT = CacheModifiedInputIterator<LOAD_MODIFIER, VecT, OffsetT>;
 
   /// Parameterized BlockLoad type for samples
-  typedef BlockLoad<SampleT, BLOCK_THREADS, SAMPLES_PER_THREAD, AgentHistogramPolicyT::LOAD_ALGORITHM> BlockLoadSampleT;
+  using BlockLoadSampleT = BlockLoad<SampleT, BLOCK_THREADS, SAMPLES_PER_THREAD, AgentHistogramPolicyT::LOAD_ALGORITHM>;
 
   /// Parameterized BlockLoad type for pixels
-  typedef BlockLoad<PixelT, BLOCK_THREADS, PIXELS_PER_THREAD, AgentHistogramPolicyT::LOAD_ALGORITHM> BlockLoadPixelT;
+  using BlockLoadPixelT = BlockLoad<PixelT, BLOCK_THREADS, PIXELS_PER_THREAD, AgentHistogramPolicyT::LOAD_ALGORITHM>;
 
   /// Parameterized BlockLoad type for vecs
-  typedef BlockLoad<VecT, BLOCK_THREADS, VECS_PER_THREAD, AgentHistogramPolicyT::LOAD_ALGORITHM> BlockLoadVecT;
+  using BlockLoadVecT = BlockLoad<VecT, BLOCK_THREADS, VECS_PER_THREAD, AgentHistogramPolicyT::LOAD_ALGORITHM>;
 
   /// Shared memory type required by this thread block
   struct _TempStorage
@@ -505,7 +505,7 @@ struct AgentHistogram
     SampleT (&samples)[PIXELS_PER_THREAD][NUM_CHANNELS],
     Int2Type<_NUM_ACTIVE_CHANNELS> num_active_channels)
   {
-    typedef PixelT AliasedPixels[PIXELS_PER_THREAD];
+    using AliasedPixels = PixelT[PIXELS_PER_THREAD];
 
     WrappedPixelIteratorT d_wrapped_pixels((PixelT*) (d_native_samples + block_offset));
 
@@ -520,7 +520,7 @@ struct AgentHistogram
     SampleT (&samples)[PIXELS_PER_THREAD][NUM_CHANNELS],
     Int2Type<1> num_active_channels)
   {
-    typedef VecT AliasedVecs[VECS_PER_THREAD];
+    using AliasedVecs = VecT[VECS_PER_THREAD];
 
     WrappedVecsIteratorT d_wrapped_vecs((VecT*) (d_native_samples + block_offset));
 
@@ -547,7 +547,7 @@ struct AgentHistogram
     Int2Type<true> is_full_tile,
     Int2Type<false> is_aligned)
   {
-    typedef SampleT AliasedSamples[SAMPLES_PER_THREAD];
+    using AliasedSamples = SampleT[SAMPLES_PER_THREAD];
 
     // Load using sample iterator
     BlockLoadSampleT(temp_storage.aliasable.sample_load)
@@ -562,7 +562,7 @@ struct AgentHistogram
     Int2Type<false> is_full_tile,
     Int2Type<true> is_aligned)
   {
-    typedef PixelT AliasedPixels[PIXELS_PER_THREAD];
+    using AliasedPixels = PixelT[PIXELS_PER_THREAD];
 
     WrappedPixelIteratorT d_wrapped_pixels((PixelT*) (d_native_samples + block_offset));
 
@@ -581,7 +581,7 @@ struct AgentHistogram
     Int2Type<false> is_full_tile,
     Int2Type<false> is_aligned)
   {
-    typedef SampleT AliasedSamples[SAMPLES_PER_THREAD];
+    using AliasedSamples = SampleT[SAMPLES_PER_THREAD];
 
     BlockLoadSampleT(temp_storage.aliasable.sample_load)
       .Load(d_wrapped_samples + block_offset, reinterpret_cast<AliasedSamples&>(samples), valid_samples);

cub/cub/agent/agent_radix_sort_downsweep.cuh

@@ -202,7 +202,7 @@ struct AgentRadixSortDownsweep
   using BlockLoadValuesT = BlockLoad<ValueT, BLOCK_THREADS, ITEMS_PER_THREAD, LOAD_ALGORITHM>;
 
   // Value exchange array type
-  typedef ValueT ValueExchangeT[TILE_ITEMS];
+  using ValueExchangeT = ValueT[TILE_ITEMS];
 
   /**
    * Shared memory storage layout

cub/cub/agent/agent_radix_sort_histogram.cuh

@@ -102,9 +102,9 @@ struct AgentRadixSortHistogram
   using bit_ordered_type       = typename traits::bit_ordered_type;
   using bit_ordered_conversion = typename traits::bit_ordered_conversion_policy;
 
-  typedef RadixSortTwiddle<IS_DESCENDING, KeyT> Twiddle;
-  typedef std::uint32_t ShmemCounterT;
-  typedef ShmemCounterT ShmemAtomicCounterT;
+  using Twiddle             = RadixSortTwiddle<IS_DESCENDING, KeyT>;
+  using ShmemCounterT       = std::uint32_t;
+  using ShmemAtomicCounterT = ShmemCounterT;
 
   using fundamental_digit_extractor_t = ShiftDigitExtractor<KeyT>;
   using digit_extractor_t = typename traits::template digit_extractor_t<fundamental_digit_extractor_t, DecomposerT>;

cub/cub/agent/agent_radix_sort_onesweep.cuh

@@ -131,7 +131,7 @@ struct AgentRadixSortOnesweep
   using fundamental_digit_extractor_t = ShiftDigitExtractor<KeyT>;
   using digit_extractor_t = typename traits::template digit_extractor_t<fundamental_digit_extractor_t, DecomposerT>;
 
-  typedef PortionOffsetT AtomicOffsetT;
+  using AtomicOffsetT = PortionOffsetT;
 
   static constexpr RadixRankAlgorithm RANK_ALGORITHM = AgentRadixSortOnesweepPolicy::RANK_ALGORITHM;
   static constexpr BlockScanAlgorithm SCAN_ALGORITHM = AgentRadixSortOnesweepPolicy::SCAN_ALGORITHM;
@@ -140,7 +140,7 @@ struct AgentRadixSortOnesweep
       ? AgentRadixSortOnesweepPolicy::STORE_ALGORITHM
       : RADIX_SORT_STORE_DIRECT;
 
-  typedef RadixSortTwiddle<IS_DESCENDING, KeyT> Twiddle;
+  using Twiddle = RadixSortTwiddle<IS_DESCENDING, KeyT>;
 
   static_assert(RANK_ALGORITHM == RADIX_RANK_MATCH || RANK_ALGORITHM == RADIX_RANK_MATCH_EARLY_COUNTS_ANY
                   || RANK_ALGORITHM == RADIX_RANK_MATCH_EARLY_COUNTS_ATOMIC_OR,
@@ -229,14 +229,8 @@ struct AgentRadixSortOnesweep
 
   struct CountsCallback
   {
-    typedef AgentRadixSortOnesweep<AgentRadixSortOnesweepPolicy,
-                                   IS_DESCENDING,
-                                   KeyT,
-                                   ValueT,
-                                   OffsetT,
-                                   PortionOffsetT,
-                                   DecomposerT>
-      AgentT;
+    using AgentT =
+      AgentRadixSortOnesweep<AgentRadixSortOnesweepPolicy, IS_DESCENDING, KeyT, ValueT, OffsetT, PortionOffsetT, DecomposerT>;
     AgentT& agent;
     int (&bins)[BINS_PER_THREAD];
     bit_ordered_type (&keys)[ITEMS_PER_THREAD];

cub/cub/agent/agent_radix_sort_upsweep.cuh

@@ -125,10 +125,10 @@ struct AgentRadixSortUpsweep
   using bit_ordered_conversion = typename traits::bit_ordered_conversion_policy;
 
   // Integer type for digit counters (to be packed into words of PackedCounters)
-  typedef unsigned char DigitCounter;
+  using DigitCounter = unsigned char;
 
   // Integer type for packing DigitCounters into columns of shared memory banks
-  typedef unsigned int PackedCounter;
+  using PackedCounter = unsigned int;
 
   static constexpr CacheLoadModifier LOAD_MODIFIER = AgentRadixSortUpsweepPolicy::LOAD_MODIFIER;
 
@@ -167,7 +167,7 @@ struct AgentRadixSortUpsweep
   };
 
   // Input iterator wrapper type (for applying cache modifier)s
-  typedef CacheModifiedInputIterator<LOAD_MODIFIER, bit_ordered_type, OffsetT> KeysItr;
+  using KeysItr = CacheModifiedInputIterator<LOAD_MODIFIER, bit_ordered_type, OffsetT>;
 
   // Digit extractor type
   using fundamental_digit_extractor_t = BFEDigitExtractor<KeyT>;

cub/cub/agent/agent_reduce_by_key.cuh

@@ -269,8 +269,8 @@ struct AgentReduceByKey
     TilePrefixCallbackOp<OffsetValuePairT, ReduceBySegmentOpT, ScanTileStateT, 0, DelayConstructorT>;
 
   // Key and value exchange types
-  typedef KeyOutputT KeyExchangeT[TILE_ITEMS + 1];
-  typedef AccumT ValueExchangeT[TILE_ITEMS + 1];
+  using KeyExchangeT   = KeyOutputT[TILE_ITEMS + 1];
+  using ValueExchangeT = AccumT[TILE_ITEMS + 1];
 
   // Shared memory type for this thread block
   union _TempStorage

cub/cub/agent/agent_rle.cuh

@@ -262,7 +262,7 @@ struct AgentRle
   using WarpExchangeOffsets = WarpExchange<OffsetT, ITEMS_PER_THREAD>;
   using WarpExchangeLengths = WarpExchange<LengthT, ITEMS_PER_THREAD>;
 
-  typedef LengthOffsetPair WarpAggregates[WARPS];
+  using WarpAggregates = LengthOffsetPair[WARPS];
 
   // Shared memory type for this thread block
   struct _TempStorage

cub/cub/agent/agent_scan.cuh

@@ -172,29 +172,29 @@ struct AgentScan
   };
 
   // Parameterized BlockLoad type
-  typedef BlockLoad<AccumT,
-                    AgentScanPolicyT::BLOCK_THREADS,
-                    AgentScanPolicyT::ITEMS_PER_THREAD,
-                    AgentScanPolicyT::LOAD_ALGORITHM>
-    BlockLoadT;
+  using BlockLoadT =
+    BlockLoad<AccumT,
+              AgentScanPolicyT::BLOCK_THREADS,
+              AgentScanPolicyT::ITEMS_PER_THREAD,
+              AgentScanPolicyT::LOAD_ALGORITHM>;
 
   // Parameterized BlockStore type
-  typedef BlockStore<AccumT,
-                     AgentScanPolicyT::BLOCK_THREADS,
-                     AgentScanPolicyT::ITEMS_PER_THREAD,
-                     AgentScanPolicyT::STORE_ALGORITHM>
-    BlockStoreT;
+  using BlockStoreT =
+    BlockStore<AccumT,
+               AgentScanPolicyT::BLOCK_THREADS,
+               AgentScanPolicyT::ITEMS_PER_THREAD,
+               AgentScanPolicyT::STORE_ALGORITHM>;
 
   // Parameterized BlockScan type
-  typedef BlockScan<AccumT, AgentScanPolicyT::BLOCK_THREADS, AgentScanPolicyT::SCAN_ALGORITHM> BlockScanT;
+  using BlockScanT = BlockScan<AccumT, AgentScanPolicyT::BLOCK_THREADS, AgentScanPolicyT::SCAN_ALGORITHM>;
 
   // Callback type for obtaining tile prefix during block scan
   using DelayConstructorT     = typename AgentScanPolicyT::detail::delay_constructor_t;
   using TilePrefixCallbackOpT = TilePrefixCallbackOp<AccumT, ScanOpT, ScanTileStateT, 0 /* PTX */, DelayConstructorT>;
 
   // Stateful BlockScan prefix callback type for managing a running total while
   // scanning consecutive tiles
-  typedef BlockScanRunningPrefixOp<AccumT, ScanOpT> RunningPrefixCallbackOp;
+  using RunningPrefixCallbackOp = BlockScanRunningPrefixOp<AccumT, ScanOpT>;
 
   // Shared memory type for this thread block
   union _TempStorage

cub/cub/agent/agent_select_if.cuh

@@ -248,7 +248,7 @@ struct AgentSelectIf
   using TilePrefixCallbackOpT = TilePrefixCallbackOp<OffsetT, cub::Sum, ScanTileStateT, 0, DelayConstructorT>;
 
   // Item exchange type
-  typedef InputT ItemExchangeT[TILE_ITEMS];
+  using ItemExchangeT = InputT[TILE_ITEMS];
 
   // Shared memory type for this thread block
   union _TempStorage

cub/cub/agent/agent_spmv_orig.cuh

@@ -224,40 +224,40 @@ struct AgentSpmv
   };
 
   /// 2D merge path coordinate type
-  typedef typename CubVector<OffsetT, 2>::Type CoordinateT;
+  using CoordinateT = typename CubVector<OffsetT, 2>::Type;
 
   /// Input iterator wrapper types (for applying cache modifiers)
 
-  typedef CacheModifiedInputIterator<AgentSpmvPolicyT::ROW_OFFSETS_SEARCH_LOAD_MODIFIER, OffsetT, OffsetT>
-    RowOffsetsSearchIteratorT;
+  using RowOffsetsSearchIteratorT =
+    CacheModifiedInputIterator<AgentSpmvPolicyT::ROW_OFFSETS_SEARCH_LOAD_MODIFIER, OffsetT, OffsetT>;
 
-  typedef CacheModifiedInputIterator<AgentSpmvPolicyT::ROW_OFFSETS_LOAD_MODIFIER, OffsetT, OffsetT> RowOffsetsIteratorT;
+  using RowOffsetsIteratorT = CacheModifiedInputIterator<AgentSpmvPolicyT::ROW_OFFSETS_LOAD_MODIFIER, OffsetT, OffsetT>;
 
-  typedef CacheModifiedInputIterator<AgentSpmvPolicyT::COLUMN_INDICES_LOAD_MODIFIER, OffsetT, OffsetT>
-    ColumnIndicesIteratorT;
+  using ColumnIndicesIteratorT =
+    CacheModifiedInputIterator<AgentSpmvPolicyT::COLUMN_INDICES_LOAD_MODIFIER, OffsetT, OffsetT>;
 
-  typedef CacheModifiedInputIterator<AgentSpmvPolicyT::VALUES_LOAD_MODIFIER, ValueT, OffsetT> ValueIteratorT;
+  using ValueIteratorT = CacheModifiedInputIterator<AgentSpmvPolicyT::VALUES_LOAD_MODIFIER, ValueT, OffsetT>;
 
-  typedef CacheModifiedInputIterator<AgentSpmvPolicyT::VECTOR_VALUES_LOAD_MODIFIER, ValueT, OffsetT>
-    VectorValueIteratorT;
+  using VectorValueIteratorT =
+    CacheModifiedInputIterator<AgentSpmvPolicyT::VECTOR_VALUES_LOAD_MODIFIER, ValueT, OffsetT>;
 
   // Tuple type for scanning (pairs accumulated segment-value with segment-index)
-  typedef KeyValuePair<OffsetT, ValueT> KeyValuePairT;
+  using KeyValuePairT = KeyValuePair<OffsetT, ValueT>;
 
   // Reduce-value-by-segment scan operator
-  typedef ReduceByKeyOp<cub::Sum> ReduceBySegmentOpT;
+  using ReduceBySegmentOpT = ReduceByKeyOp<cub::Sum>;
 
   // BlockReduce specialization
-  typedef BlockReduce<ValueT, BLOCK_THREADS, BLOCK_REDUCE_WARP_REDUCTIONS> BlockReduceT;
+  using BlockReduceT = BlockReduce<ValueT, BLOCK_THREADS, BLOCK_REDUCE_WARP_REDUCTIONS>;
 
   // BlockScan specialization
-  typedef BlockScan<KeyValuePairT, BLOCK_THREADS, AgentSpmvPolicyT::SCAN_ALGORITHM> BlockScanT;
+  using BlockScanT = BlockScan<KeyValuePairT, BLOCK_THREADS, AgentSpmvPolicyT::SCAN_ALGORITHM>;
 
   // BlockScan specialization
-  typedef BlockScan<ValueT, BLOCK_THREADS, AgentSpmvPolicyT::SCAN_ALGORITHM> BlockPrefixSumT;
+  using BlockPrefixSumT = BlockScan<ValueT, BLOCK_THREADS, AgentSpmvPolicyT::SCAN_ALGORITHM>;
 
   // BlockExchange specialization
-  typedef BlockExchange<ValueT, BLOCK_THREADS, ITEMS_PER_THREAD> BlockExchangeT;
+  using BlockExchangeT = BlockExchange<ValueT, BLOCK_THREADS, ITEMS_PER_THREAD>;
 
   /// Merge item type (either a non-zero value or a row-end offset)
   union MergeItem

cub/cub/agent/single_pass_scan_operators.cuh

@@ -861,7 +861,7 @@ struct ReduceByKeyScanTileState;
 template <typename ValueT, typename KeyT>
 struct ReduceByKeyScanTileState<ValueT, KeyT, false> : ScanTileState<KeyValuePair<KeyT, ValueT>>
 {
-  typedef ScanTileState<KeyValuePair<KeyT, ValueT>> SuperClass;
+  using SuperClass = ScanTileState<KeyValuePair<KeyT, ValueT>>;
 
   /// Constructor
   _CCCL_HOST_DEVICE _CCCL_FORCEINLINE ReduceByKeyScanTileState()
@@ -1082,7 +1082,7 @@ template <typename T,
 struct TilePrefixCallbackOp
 {
   // Parameterized warp reduce
-  typedef WarpReduce<T, CUB_PTX_WARP_THREADS> WarpReduceT;
+  using WarpReduceT = WarpReduce<T, (1 << (5))>;
 
   // Temporary storage type
   struct _TempStorage
@@ -1098,7 +1098,7 @@ struct TilePrefixCallbackOp
   {};
 
   // Type of status word
-  typedef typename ScanTileStateT::StatusWord StatusWord;
+  using StatusWord = typename ScanTileStateT::StatusWord;
 
   // Fields
   _TempStorage& temp_storage; ///< Reference to a warp-reduction instance

cub/cub/block/block_discontinuity.cuh

@@ -84,7 +84,7 @@ CUB_NAMESPACE_BEGIN
 //!    __global__ void ExampleKernel(...)
 //!    {
 //!        // Specialize BlockDiscontinuity for a 1D block of 128 threads of type int
-//!        typedef cub::BlockDiscontinuity<int, 128> BlockDiscontinuity;
+//!        using BlockDiscontinuity = cub::BlockDiscontinuity<int, 128>;
 //!
 //!        // Allocate shared memory for BlockDiscontinuity
 //!        __shared__ typename BlockDiscontinuity::TempStorage temp_storage;
@@ -377,7 +377,7 @@ public:
   //!    __global__ void ExampleKernel(...)
   //!    {
   //!        // Specialize BlockDiscontinuity for a 1D block of 128 threads of type int
-  //!        typedef cub::BlockDiscontinuity<int, 128> BlockDiscontinuity;
+  //!        using BlockDiscontinuity = cub::BlockDiscontinuity<int, 128>;
   //!
   //!        // Allocate shared memory for BlockDiscontinuity
   //!        __shared__ typename BlockDiscontinuity::TempStorage temp_storage;
@@ -450,7 +450,7 @@ public:
   //!    __global__ void ExampleKernel(...)
   //!    {
   //!        // Specialize BlockDiscontinuity for a 1D block of 128 threads of type int
-  //!        typedef cub::BlockDiscontinuity<int, 128> BlockDiscontinuity;
+  //!        using BlockDiscontinuity = cub::BlockDiscontinuity<int, 128>;
   //!
   //!        // Allocate shared memory for BlockDiscontinuity
   //!        __shared__ typename BlockDiscontinuity::TempStorage temp_storage;
@@ -539,7 +539,7 @@ public:
   //!    __global__ void ExampleKernel(...)
   //!    {
   //!        // Specialize BlockDiscontinuity for a 1D block of 128 threads of type int
-  //!        typedef cub::BlockDiscontinuity<int, 128> BlockDiscontinuity;
+  //!        using BlockDiscontinuity = cub::BlockDiscontinuity<int, 128>;
   //!
   //!        // Allocate shared memory for BlockDiscontinuity
   //!        __shared__ typename BlockDiscontinuity::TempStorage temp_storage;
@@ -627,7 +627,7 @@ public:
   //!    __global__ void ExampleKernel(...)
   //!    {
   //!        // Specialize BlockDiscontinuity for a 1D block of 128 threads of type int
-  //!        typedef cub::BlockDiscontinuity<int, 128> BlockDiscontinuity;
+  //!        using BlockDiscontinuity = cub::BlockDiscontinuity<int, 128>;
   //!
   //!        // Allocate shared memory for BlockDiscontinuity
   //!        __shared__ typename BlockDiscontinuity::TempStorage temp_storage;
@@ -732,7 +732,7 @@ public:
   //!    __global__ void ExampleKernel(...)
   //!    {
   //!        // Specialize BlockDiscontinuity for a 1D block of 128 threads of type int
-  //!        typedef cub::BlockDiscontinuity<int, 128> BlockDiscontinuity;
+  //!        using BlockDiscontinuity = cub::BlockDiscontinuity<int, 128>;
   //!
   //!        // Allocate shared memory for BlockDiscontinuity
   //!        __shared__ typename BlockDiscontinuity::TempStorage temp_storage;
@@ -850,7 +850,7 @@ public:
   //!    __global__ void ExampleKernel(...)
   //!    {
   //!        // Specialize BlockDiscontinuity for a 1D block of 128 threads of type int
-  //!        typedef cub::BlockDiscontinuity<int, 128> BlockDiscontinuity;
+  //!        using BlockDiscontinuity = cub::BlockDiscontinuity<int, 128>;
   //!
   //!        // Allocate shared memory for BlockDiscontinuity
   //!        __shared__ typename BlockDiscontinuity::TempStorage temp_storage;
@@ -979,7 +979,7 @@ public:
   //!    __global__ void ExampleKernel(...)
   //!    {
   //!        // Specialize BlockDiscontinuity for a 1D block of 128 threads of type int
-  //!        typedef cub::BlockDiscontinuity<int, 128> BlockDiscontinuity;
+  //!        using BlockDiscontinuity = cub::BlockDiscontinuity<int, 128>;
   //!
   //!        // Allocate shared memory for BlockDiscontinuity
   //!        __shared__ typename BlockDiscontinuity::TempStorage temp_storage;
@@ -1108,7 +1108,7 @@ public:
   //!    __global__ void ExampleKernel(...)
   //!    {
   //!        // Specialize BlockDiscontinuity for a 1D block of 128 threads of type int
-  //!        typedef cub::BlockDiscontinuity<int, 128> BlockDiscontinuity;
+  //!        using BlockDiscontinuity = cub::BlockDiscontinuity<int, 128>;
   //!
   //!        // Allocate shared memory for BlockDiscontinuity
   //!        __shared__ typename BlockDiscontinuity::TempStorage temp_storage;