Remove cluster_dims references from Intel's backend to match upstream

Author: slawblauciak

third_party/intel/backend/compiler.py

@@ -246,12 +246,6 @@ def make_ttir(cls, mod, metadata, opt):
     @classmethod
     @track
     def make_ttgir(cls, mod, metadata, opt, properties):
-        cluster_info = intel.ClusterInfo()
-        if opt.cluster_dims is not None:
-            cluster_info.clusterDimX = opt.cluster_dims[0]
-            cluster_info.clusterDimY = opt.cluster_dims[1]
-            cluster_info.clusterDimZ = opt.cluster_dims[2]
-
         # Annotate module with information required by subsequent transformations.
         pm = ir.pass_manager(mod.context)
         pm.enable_debug()
@@ -303,7 +297,6 @@ def make_ttgir(cls, mod, metadata, opt, properties):
             intel.passes.ttgpuir.add_optimize_reduction_locality(pm)
         intel.passes.arith.add_arith_emulate_unsupported_floats(pm, ["bf16"], "f32")
         pm.run(mod, 'make_ttgir')
-        metadata["cluster_dims"] = (cluster_info.clusterDimX, cluster_info.clusterDimY, cluster_info.clusterDimZ)
         return mod
 
     def gluon_to_ttgir(self, src, metadata, options):

third_party/intel/backend/driver.py

@@ -738,16 +738,6 @@ def format_of(ty):
   int threads_per_warp = PyLong_AsLong(threads_per_warp_attr);
   Py_DECREF(threads_per_warp_attr);
 
-  // extract cluster dims
-  PyObject *clusterDim =  PyObject_GetAttrString(kernel_metadata, "cluster_dims");
-  if (!PyTuple_Check(kernel_metadata)) {{
-    PyErr_SetString(PyExc_TypeError, "kernel_metadata.cluster_dims must be a tuple");
-    return NULL;
-  }}
-  int clusterDimX   = PyLong_AsLong(PyTuple_GetItem(clusterDim, 0));
-  int clusterDimY   = PyLong_AsLong(PyTuple_GetItem(clusterDim, 1));
-  int clusterDimZ   = PyLong_AsLong(PyTuple_GetItem(clusterDim, 2));
-  Py_DECREF(clusterDim);
   // extract launch metadata
   if (launch_enter_hook != Py_None){{
     PyObject* ret = PyObject_CallOneArg(launch_enter_hook, launch_metadata);

third_party/intel/include/Dialect/TritonIntelGPU/Transforms/Passes.h

@@ -13,14 +13,6 @@
 
 namespace mlir::triton::gpu::intel {
 
-// Used by Triton runtime
-struct ClusterInfo {
-  ClusterInfo() = default;
-  unsigned clusterDimX = 1u;
-  unsigned clusterDimY = 1u;
-  unsigned clusterDimZ = 1u;
-};
-
 /// Split barrier scope
 enum SplitBarrierScope { None = 0, Workgroup = 1, Subgroup = 2 };
 

third_party/intel/triton_xpu.cc

@@ -136,19 +136,6 @@ void init_triton_intel(py::module &&m) {
   init_triton_intel_passes_ttgpuir(passes.def_submodule("ttgpuir"));
   init_triton_intel_passes_arith(passes.def_submodule("arith"));
 
-  // cluster info
-  py::class_<gpu::intel::ClusterInfo>(m, "ClusterInfo")
-      .def(py::init<>())
-      .def_readwrite("clusterDimX", &gpu::intel::ClusterInfo::clusterDimX)
-      .def_readwrite("clusterDimY", &gpu::intel::ClusterInfo::clusterDimY)
-      .def_readwrite("clusterDimZ", &gpu::intel::ClusterInfo::clusterDimZ)
-      .def("__repr__", [](gpu::intel::ClusterInfo &self) {
-        std::ostringstream oss;
-        oss << "(" << self.clusterDimX << ", " << self.clusterDimY << ", "
-            << self.clusterDimZ << ")";
-        return oss.str();
-      });
-
   // Split barrier scope enum
   py::enum_<gpu::intel::SplitBarrierScope>(m, "SplitBarrierScope")
       .value("none", gpu::intel::SplitBarrierScope::None)