Remove MatmulParams::rotate_ldmatrix_out_of_main_loop

I can't find any commit in which this option was ever actually used.
This is the commit where the option was originally introduced:
https://github.com/csarofeen/pytorch/pull/2488/files#diff-e7a5a84a2cfeddeb15669f07105bdb3722a796600ea9e1f2eb25afb29283457eR22
We've gone this long without the ability to disable loop rotation, so
either we should change the condition in the schedulers to respect it,
or just remove it.

csrc/python_frontend/python_bindings.cpp

@@ -651,7 +651,6 @@ void defineHeuristicParamBindings(py::module& nvfuser) {
       .PARAM(MatmulParams, circular_buffer_options)
       .PARAM(MatmulParams, supported_vec_size)
       .PARAM(MatmulParams, async_gmem_load_operands)
-      .PARAM(MatmulParams, rotate_ldmatrix_out_of_main_loop)
       .PARAM(MatmulParams, grid_swizzle_factor)
       .PARAM(MatmulParams, use_smem_epilogue)
       .PARAM(MatmulParams, promote_prologue_smem_reuse)

csrc/scheduler/matmul_heuristic.h

@@ -138,9 +138,6 @@ class MatmulParams : public HeuristicParams {
     }
   } supported_vec_size;
 
-  //! Whether to rotate the ldmatrix out of the main loop
-  bool rotate_ldmatrix_out_of_main_loop = true;
-
   //! (Ampere+) Use cp.async to load operands.
   bool async_gmem_load_operands = false;
 
@@ -191,8 +188,6 @@ class MatmulParams : public HeuristicParams {
        << circular_buffer_options.toString() << "\n"
        << supported_vec_size.toString() << "\n"
        << nvfuser::toString(tile_sizes) << "\n"
-       << "Rotate ldmatrix out of main loop: "
-       << (rotate_ldmatrix_out_of_main_loop ? "true" : "false") << "\n"
        << "Async global mem load: "
        << (async_gmem_load_operands ? "true" : "false") << "\n"
        << "Indexing mode: "
@@ -216,9 +211,8 @@ class MatmulParams : public HeuristicParams {
 
   size_t hash() const override {
     // combine boolean flags for hashing
-    size_t attr_hash = (static_cast<size_t>(promote_prologue_smem_reuse) << 3) |
-        (static_cast<size_t>(use_smem_epilogue) << 2) |
-        (static_cast<size_t>(rotate_ldmatrix_out_of_main_loop) << 1) |
+    size_t attr_hash = (static_cast<size_t>(promote_prologue_smem_reuse) << 2) |
+        (static_cast<size_t>(use_smem_epilogue) << 1) |
         (static_cast<size_t>(async_gmem_load_operands));
 
     // combined hash
@@ -240,8 +234,6 @@ class MatmulParams : public HeuristicParams {
 
     return other->cparams == cparams && other->mma_macro == mma_macro &&
         other->async_gmem_load_operands == async_gmem_load_operands &&
-        other->rotate_ldmatrix_out_of_main_loop ==
-        rotate_ldmatrix_out_of_main_loop &&
         other->tile_sizes == tile_sizes &&
         other->circular_buffer_options == circular_buffer_options &&
         other->supported_vec_size == supported_vec_size &&

csrc/scheduler/matmul_heuristic_plugin.cpp

@@ -146,8 +146,6 @@ void copyParamsToConfig(KernelConfig* config, const MatmulParams* mparams) {
                                                                            : 1;
   config->circular_buffer_smem_read =
       mparams->circular_buffer_options.circular_buffer_smem_read;
-  config->rotate_ldmatrix_out_of_main_loop =
-      mparams->rotate_ldmatrix_out_of_main_loop;
   config->problem.supported_vec_size.a = (uint8_t)mparams->supported_vec_size.a;
   config->problem.supported_vec_size.b = (uint8_t)mparams->supported_vec_size.b;
   config->problem.supported_vec_size.epilogue =
@@ -190,8 +188,6 @@ void copyConfigToParams(MatmulParams* mparams, const KernelConfig* config) {
   }
   mparams->circular_buffer_options.circular_buffer_smem_read =
       config->circular_buffer_smem_read;
-  mparams->rotate_ldmatrix_out_of_main_loop =
-      config->rotate_ldmatrix_out_of_main_loop;
 
   // enable circular buffering if configured
   mparams->circular_buffer_options.circular_buffer_smem_write =

csrc/scheduler/matmul_heuristic_plugin_api.h

@@ -77,7 +77,6 @@ struct KernelConfig {
   uint8_t grid_swizzle_factor = 0;
   uint8_t cta_order = 0;
   bool circular_buffer_smem_read = true;
-  bool rotate_ldmatrix_out_of_main_loop = true;
   bool async_gmem_load_operands = true;
 
  public: