[AMD] Explicitly set ConvertToBufferOps small tensor option (#8470)

Recently, a new pass option is added to bufferop pass. However, the
driver side (compiler.py/compiler.cc) does not feed the option-value
explicitly to the pass. Default value is used for the option-value.

To resolve this problem, all pass-options are fed explicitly by driver,
i.e., no default value.

Author: yangshuxin

third_party/amd/backend/compiler.py

@@ -243,7 +243,7 @@ def make_ttgir(mod, metadata, options):
         if knobs.amd.use_buffer_ops:
             amd.passes.ttgpuir.add_canonicalize_pointers(pm)
             passes.common.add_canonicalizer(pm)
-            amd.passes.ttgpuir.add_convert_to_buffer_ops(pm, options.arch, knobs.amd.use_buffer_atomics)
+            amd.passes.ttgpuir.add_convert_to_buffer_ops(pm, options.arch, knobs.amd.use_buffer_atomics, False)
 
         amd.passes.ttgpuir.add_fold_true_cmpi(pm)
         passes.common.add_canonicalizer(pm)

third_party/amd/python/triton_amd.cc

@@ -83,9 +83,9 @@ void init_triton_amd_passes_ttgpuir(py::module &&m) {
     pm.addNestedPass<mlir::triton::FuncOp>(
         mlir::createTritonAMDGPUCanonicalizePointers());
   });
-  ADD_PASS_OPTION_WRAPPER_2("add_convert_to_buffer_ops",
+  ADD_PASS_OPTION_WRAPPER_3("add_convert_to_buffer_ops",
                             mlir::createTritonAMDGPUConvertToBufferOps,
-                            const std::string &, bool);
+                            const std::string &, bool, bool);
   ADD_PASS_WRAPPER_0("add_reorder_instructions",
                      mlir::createTritonAMDGPUReorderInstructions);
   ADD_PASS_WRAPPER_0("add_fold_true_cmpi", mlir::createTritonAMDFoldTrueCmpI);