[hip-kernel-provider] Remove HIP_PLUGIN_BN_MAXN macro (ROCm#8095)

EwanC · web-flow · commit e868a35c6762 · 2026-06-12T03:09:02.000Z
## Motivation

In the hip-kernel-provider, the `HIP_PLUGIN_BN_MAXN` macro is never
changed from the default value by plans, this is also the case in the
original MIOpen code. Therefore to make it clear that this value is
invariant, this PR folds the default value into the batchnorm
configuration code as a constant expression.

## Technical Details

* Replace macro with hardcoded literatal
* Change related macro guarded code to constexpr
* Add comment documenting relationship between `max_n` and kernel
selection heuristics in BN plan.

## Test Plan

Run `ninja check` in hip-kernel-provider build

## Test Result

Passes on MI100
diff --git a/dnn-providers/hip-kernel-provider/src/engines/hip_mlops_engine/kernels/batchnorm/BatchNormBwdSpatial.cpp b/dnn-providers/hip-kernel-provider/src/engines/hip_mlops_engine/kernels/batchnorm/BatchNormBwdSpatial.cpp
@@ -679,7 +679,7 @@ struct BatchNormBwdSpatialImpl<3, FpType, FpPrecType, FpAccumType>
         FpPrecType ds = 0;
         FpPrecType db = 0;
 
-        // maybe unused if HIP_PLUGIN_BN_N >= HIP_PLUGIN_BN_MAXN
+        // Unused if hip_plugin_bn_config::n >= hip_plugin_bn_config::max_n
         FpPrecType batchvalues[hip_plugin_bn_config::n];
         FpPrecType dyvalues[hip_plugin_bn_config::n];
 
diff --git a/dnn-providers/hip-kernel-provider/src/engines/hip_mlops_engine/kernels/batchnorm/BatchNormFwdTrainSpatial.cpp b/dnn-providers/hip-kernel-provider/src/engines/hip_mlops_engine/kernels/batchnorm/BatchNormFwdTrainSpatial.cpp
@@ -526,9 +526,8 @@ struct BatchNormFwdTrainSpatialImpl<3, FpType, FpPrecType, FpAccumType>
         unsigned int grpid = blockIdx.x;
         unsigned int cidx = grpid * hip_plugin_bn_config::hw;
 
-#if(HIP_PLUGIN_BN_N < HIP_PLUGIN_BN_MAXN)
+        // Unused if hip_plugin_bn_config::n >= hip_plugin_bn_config::max_n
         FpType minibatch[HIP_PLUGIN_BN_N];
-#endif
 
         if(lid == 0)
         {
@@ -545,9 +544,10 @@ struct BatchNormFwdTrainSpatialImpl<3, FpType, FpPrecType, FpAccumType>
                     mean += xin;
                     variance = fma(xin, xin, variance);
 
-#if(HIP_PLUGIN_BN_N < HIP_PLUGIN_BN_MAXN)
-                    minibatch[n] = (*(in + index));
-#endif
+                    if constexpr(hip_plugin_bn_config::n < hip_plugin_bn_config::max_n)
+                    {
+                        minibatch[n] = (*(in + index));
+                    }
                 }};
         }
         __syncthreads();
@@ -592,12 +592,15 @@ struct BatchNormFwdTrainSpatialImpl<3, FpType, FpPrecType, FpAccumType>
             static_unroll_count<unsigned int, 0, hip_plugin_bn_config::n, 1, 2>{
                 [&](unsigned int n) { // apply normalization
                     index = n * hip_plugin_bn_config::chw + cidx + lid;
-#if(HIP_PLUGIN_BN_N < HIP_PLUGIN_BN_MAXN)
-                    inhat = (cast<FpPrecType>(minibatch[n]) - mean)
-                            * invVariance; // (in[index] - mean) * invVariance;
-#else
-                    inhat = (cast<FpPrecType>(*(in + index)) - mean) * invVariance;
-#endif
+                    if constexpr(hip_plugin_bn_config::n < hip_plugin_bn_config::max_n)
+                    {
+                        inhat = (cast<FpPrecType>(minibatch[n]) - mean)
+                                * invVariance; // (in[index] - mean) * invVariance;
+                    }
+                    else
+                    {
+                        inhat = (cast<FpPrecType>(*(in + index)) - mean) * invVariance;
+                    }
                     out[index] = cast<FpType>(
                         hip_kernel_provider::applyActivation<FpPrecType,
                                                              hip_kernel_provider::ActivationMode{
diff --git a/dnn-providers/hip-kernel-provider/src/engines/hip_mlops_engine/kernels/batchnorm/Configuration.hpp b/dnn-providers/hip-kernel-provider/src/engines/hip_mlops_engine/kernels/batchnorm/Configuration.hpp
@@ -150,7 +150,6 @@ template <typename HipKernelConfig,
           typename Architecture,
           int Variant,
           int NCHW,
-          int MaxN,
           int NElements,
           int N,
           int C,
@@ -170,7 +169,6 @@ struct proto_config
     static_assert(Vectorize == 0 || Vectorize == 1, "Vectorize must be 0 or 1");
     static_assert(UseNodpp == 0 || UseNodpp == 1, "UseNodpp must be 0 or 1");
     static_assert(NCHW >= 0, "HIP_PLUGIN_BN_NCHW should be always >= 0");
-    static_assert(MaxN >= 0, "HIP_PLUGIN_BN_MAXN should be always >= 0");
     static_assert(C >= 0, "HIP_PLUGIN_BN_C should be always >= 0");
     static_assert(N >= 0, "HIP_PLUGIN_BN_N should be always >= 0");
     static_assert(HW >= 0, "HIP_PLUGIN_BN_HW should be always >= 0");
@@ -196,7 +194,14 @@ struct proto_config
                      : 0); // TODO: not sure if 0 should be the default value of this.
     static constexpr auto launch_dim = LaunchDim{};
     static constexpr unsigned int nchw = static_cast<unsigned int>(NCHW);
-    static constexpr unsigned int max_n = static_cast<unsigned int>(MaxN);
+
+    // `max_n` is a limit on the number of batch elements from the x tensor that can be cached in per-thread memory as part of
+    // variant 3 kernels. If batchsize (n_elements) exceeds this limit then the kernel doesn't cache the global memory accesses.
+    // The constant of 65 is related to the heuristic used in `defaultConfigSpatialSingle` where Batchnorm plans select the kernel
+    // variant to use. In BN Fwd training, the heuristics only selects variant 3 when N <= 32, so the caching optimization is always
+    // used. In BN Bwd training, the heuristics only selects variant 3 when N > 64, so the caching optimization is never used. If
+    // there are future changes to how the heuristic selects variant 3 kernels, then it may be worth revisiting this caching limit.
+    static constexpr unsigned int max_n = 65;
     static constexpr unsigned int n_elements = static_cast<unsigned int>(NElements);
     static constexpr unsigned int n = static_cast<unsigned int>(N);
     static constexpr unsigned int c = static_cast<unsigned int>(C);
@@ -267,7 +272,6 @@ using config = hip_kernel_provider::batchnorm::detail::proto_config<
                                                                 HIP_PLUGIN_GFX115X>,
     HIP_PLUGIN_BN_VARIANT,
     HIP_PLUGIN_BN_NCHW,
-    HIP_PLUGIN_BN_MAXN,
     HIP_PLUGIN_BN_N_ELEMENTS,
     HIP_PLUGIN_BN_N,
     HIP_PLUGIN_BN_C,
diff --git a/dnn-providers/hip-kernel-provider/src/engines/hip_mlops_engine/plans/batchnorm/BatchnormBwdPlan.cpp b/dnn-providers/hip-kernel-provider/src/engines/hip_mlops_engine/plans/batchnorm/BatchnormBwdPlan.cpp
@@ -293,7 +293,6 @@ void BatchnormBwdPlan::compile(const IKernelCompiler& kernelCompiler,
         options.update("HIP_PLUGIN_BN_GRP1", ylocalsize);
         options.update("HIP_PLUGIN_BN_GRP2", zlocalsize);
         options.update("HIP_PLUGIN_BN_LDS_SIZE", ldsSize);
-        options.update("HIP_PLUGIN_BN_MAXN", 65);
         options.update("HIP_PLUGIN_BN_VEC_SIZE", config.vectorsize);
 
         _compiledProgram = kernelCompiler.compile("BatchNormBwdSpatial.cpp", options);
diff --git a/dnn-providers/hip-kernel-provider/src/engines/hip_mlops_engine/plans/batchnorm/BatchnormKernelCompileOptions.hpp b/dnn-providers/hip-kernel-provider/src/engines/hip_mlops_engine/plans/batchnorm/BatchnormKernelCompileOptions.hpp
@@ -54,7 +54,6 @@ class BatchnormKernelCompileOptions : public KernelCompileOptions
         add("HIP_PLUGIN_BN_VARIANT", 255);
         add("HIP_PLUGIN_BN_USESAVED", 0);
         add("HIP_PLUGIN_BN_NCHW", 1);
-        add("HIP_PLUGIN_BN_MAXN", 65);
         add("HIP_PLUGIN_BN_VECTORIZE", 0);
         add("HIP_PLUGIN_BN_VEC_SIZE", 1);
         add("HIP_PLUGIN_BN_STASH_METHOD", 0);
diff --git a/dnn-providers/hip-kernel-provider/src/tests/engines/hip_mlops_engine/plans/Batchnorm/TestBatchnormBwdPlan.cpp b/dnn-providers/hip-kernel-provider/src/tests/engines/hip_mlops_engine/plans/Batchnorm/TestBatchnormBwdPlan.cpp
@@ -189,7 +189,6 @@ TEST(TestBatchnormBwdPlan, CompileDefaultSetsCorrectDefines)
     EXPECT_TRUE(hasOption("-DHIP_PLUGIN_BN_GRP1=1"));
     EXPECT_TRUE(hasOption("-DHIP_PLUGIN_BN_GRP2=1"));
     EXPECT_TRUE(hasOption("-DHIP_PLUGIN_BN_N_ELEMENTS=HIP_PLUGIN_BN_N"));
-    EXPECT_TRUE(hasOption("-DHIP_PLUGIN_BN_MAXN=65"));
     EXPECT_TRUE(hasOption("-DHIP_PLUGIN_USE_AMDGCN=0"));
     EXPECT_TRUE(hasOption("-DHIP_PLUGIN_BN_LOOP_UNROLL_MAXN=768"));
     EXPECT_TRUE(hasOption("-DHIP_PLUGIN_BN_LDS_SIZE=1024"));
diff --git a/dnn-providers/hip-kernel-provider/src/tests/engines/hip_mlops_engine/plans/Batchnorm/TestBatchnormFwdInferencePlan.cpp b/dnn-providers/hip-kernel-provider/src/tests/engines/hip_mlops_engine/plans/Batchnorm/TestBatchnormFwdInferencePlan.cpp
@@ -310,7 +310,6 @@ TEST(TestBatchnormFwdInferencePlanFp32, CompileDefaultSetsCorrectDefines)
     EXPECT_TRUE(hasOption("-DHIP_PLUGIN_BN_GRP1=256"));
     EXPECT_TRUE(hasOption("-DHIP_PLUGIN_BN_GRP2=1"));
     EXPECT_TRUE(hasOption("-DHIP_PLUGIN_BN_N_ELEMENTS=HIP_PLUGIN_BN_N"));
-    EXPECT_TRUE(hasOption("-DHIP_PLUGIN_BN_MAXN=65"));
     EXPECT_TRUE(hasOption("-DHIP_PLUGIN_USE_AMDGCN=0"));
     EXPECT_TRUE(hasOption("-DHIP_PLUGIN_BN_LOOP_UNROLL_MAXN=768"));
     EXPECT_TRUE(hasOption("-DHIP_PLUGIN_BN_LDS_SIZE=256"));
diff --git a/dnn-providers/hip-kernel-provider/src/tests/engines/hip_mlops_engine/plans/Batchnorm/TestBatchnormFwdInferenceWithVariancePlan.cpp b/dnn-providers/hip-kernel-provider/src/tests/engines/hip_mlops_engine/plans/Batchnorm/TestBatchnormFwdInferenceWithVariancePlan.cpp
@@ -317,7 +317,6 @@ TEST(TestBatchnormFwdInferenceWithVariancePlanFp32, CompileDefaultSetsCorrectDef
     EXPECT_TRUE(hasOption("-DHIP_PLUGIN_BN_GRP1=256"));
     EXPECT_TRUE(hasOption("-DHIP_PLUGIN_BN_GRP2=1"));
     EXPECT_TRUE(hasOption("-DHIP_PLUGIN_BN_N_ELEMENTS=HIP_PLUGIN_BN_N"));
-    EXPECT_TRUE(hasOption("-DHIP_PLUGIN_BN_MAXN=65"));
     EXPECT_TRUE(hasOption("-DHIP_PLUGIN_USE_AMDGCN=0"));
     EXPECT_TRUE(hasOption("-DHIP_PLUGIN_BN_LOOP_UNROLL_MAXN=768"));
     EXPECT_TRUE(hasOption("-DHIP_PLUGIN_BN_LDS_SIZE=256"));
diff --git a/dnn-providers/hip-kernel-provider/src/tests/engines/hip_mlops_engine/plans/Batchnorm/TestBatchnormFwdTrainingActivPlan.cpp b/dnn-providers/hip-kernel-provider/src/tests/engines/hip_mlops_engine/plans/Batchnorm/TestBatchnormFwdTrainingActivPlan.cpp
@@ -897,7 +897,6 @@ TEST(TestBatchnormFwdTrainingActivPlan, CompileDefaultSetsCorrectDefines)
     EXPECT_TRUE(hasOption("-DHIP_PLUGIN_BN_NGRPS=1"));
     EXPECT_TRUE(hasOption("-DHIP_PLUGIN_BN_NGRPS2=1"));
     EXPECT_TRUE(hasOption("-DHIP_PLUGIN_BN_N_ELEMENTS=1"));
-    EXPECT_TRUE(hasOption("-DHIP_PLUGIN_BN_MAXN=65"));
     EXPECT_TRUE(hasOption("-DHIP_PLUGIN_USE_AMDGCN=0"));
     EXPECT_TRUE(hasOption("-DHIP_PLUGIN_BN_LOOP_UNROLL_MAXN=768"));
     EXPECT_TRUE(hasOption("-DHIP_PLUGIN_BN_LDS_SIZE=1024"));
diff --git a/dnn-providers/hip-kernel-provider/src/tests/engines/hip_mlops_engine/plans/Batchnorm/TestBatchnormFwdTrainingPlan.cpp b/dnn-providers/hip-kernel-provider/src/tests/engines/hip_mlops_engine/plans/Batchnorm/TestBatchnormFwdTrainingPlan.cpp
@@ -351,7 +351,6 @@ TEST(TestBatchnormFwdTrainingPlan, CompileDefaultSetsCorrectDefines)
     EXPECT_TRUE(hasOption("-DHIP_PLUGIN_BN_GRP1_FINAL=1"));
     EXPECT_TRUE(hasOption("-DHIP_PLUGIN_BN_GRP2_FINAL=1"));
     EXPECT_TRUE(hasOption("-DHIP_PLUGIN_BN_N_ELEMENTS=1"));
-    EXPECT_TRUE(hasOption("-DHIP_PLUGIN_BN_MAXN=65"));
     EXPECT_TRUE(hasOption("-DHIP_PLUGIN_USE_AMDGCN=0"));
     EXPECT_TRUE(hasOption("-DHIP_PLUGIN_BN_LOOP_UNROLL_MAXN=768"));
     EXPECT_TRUE(hasOption("-DHIP_PLUGIN_BN_LDS_SIZE=1024"));
