diff --git a/cub/cub/warp/warp_reduce.cuh b/cub/cub/warp/warp_reduce.cuh index 34a46a26fa..2901f56bbc 100644 --- a/cub/cub/warp/warp_reduce.cuh +++ b/cub/cub/warp/warp_reduce.cuh @@ -668,8 +668,31 @@ private: using _TempStorage = cub::NullType; public: - struct TempStorage : Uninitialized<_TempStorage> - {}; + struct InternalWarpReduce + { + struct TempStorage : Uninitialized<_TempStorage> + {}; + + __device__ __forceinline__ InternalWarpReduce(TempStorage & /*temp_storage */) {} + + template + __device__ __forceinline__ T Reduce(T input, + int /* valid_items */, + ReductionOp /* reduction_op */) + { + return input; + } + + template + __device__ __forceinline__ T SegmentedReduce(T input, + FlagT /* flag */, + ReductionOp /* reduction_op */) + { + return input; + } + }; + + using TempStorage = typename InternalWarpReduce::TempStorage; __device__ __forceinline__ WarpReduce(TempStorage & /*temp_storage */) {} diff --git a/cub/test/catch2_test_block_reduce.cu b/cub/test/catch2_test_block_reduce.cu index f2ce06d4d2..06cfe32f03 100644 --- a/cub/test/catch2_test_block_reduce.cu +++ b/cub/test/catch2_test_block_reduce.cu @@ -137,7 +137,7 @@ struct max_full_tile_op_t using types = c2h::type_list; using vec_types = c2h::type_list; -// %PARAM% TEST_DIM_X dimx 7:32:65:128 +// %PARAM% TEST_DIM_X dimx 1:7:32:65:128 // %PARAM% TEST_DIM_YZ dimyz 1:2 using block_dim_xs = c2h::enum_type_list;