diff --git a/benchmarks/cpp/nvfuser/batch_norm.cpp b/benchmarks/cpp/nvfuser/batch_norm.cpp index 57e889b19fb8d..265f293407ca2 100644 --- a/benchmarks/cpp/nvfuser/batch_norm.cpp +++ b/benchmarks/cpp/nvfuser/batch_norm.cpp @@ -78,10 +78,10 @@ static void NvFuserScheduler_BatchNorm( const float kEps = 1e-5; std::vector input_shape{ - benchmark_state.range(0), - benchmark_state.range(1), - benchmark_state.range(2), - benchmark_state.range(2)}; + benchmark_state.range(0) + SIZE_OFFSET, + benchmark_state.range(1) + SIZE_OFFSET, + benchmark_state.range(2) + SIZE_OFFSET, + benchmark_state.range(2) + SIZE_OFFSET}; // inputs at::manual_seed(0); @@ -117,10 +117,10 @@ static void Baseline_BatchNorm( const float kMomentum = 0.1; const float kEps = 1e-5; std::vector input_shape{ - benchmark_state.range(0), - benchmark_state.range(1), - benchmark_state.range(2), - benchmark_state.range(2)}; + benchmark_state.range(0) + SIZE_OFFSET, + benchmark_state.range(1) + SIZE_OFFSET, + benchmark_state.range(2) + SIZE_OFFSET, + benchmark_state.range(2) + SIZE_OFFSET}; // inputs at::manual_seed(0); @@ -152,9 +152,11 @@ static void Baseline_BatchNorm( clearL2Cache(); cudaDeviceSynchronize(); + + CudaKernelTimer timer; for (auto _ : benchmark_state) { - CudaKernelTimer timer; - auto output = at::batch_norm( + timer.restart(); + auto output = at::_ops::_batch_norm_impl_index::call( at_x, ato_weight, ato_bias, @@ -225,6 +227,36 @@ NVFUSER_BENCHMARK_RUN(NvFuserScheduler_BatchNorm_fp16) ->Unit(benchmark::kMicrosecond) ->UseManualTime(); +NVFUSER_BENCHMARK_RUN(NvFuserScheduler_BatchNorm_fp16) + // ->RangeMultiplier(2) + ->Ranges({{64, 256}, {64, 64}, {7, 112}}) + ->Unit(benchmark::kMicrosecond) + ->UseManualTime(); + +NVFUSER_BENCHMARK_RUN(NvFuserScheduler_BatchNorm_fp16) + // ->RangeMultiplier(2) + ->Ranges({{64, 256}, {256, 256}, {7, 56}}) + ->Unit(benchmark::kMicrosecond) + ->UseManualTime(); + +NVFUSER_BENCHMARK_RUN(NvFuserScheduler_BatchNorm_fp16) + // ->RangeMultiplier(2) + ->Ranges({{64, 256}, {512, 512}, {7, 28}}) + ->Unit(benchmark::kMicrosecond) + ->UseManualTime(); + +NVFUSER_BENCHMARK_RUN(NvFuserScheduler_BatchNorm_fp16) + // ->RangeMultiplier(2) + ->Ranges({{64, 256}, {1024, 1024}, {7, 14}}) + ->Unit(benchmark::kMicrosecond) + ->UseManualTime(); + +NVFUSER_BENCHMARK_RUN(NvFuserScheduler_BatchNorm_fp16) + // ->RangeMultiplier(2) + ->Ranges({{64, 256}, {2048, 2048}, {7, 7}}) + ->Unit(benchmark::kMicrosecond) + ->UseManualTime(); + //------------------------------------------------------------------------------ BENCHMARK(Baseline_BatchNorm_cuDNN_fp32) @@ -251,3 +283,33 @@ BENCHMARK(Baseline_BatchNorm_cuDNN_fp16) ->Ranges({{2, 64}, {2, 32}, {2, 256}}) ->Unit(benchmark::kMicrosecond) ->UseManualTime(); + +BENCHMARK(Baseline_BatchNorm_cuDNN_fp16) + // ->RangeMultiplier(2) + ->Ranges({{64, 256}, {64, 64}, {7, 112}}) + ->Unit(benchmark::kMicrosecond) + ->UseManualTime(); + +BENCHMARK(Baseline_BatchNorm_cuDNN_fp16) + // ->RangeMultiplier(2) + ->Ranges({{64, 256}, {256, 256}, {7, 56}}) + ->Unit(benchmark::kMicrosecond) + ->UseManualTime(); + +BENCHMARK(Baseline_BatchNorm_cuDNN_fp16) + // ->RangeMultiplier(2) + ->Ranges({{64, 256}, {512, 512}, {7, 28}}) + ->Unit(benchmark::kMicrosecond) + ->UseManualTime(); + +BENCHMARK(Baseline_BatchNorm_cuDNN_fp16) + // ->RangeMultiplier(2) + ->Ranges({{64, 256}, {1024, 1024}, {7, 14}}) + ->Unit(benchmark::kMicrosecond) + ->UseManualTime(); + +BENCHMARK(Baseline_BatchNorm_cuDNN_fp16) + // ->RangeMultiplier(2) + ->Ranges({{64, 256}, {2048, 2048}, {7, 7}}) + ->Unit(benchmark::kMicrosecond) + ->UseManualTime(); diff --git a/benchmarks/cpp/nvfuser/batch_norm_backward.cpp b/benchmarks/cpp/nvfuser/batch_norm_backward.cpp index 77a09564de5d2..88c26b848ba9e 100644 --- a/benchmarks/cpp/nvfuser/batch_norm_backward.cpp +++ b/benchmarks/cpp/nvfuser/batch_norm_backward.cpp @@ -89,10 +89,10 @@ static void NvFuserScheduler_BatchNorm_BWD( const float kEps = 1e-5; std::vector input_shape{ - benchmark_state.range(0), - benchmark_state.range(1), - benchmark_state.range(2), - benchmark_state.range(2)}; + benchmark_state.range(0) + SIZE_OFFSET, + benchmark_state.range(1) + SIZE_OFFSET, + benchmark_state.range(2) + SIZE_OFFSET, + benchmark_state.range(2) + SIZE_OFFSET}; at::manual_seed(0); auto options = @@ -130,10 +130,10 @@ static void Baseline_BatchNorm_BWD( const float kMomentum = 0.1; const float kEps = 1e-5; std::vector input_shape{ - benchmark_state.range(0), - benchmark_state.range(1), - benchmark_state.range(2), - benchmark_state.range(2)}; + benchmark_state.range(0) + SIZE_OFFSET, + benchmark_state.range(1) + SIZE_OFFSET, + benchmark_state.range(2) + SIZE_OFFSET, + benchmark_state.range(2) + SIZE_OFFSET}; at::manual_seed(0); auto options = @@ -166,13 +166,12 @@ static void Baseline_BatchNorm_BWD( kMomentum, kEps, true); - cudaDeviceSynchronize(); - - // Sync everything up before we start clearL2Cache(); cudaDeviceSynchronize(); + + CudaKernelTimer timer; for (auto _ : benchmark_state) { - CudaKernelTimer timer; + timer.restart(); at::_ops::cudnn_batch_norm_backward::call( input, @@ -186,7 +185,6 @@ static void Baseline_BatchNorm_BWD( std::get<3>(fwd_result)); benchmark_state.SetIterationTime(timer.elapsed() / 1000.0); - cudaDeviceSynchronize(); clearL2Cache(); cudaDeviceSynchronize(); } @@ -249,6 +247,36 @@ NVFUSER_BENCHMARK_RUN(NvFuserScheduler_BatchNorm_BWD_fp16) ->Unit(benchmark::kMicrosecond) ->UseManualTime(); +NVFUSER_BENCHMARK_RUN(NvFuserScheduler_BatchNorm_BWD_fp16) + // ->RangeMultiplier(2) + ->Ranges({{64, 256}, {64, 64}, {7, 112}}) + ->Unit(benchmark::kMicrosecond) + ->UseManualTime(); + +NVFUSER_BENCHMARK_RUN(NvFuserScheduler_BatchNorm_BWD_fp16) + // ->RangeMultiplier(2) + ->Ranges({{64, 256}, {256, 256}, {7, 56}}) + ->Unit(benchmark::kMicrosecond) + ->UseManualTime(); + +NVFUSER_BENCHMARK_RUN(NvFuserScheduler_BatchNorm_BWD_fp16) + // ->RangeMultiplier(2) + ->Ranges({{64, 256}, {512, 512}, {7, 28}}) + ->Unit(benchmark::kMicrosecond) + ->UseManualTime(); + +NVFUSER_BENCHMARK_RUN(NvFuserScheduler_BatchNorm_BWD_fp16) + // ->RangeMultiplier(2) + ->Ranges({{64, 256}, {1024, 1024}, {7, 14}}) + ->Unit(benchmark::kMicrosecond) + ->UseManualTime(); + +NVFUSER_BENCHMARK_RUN(NvFuserScheduler_BatchNorm_BWD_fp16) + // ->RangeMultiplier(2) + ->Ranges({{64, 256}, {2048, 2048}, {7, 7}}) + ->Unit(benchmark::kMicrosecond) + ->UseManualTime(); + //------------------------------------------------------------------------------ BENCHMARK(Baseline_BatchNorm_BWD_cuDNN_fp32) @@ -275,3 +303,33 @@ BENCHMARK(Baseline_BatchNorm_BWD_cuDNN_fp16) ->Ranges({{2, 64}, {2, 32}, {2, 256}}) ->Unit(benchmark::kMicrosecond) ->UseManualTime(); + +BENCHMARK(Baseline_BatchNorm_BWD_cuDNN_fp16) + // ->RangeMultiplier(2) + ->Ranges({{64, 256}, {64, 64}, {7, 112}}) + ->Unit(benchmark::kMicrosecond) + ->UseManualTime(); + +BENCHMARK(Baseline_BatchNorm_BWD_cuDNN_fp16) + // ->RangeMultiplier(2) + ->Ranges({{64, 256}, {256, 256}, {7, 56}}) + ->Unit(benchmark::kMicrosecond) + ->UseManualTime(); + +BENCHMARK(Baseline_BatchNorm_BWD_cuDNN_fp16) + // ->RangeMultiplier(2) + ->Ranges({{64, 256}, {512, 512}, {7, 28}}) + ->Unit(benchmark::kMicrosecond) + ->UseManualTime(); + +BENCHMARK(Baseline_BatchNorm_BWD_cuDNN_fp16) + // ->RangeMultiplier(2) + ->Ranges({{64, 256}, {1024, 1024}, {7, 14}}) + ->Unit(benchmark::kMicrosecond) + ->UseManualTime(); + +BENCHMARK(Baseline_BatchNorm_BWD_cuDNN_fp16) + // ->RangeMultiplier(2) + ->Ranges({{64, 256}, {2048, 2048}, {7, 7}}) + ->Unit(benchmark::kMicrosecond) + ->UseManualTime(); \ No newline at end of file diff --git a/benchmarks/cpp/nvfuser/bert.cpp b/benchmarks/cpp/nvfuser/bert.cpp index a1dd58d5646a3..9d4213dbf7543 100644 --- a/benchmarks/cpp/nvfuser/bert.cpp +++ b/benchmarks/cpp/nvfuser/bert.cpp @@ -113,10 +113,10 @@ static void MagicScheduler_DivMaxSoftDropFwd( Fusion fusion; FusionGuard fg(&fusion); - auto w = benchmark_state.range(0); - auto x = benchmark_state.range(1); - auto y = benchmark_state.range(2); - auto z = benchmark_state.range(3); + auto w = benchmark_state.range(0) + SIZE_OFFSET; + auto x = benchmark_state.range(1) + SIZE_OFFSET; + auto y = benchmark_state.range(2) + SIZE_OFFSET; + auto z = benchmark_state.range(3) + SIZE_OFFSET; setupDivMaxSoftmaxDropoutForward(&fusion, dtype); @@ -171,10 +171,10 @@ static void MagicScheduler_DivMaxSoftDropBwd( Fusion fusion; FusionGuard fg(&fusion); - auto w = benchmark_state.range(0); - auto x = benchmark_state.range(1); - auto y = benchmark_state.range(2); - auto z = benchmark_state.range(3); + auto w = benchmark_state.range(0) + SIZE_OFFSET; + auto x = benchmark_state.range(1) + SIZE_OFFSET; + auto y = benchmark_state.range(2) + SIZE_OFFSET; + auto z = benchmark_state.range(3) + SIZE_OFFSET; setupDivMaxSoftmaxDropoutBackward(&fusion, dtype); @@ -286,9 +286,9 @@ static void MagicScheduler_BiasDropoutAddLayernormFwd( Fusion fusion; FusionGuard fg(&fusion); - auto x = benchmark_state.range(0); - auto y = benchmark_state.range(1); - auto z = benchmark_state.range(2); + auto x = benchmark_state.range(0) + SIZE_OFFSET; + auto y = benchmark_state.range(1) + SIZE_OFFSET; + auto z = benchmark_state.range(2) + SIZE_OFFSET; setupBiasDropoutAddLayernormFwd(&fusion, dtype); @@ -402,9 +402,9 @@ static void MagicScheduler_BiasDropoutAddLayernormBwd1( Fusion fusion; FusionGuard fg(&fusion); - auto x = benchmark_state.range(0); - auto y = benchmark_state.range(1); - auto z = benchmark_state.range(2); + auto x = benchmark_state.range(0) + SIZE_OFFSET; + auto y = benchmark_state.range(1) + SIZE_OFFSET; + auto z = benchmark_state.range(2) + SIZE_OFFSET; setupBiasDropoutAddLayernormBwd1(&fusion, dtype); @@ -513,9 +513,9 @@ static void MagicScheduler_BiasDropoutAddLayernormBwd2( Fusion fusion; FusionGuard fg(&fusion); - auto x = benchmark_state.range(0); - auto y = benchmark_state.range(1); - auto z = benchmark_state.range(2); + auto x = benchmark_state.range(0) + SIZE_OFFSET; + auto y = benchmark_state.range(1) + SIZE_OFFSET; + auto z = benchmark_state.range(2) + SIZE_OFFSET; setupBiasDropoutAddLayernormBwd2(&fusion, dtype); @@ -606,9 +606,9 @@ static void MagicScheduler_BiasDropoutAddLayernormBwd3( Fusion fusion; FusionGuard fg(&fusion); - auto x = benchmark_state.range(0); - auto y = benchmark_state.range(1); - auto z = benchmark_state.range(2); + auto x = benchmark_state.range(0) + SIZE_OFFSET; + auto y = benchmark_state.range(1) + SIZE_OFFSET; + auto z = benchmark_state.range(2) + SIZE_OFFSET; setupBiasDropoutAddLayernormBwd3(&fusion, dtype); diff --git a/benchmarks/cpp/nvfuser/broadcast.cpp b/benchmarks/cpp/nvfuser/broadcast.cpp index d693ff68bf85a..e6b1b6e4bdd42 100644 --- a/benchmarks/cpp/nvfuser/broadcast.cpp +++ b/benchmarks/cpp/nvfuser/broadcast.cpp @@ -51,8 +51,8 @@ static void NvFuserScheduler_Broadcast( FusionExecutorCache* fusion_executor_cache, DataType dtype, int bcast_dim) { - auto bcast_size = benchmark_state.range(0); - auto iter_size = benchmark_state.range(1); + auto bcast_size = benchmark_state.range(0) + SIZE_OFFSET; + auto iter_size = benchmark_state.range(1) + SIZE_OFFSET; at::manual_seed(0); auto options = @@ -99,8 +99,8 @@ static void Baseline_Broadcast( benchmark::State& benchmark_state, DataType dtype, int bcast_dim) { - auto bcast_size = benchmark_state.range(0); - auto iter_size = benchmark_state.range(1); + auto bcast_size = benchmark_state.range(0) + SIZE_OFFSET; + auto iter_size = benchmark_state.range(1) + SIZE_OFFSET; at::manual_seed(0); auto options = diff --git a/benchmarks/cpp/nvfuser/instance_norm.cpp b/benchmarks/cpp/nvfuser/instance_norm.cpp index 007291d75f5f1..2e03186146ddd 100644 --- a/benchmarks/cpp/nvfuser/instance_norm.cpp +++ b/benchmarks/cpp/nvfuser/instance_norm.cpp @@ -71,10 +71,10 @@ static void NvFuserScheduler_InstanceNorm( TORCH_INTERNAL_ASSERT(dtype == DataType::Float || dtype == DataType::Half); std::vector input_shape{ - benchmark_state.range(0), - benchmark_state.range(2), - benchmark_state.range(1), - benchmark_state.range(1)}; + benchmark_state.range(0) + SIZE_OFFSET, + benchmark_state.range(2) + SIZE_OFFSET, + benchmark_state.range(1) + SIZE_OFFSET, + benchmark_state.range(1) + SIZE_OFFSET}; // inputs at::manual_seed(0); @@ -112,10 +112,10 @@ static void Baseline_InstanceNorm( TORCH_INTERNAL_ASSERT(dtype == DataType::Float || dtype == DataType::Half); std::vector input_shape{ - benchmark_state.range(0), - benchmark_state.range(2), - benchmark_state.range(1), - benchmark_state.range(1)}; + benchmark_state.range(0) + SIZE_OFFSET, + benchmark_state.range(2) + SIZE_OFFSET, + benchmark_state.range(1) + SIZE_OFFSET, + benchmark_state.range(1) + SIZE_OFFSET}; const float kMomentum = 0.1; const float kEps = 1e-5; const auto aten_dtype = data_type_to_aten(dtype); diff --git a/benchmarks/cpp/nvfuser/layer_norm.cpp b/benchmarks/cpp/nvfuser/layer_norm.cpp index 7500ac8525b6b..3dde2518c7131 100644 --- a/benchmarks/cpp/nvfuser/layer_norm.cpp +++ b/benchmarks/cpp/nvfuser/layer_norm.cpp @@ -60,7 +60,8 @@ static void NvFuserScheduler_LayerNorm( TORCH_INTERNAL_ASSERT(dtype == DataType::Float || dtype == DataType::Half); std::vector input_shape{ - benchmark_state.range(0), benchmark_state.range(1)}; + benchmark_state.range(0) + SIZE_OFFSET, + benchmark_state.range(1) + SIZE_OFFSET}; const float kEps = 1e-5; // inputs @@ -89,7 +90,8 @@ static void Baseline_LayerNorm( TORCH_INTERNAL_ASSERT(dtype == DataType::Float || dtype == DataType::Half); std::vector input_shape{ - benchmark_state.range(0), benchmark_state.range(1)}; + benchmark_state.range(0) + SIZE_OFFSET, + benchmark_state.range(1) + SIZE_OFFSET}; const int kReductionAxis = 1; std::vector norm_shape; for (int idx = kReductionAxis; idx < input_shape.size(); ++idx) { diff --git a/benchmarks/cpp/nvfuser/layer_norm_backward.cpp b/benchmarks/cpp/nvfuser/layer_norm_backward.cpp index 045465e712539..074aad3ec9bf5 100644 --- a/benchmarks/cpp/nvfuser/layer_norm_backward.cpp +++ b/benchmarks/cpp/nvfuser/layer_norm_backward.cpp @@ -82,7 +82,8 @@ static void NvFuserScheduler_LayerNorm_BWD( TORCH_INTERNAL_ASSERT(dtype == DataType::Float || dtype == DataType::Half); std::vector input_shape{ - benchmark_state.range(0), benchmark_state.range(1)}; + benchmark_state.range(0) + SIZE_OFFSET, + benchmark_state.range(1) + SIZE_OFFSET}; // inputs at::manual_seed(0); @@ -115,7 +116,8 @@ static void Baseline_LayerNorm_BWD( TORCH_INTERNAL_ASSERT(dtype == DataType::Float || dtype == DataType::Half); std::vector input_shape{ - benchmark_state.range(0), benchmark_state.range(1)}; + benchmark_state.range(0) + SIZE_OFFSET, + benchmark_state.range(1) + SIZE_OFFSET}; const int kReductionAxis = 1; std::vector norm_shape; for (int idx = kReductionAxis; idx < input_shape.size(); ++idx) { diff --git a/benchmarks/cpp/nvfuser/reduction.cpp b/benchmarks/cpp/nvfuser/reduction.cpp index c25097963dbc8..db5217c5bc870 100644 --- a/benchmarks/cpp/nvfuser/reduction.cpp +++ b/benchmarks/cpp/nvfuser/reduction.cpp @@ -50,8 +50,8 @@ static void NvFuserScheduler_Reduction( FusionExecutorCache* fusion_executor_cache, DataType dtype, int reduction_dim) { - auto reduction_size = benchmark_state.range(0); - auto iter_size = benchmark_state.range(1); + auto reduction_size = benchmark_state.range(0) + SIZE_OFFSET; + auto iter_size = benchmark_state.range(1) + SIZE_OFFSET; at::manual_seed(0); auto options = @@ -95,8 +95,8 @@ static void Baseline_Reduction( benchmark::State& benchmark_state, DataType dtype, int reduction_dim) { - auto reduction_size = benchmark_state.range(0); - auto iter_size = benchmark_state.range(1); + auto reduction_size = benchmark_state.range(0) + SIZE_OFFSET; + auto iter_size = benchmark_state.range(1) + SIZE_OFFSET; at::manual_seed(0); auto options = diff --git a/benchmarks/cpp/nvfuser/scale_bias_relu.cpp b/benchmarks/cpp/nvfuser/scale_bias_relu.cpp index 47ed9047f1592..9efcb10c4eec5 100644 --- a/benchmarks/cpp/nvfuser/scale_bias_relu.cpp +++ b/benchmarks/cpp/nvfuser/scale_bias_relu.cpp @@ -112,15 +112,16 @@ static void NvFuserScheduler_SBR( DataType dtype) { // N, H, W, C format std::vector input_shape{ - benchmark_state.range(0), - benchmark_state.range(1), - benchmark_state.range(1), - benchmark_state.range(2)}; + benchmark_state.range(0) + SIZE_OFFSET, + benchmark_state.range(1) + SIZE_OFFSET, + benchmark_state.range(1) + SIZE_OFFSET, + benchmark_state.range(2) + SIZE_OFFSET}; std::vector bcast_shape{1, 1, 1, -1}; // inputs at::manual_seed(0); - std::vector static_bcast_shape{1, 1, 1, benchmark_state.range(2)}; + std::vector static_bcast_shape{ + 1, 1, 1, benchmark_state.range(2) + SIZE_OFFSET}; auto options = at::TensorOptions().dtype(data_type_to_aten(dtype)).device(at::kCUDA, 0); at::Tensor at_x = at::randn(input_shape, options); @@ -168,11 +169,11 @@ static void NvFuserScheduler_SBR( static void Baseline_SBR(benchmark::State& benchmark_state, DataType dtype) { // N, H, W, C format std::vector input_shape{ - benchmark_state.range(0), - benchmark_state.range(1), - benchmark_state.range(1), - benchmark_state.range(2)}; - std::vector bcast_shape{benchmark_state.range(2)}; + benchmark_state.range(0) + SIZE_OFFSET, + benchmark_state.range(1) + SIZE_OFFSET, + benchmark_state.range(1) + SIZE_OFFSET, + benchmark_state.range(2) + SIZE_OFFSET}; + std::vector bcast_shape{benchmark_state.range(2) + SIZE_OFFSET}; // inputs at::manual_seed(0); @@ -214,11 +215,11 @@ static void NvFuserScheduler_SBR_Norm( DataType dtype) { // N, H, W, C format std::vector input_shape{ - benchmark_state.range(0), - benchmark_state.range(1), - benchmark_state.range(1), - benchmark_state.range(2)}; - std::vector bcast_shape{benchmark_state.range(2)}; + benchmark_state.range(0) + SIZE_OFFSET, + benchmark_state.range(1) + SIZE_OFFSET, + benchmark_state.range(1) + SIZE_OFFSET, + benchmark_state.range(2) + SIZE_OFFSET}; + std::vector bcast_shape{benchmark_state.range(2) + SIZE_OFFSET}; // inputs at::manual_seed(0); @@ -274,11 +275,12 @@ static void Baseline_SBR_Norm( DataType dtype) { // N, H, W, C format std::vector input_shape{ - benchmark_state.range(0), - benchmark_state.range(1), - benchmark_state.range(1), - benchmark_state.range(2)}; - std::vector bcast_shape{1, 1, 1, benchmark_state.range(2)}; + benchmark_state.range(0) + SIZE_OFFSET, + benchmark_state.range(1) + SIZE_OFFSET, + benchmark_state.range(1) + SIZE_OFFSET, + benchmark_state.range(2) + SIZE_OFFSET}; + std::vector bcast_shape{ + 1, 1, 1, benchmark_state.range(2) + SIZE_OFFSET}; // inputs at::manual_seed(0); diff --git a/benchmarks/cpp/nvfuser/softmax.cpp b/benchmarks/cpp/nvfuser/softmax.cpp index 3964e03671fab..8fdb5c226edae 100644 --- a/benchmarks/cpp/nvfuser/softmax.cpp +++ b/benchmarks/cpp/nvfuser/softmax.cpp @@ -52,8 +52,8 @@ static void NvFuserScheduler_Softmax( auto options = at::TensorOptions().dtype(data_type_to_aten(dtype)).device(at::kCUDA, 0); - auto reduction_size = benchmark_state.range(0); - auto iter_size = benchmark_state.range(1); + auto reduction_size = benchmark_state.range(0) + SIZE_OFFSET; + auto iter_size = benchmark_state.range(1) + SIZE_OFFSET; at::Tensor aten_input = (reduction_axis ? at::randn({iter_size, reduction_size}, options) @@ -72,7 +72,8 @@ static void NvFuserScheduler_Softmax( static void Softmax_WarpReduceReference(benchmark::State& benchmark_state) { auto dtype = DataType::Float; std::vector input_shape{ - benchmark_state.range(0), benchmark_state.range(1)}; + benchmark_state.range(0) + SIZE_OFFSET, + benchmark_state.range(1) + SIZE_OFFSET}; auto fusion_ptr = std::make_unique(); auto fusion = fusion_ptr.get(); @@ -117,7 +118,8 @@ static void Softmax_WarpReduceReference(benchmark::State& benchmark_state) { static void Softmax_WarpReduce(benchmark::State& benchmark_state) { auto dtype = DataType::Float; std::vector input_shape{ - benchmark_state.range(0), benchmark_state.range(1)}; + benchmark_state.range(0) + SIZE_OFFSET, + benchmark_state.range(1) + SIZE_OFFSET}; auto fusion_ptr = std::make_unique(); auto fusion = fusion_ptr.get(); @@ -170,13 +172,13 @@ static void Softmax_WarpReduce(benchmark::State& benchmark_state) { } BENCHMARK(Softmax_WarpReduce) - ->RangeMultiplier(2) + // ->RangeMultiplier(2) ->Ranges({{8, 8}, {16 * 197, 16 * 197}}) ->Unit(benchmark::kMicrosecond) ->UseManualTime(); BENCHMARK(Softmax_WarpReduceReference) - ->RangeMultiplier(2) + // ->RangeMultiplier(2) ->Ranges({{8, 8}, {16 * 197, 16 * 197}}) ->Unit(benchmark::kMicrosecond) ->UseManualTime(); @@ -191,8 +193,8 @@ static void Baseline_Softmax( auto options = at::TensorOptions().dtype(data_type_to_aten(dtype)).device(at::kCUDA, 0); - auto reduction_size = benchmark_state.range(0); - auto iter_size = benchmark_state.range(1); + auto reduction_size = benchmark_state.range(0) + SIZE_OFFSET; + auto iter_size = benchmark_state.range(1) + SIZE_OFFSET; at::Tensor aten_input = (reduction_axis ? at::randn({iter_size, reduction_size}, options) diff --git a/benchmarks/cpp/nvfuser/softmax_backward.cpp b/benchmarks/cpp/nvfuser/softmax_backward.cpp index 1bf2e623291a2..7664a0976e430 100644 --- a/benchmarks/cpp/nvfuser/softmax_backward.cpp +++ b/benchmarks/cpp/nvfuser/softmax_backward.cpp @@ -58,8 +58,8 @@ static void NvFuserScheduler_Softmax_BWD( auto options = at::TensorOptions().dtype(data_type_to_aten(dtype)).device(at::kCUDA, 0); - auto reduction_size = benchmark_state.range(0); - auto iter_size = benchmark_state.range(1); + auto reduction_size = benchmark_state.range(0) + SIZE_OFFSET; + auto iter_size = benchmark_state.range(1) + SIZE_OFFSET; at::Tensor input = (reduction_axis ? at::randn({iter_size, reduction_size}, options) @@ -92,8 +92,8 @@ static void Baseline_Softmax_BWD( auto options = at::TensorOptions().dtype(data_type_to_aten(dtype)).device(at::kCUDA, 0); - auto reduction_size = benchmark_state.range(0); - auto iter_size = benchmark_state.range(1); + auto reduction_size = benchmark_state.range(0) + SIZE_OFFSET; + auto iter_size = benchmark_state.range(1) + SIZE_OFFSET; at::Tensor input = (reduction_axis ? at::randn({iter_size, reduction_size}, options) diff --git a/benchmarks/cpp/nvfuser/softmax_dropout.cpp b/benchmarks/cpp/nvfuser/softmax_dropout.cpp index 828940933f418..2a923317be1a4 100644 --- a/benchmarks/cpp/nvfuser/softmax_dropout.cpp +++ b/benchmarks/cpp/nvfuser/softmax_dropout.cpp @@ -76,7 +76,8 @@ static void NvFuserScheduler_SoftmaxDropout( TORCH_INTERNAL_ASSERT(dtype == DataType::Float || dtype == DataType::Half); // reduce across 1, [256, 12, 100, 8] - std::vector input_shape{256, 12, 100, benchmark_state.range(0)}; + std::vector input_shape{ + 256, 12, 100, benchmark_state.range(0) + SIZE_OFFSET}; constexpr int kHiddenSize = 768; constexpr int kNumAttentionHeads = 12; @@ -113,7 +114,8 @@ static void Baseline_Softmax_Dropout( benchmark::State& benchmark_state, const int kReductionAxis, DataType dtype) { - std::vector input_shape{256, 12, 100, benchmark_state.range(0)}; + std::vector input_shape{ + 256, 12, 100, benchmark_state.range(0) + SIZE_OFFSET}; constexpr int kHiddenSize = 768; constexpr int kNumAttentionHeads = 12; diff --git a/benchmarks/cpp/nvfuser/utils.h b/benchmarks/cpp/nvfuser/utils.h index b4a2f3a7a9164..4760237bb004d 100644 --- a/benchmarks/cpp/nvfuser/utils.h +++ b/benchmarks/cpp/nvfuser/utils.h @@ -18,6 +18,10 @@ using namespace torch::jit::fuser::cuda; +// Easy flag to tick if we want to increment or decrement from pow2 sizes in +// benchmarks +#define SIZE_OFFSET 0 + std::string toString(ReductionParams rparams); std::string toString(PointwiseParams params); std::string toString(LaunchParams lparams);