diff --git a/.github/workflows/enflame3.6-gcu400-build-and-test.yml b/.github/workflows/enflame3.6-gcu400-build-and-test.yml
index adc2ebbdb..f91f2681c 100644
--- a/.github/workflows/enflame3.6-gcu400-build-and-test.yml
+++ b/.github/workflows/enflame3.6-gcu400-build-and-test.yml
@@ -88,7 +88,8 @@ jobs:
 
           python3 -m pytest -s python/test/tle \
               --ignore=python/test/tle/unit/test_tle_distributed_d2d.py \
-              --ignore=python/test/tle/unit/test_tle_get_local_pe.py
+              --ignore=python/test/tle/unit/test_tle_get_local_pe.py \
+              --ignore=python/test/tle/unit/test_tle_d2d_barrier.py
 
           ## tle raw test
           python3 -m pytest -s third_party/enflame/python/test/tle/raw
diff --git a/.github/workflows/nv3.6-build-and-test.yml b/.github/workflows/nv3.6-build-and-test.yml
index 73f660c85..39d190274 100644
--- a/.github/workflows/nv3.6-build-and-test.yml
+++ b/.github/workflows/nv3.6-build-and-test.yml
@@ -72,6 +72,7 @@ jobs:
         run: |
           set -x
           source ~/env-3.6.sh
+          export FLAGCX_RECOMPILE=ON
           export USE_FLAGCX=ON
           MAX_JOBS=32 python3 -m pip install . --no-build-isolation
 
@@ -112,9 +113,11 @@ jobs:
           python3 -m pytest -s python/test/tle/integration
           python3 -m pytest -s python/test/tle/unit \
               --ignore=python/test/tle/unit/test_tle_distributed_d2d.py \
-              --ignore=python/test/tle/unit/test_tle_get_local_pe.py
+              --ignore=python/test/tle/unit/test_tle_get_local_pe.py \
+              --ignore=python/test/tle/unit/test_tle_d2d_barrier.py
           CUDA_LAUNCH_BLOCKING=1 bash python/test/tle/unit/test_tle_distributed_d2d.sh
           CUDA_LAUNCH_BLOCKING=1 bash python/test/tle/unit/test_tle_get_local_pe.sh
+          CUDA_LAUNCH_BLOCKING=1 bash python/test/tle/unit/test_tle_d2d_barrier.sh
           ## flagtree hints python tutorials
           python3 python/tutorials/hints/01/01-vector-add.py --only_unit_test
           # python3 python/tutorials/hints/02/02-fused-softmax.py --only_unit_test
diff --git a/.gitmodules b/.gitmodules
deleted file mode 100644
index c9f02c419..000000000
--- a/.gitmodules
+++ /dev/null
@@ -1,23 +0,0 @@
-# Copyright 2026 FlagOS Contributors
-#
-# Permission is hereby granted, free of charge, to any person obtaining a copy
-# of this software and associated documentation files (the "Software"), to deal
-# in the Software without restriction, including without limitation the rights
-# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
-# copies of the Software, and to permit persons to whom the Software is
-# furnished to do so, subject to the following conditions:
-#
-# The above copyright notice and this permission notice shall be included in all
-# copies or substantial portions of the Software.
-#
-# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
-# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
-# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
-# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
-# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
-# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
-# SOFTWARE.
-
-[submodule "third_party/tileir/third_party/cuda-tile"]
-	path = third_party/tileir/third_party/cuda-tile
-	url = https://github.com/NVIDIA/cuda-tile
diff --git a/packaging/rpm/helpers/Dockerfile.rpm b/packaging/rpm/helpers/Dockerfile.rpm
index 49d8e05b9..8a520c346 100644
--- a/packaging/rpm/helpers/Dockerfile.rpm
+++ b/packaging/rpm/helpers/Dockerfile.rpm
@@ -41,14 +41,12 @@ COPY test /src/test
 COPY unittest /src/unittest
 COPY CMakeLists.txt pyproject.toml setup.py MANIFEST.in LICENSE README.md /src/
 
-# FLAGTREE_DEFAULT_BACKENDS=nvidia,amd matches the historical package
-# contents and keeps the tileir backend (GCC >= 13 + cuda-tile
+# Keep the tileir backend (GCC >= 13 + cuda-tile
 # submodule) out of the distro build.
 # The trailing rm keeps only /wheels in this layer: the CMake build
 # tree and the LLVM tarball cache (~/.triton) would otherwise add
 # several GB and push CI runners into ENOSPC.
 RUN unset FLAGTREE_BACKEND && \
-    FLAGTREE_DEFAULT_BACKENDS=nvidia,amd \
     MAX_JOBS=4 python3 -m pip wheel . --no-build-isolation --no-deps -w /wheels -v && \
     ls -lh /wheels/ && \
     test -n "$(ls /wheels/flagtree-*.whl 2>/dev/null)" && \
diff --git a/python/setup_tools/setup_helper.py b/python/setup_tools/setup_helper.py
index 1bfc29d32..ebba5f217 100644
--- a/python/setup_tools/setup_helper.py
+++ b/python/setup_tools/setup_helper.py
@@ -519,6 +519,8 @@ def uninstall_triton():
 
 download_flagtree_third_party("flagcx", condition=(not flagtree_backend), hook="handle_flagcx", required=True)
 
+download_flagtree_third_party("tileir", condition=(flagtree_backend == "tileir"), required=True)
+
 handle_flagtree_backend()
 
 # iluvatar
diff --git a/python/setup_tools/utils/__init__.py b/python/setup_tools/utils/__init__.py
index 6f27066e7..af8ca5395 100644
--- a/python/setup_tools/utils/__init__.py
+++ b/python/setup_tools/utils/__init__.py
@@ -60,6 +60,10 @@ def _register_submodule(self, submodule):
     },
     {"name": "flir", "url": "https://github.com/FlagTree/flir.git"},
     {"name": "flagcx", "url": "https://github.com/flagos-ai/FlagCX.git", "relative_path": "tle/third_party/flagcx"},
+    {
+        "name": "tileir", "url": "https://github.com/NVIDIA/cuda-tile", "relative_path": "tileir/third_party/cuda-tile",
+        "commit_id": "2e5ccba66fb3afdba34b26cf358418283027c248"
+    },
 ))
 
 
diff --git a/python/setup_tools/utils/default.py b/python/setup_tools/utils/default.py
index 701339063..f1d10f949 100644
--- a/python/setup_tools/utils/default.py
+++ b/python/setup_tools/utils/default.py
@@ -27,7 +27,7 @@
 
 
 def printinfo(msgs):
-    print(f" [TLE-DIST-INFO]: {msgs}. \n")
+    print(f" [TLE-DIST-INFO]: {msgs}.")
 
 
 class FlagCXRegistrar:
@@ -37,6 +37,29 @@ def __init__(self, external):
         self.shared_lib_name = "libflagcx.so"
         self._set_path(external)
 
+    def _is_flagcx_recompile_required(self):
+        ENVS = ("FLAGCX_RECOMPILE", "DIST_RECOMPILE")
+        for env in ENVS:
+            env_value = os.environ.get(env, "0")
+            if env_value in ("1", "true", "True", "ON"):
+                printinfo(f"Recompiling FlagCX due to {env}={env_value}\n ")
+                return env_value
+        return False
+
+    def _is_cache_available(self):
+        ENVS = ("FLAGCX_CACHE", "DIST_CACHE")
+
+        for env in ENVS:
+            env_value = os.environ.get(env, "1")
+            if env_value in ("0", "false", "False", "OFF", "clean"):
+                printinfo(f"Skipping using cache at {self.cache_lib_dir} due to {env}={env_value}\n ")
+            if env_value in ("clean"):
+                shutil.rmtree(self.cache_lib_dir, ignore_errors=True)
+                os.makedirs(self.cache_lib_dir, exist_ok=True)
+                printinfo(f"Cache {self.cache_lib_dir} cleaned due to {env}={env_value}\n ")
+            return False
+        return True
+
     def _set_path(self, external):
         submodule = external['backend']
         flagtree_cache = external['cache']
@@ -70,15 +93,18 @@ def get_compile_cmds(self):
 
     def _compile_and_cache(self):
         cmds = self.get_compile_cmds()
+        is_unused_cache = self._is_cache_available()
+        is_recompile = self._is_flagcx_recompile_required()
+        is_unused_cache = is_unused_cache or is_recompile
 
         for lib_name, cmd in cmds.items():
             cache_path = getattr(self, f"{lib_name.split('.')[0]}_cache_path")
             src_path = getattr(self, f"{lib_name.split('.')[0]}_src_path")
             runtime_path = self._get_runtime_path(lib_name)
-            if cache_path.exists():
+            if cache_path.exists() and not is_unused_cache:
                 printinfo(f"{lib_name} already exists in cache, skipping compilation ...")
                 shutil.copy(cache_path, runtime_path)
-            elif src_path.exists():
+            elif src_path.exists() and not is_recompile:
                 printinfo(f"{lib_name} already exists in build directory, copying to cache...")
                 shutil.copy(src_path, cache_path)
                 shutil.copy(src_path, runtime_path)
@@ -93,7 +119,7 @@ def _compile_and_cache(self):
                 shutil.copy(src_path, cache_path)
                 shutil.copy(src_path, runtime_path)
                 printinfo(f"{lib_name} copied from {src_path} to cache at {cache_path}")
-                printinfo(f"[32m{lib_name} copied from {src_path} to cache at {runtime_path}")
+                printinfo(f"{lib_name} copied from {src_path} to cache at {runtime_path}")
 
     def _copy_required_files(self):
         dst = Path(self.flagtree_dir) / "python" / "triton" / "experimental" / "tle" / "language" / "flagcx_wrapper.py"
diff --git a/python/setup_tools/utils/tools.py b/python/setup_tools/utils/tools.py
index a77025aab..0907a0851 100644
--- a/python/setup_tools/utils/tools.py
+++ b/python/setup_tools/utils/tools.py
@@ -42,12 +42,7 @@ def _get_flagtree_root() -> str:
 
 @dataclass
 class FlagtreeConfigs:
-    # Overridable for distro/packaging builds whose toolchain can't build
-    # every default backend (e.g. tileir requires GCC >= 13 and the
-    # cuda-tile submodule):
-    #   FLAGTREE_DEFAULT_BACKENDS=nvidia,amd pip wheel .
-    default_backends: tuple = field(default_factory=lambda: tuple(
-        b for b in os.environ.get("FLAGTREE_DEFAULT_BACKENDS", "nvidia,amd,tileir").replace(" ", "").split(",") if b))
+    default_backends: tuple = ("nvidia", "amd")
     plugin_backends: tuple = ("cambricon", "ascend", "aipu", "tsingmicro", "enflame", "hcu", "thrive")
     use_cuda_toolkit_backends: tuple = ('aipu', 'tileir')
     language_extra_backends: tuple = ('xpu', 'mthreads', "cambricon")
@@ -69,12 +64,12 @@ class FlagtreeConfigs:
     }))
 
     def __post_init__(self):
+        backends = list(self.default_backends)
+        _backends = [backend for backend in backends if os.environ.get(f"USE_{backend.upper()}", "ON").upper() != "OFF"]
+        self.default_backends = tuple(_backends)
         self.flagtree_submodule_dir = os.path.join(self.flagtree_root_dir, "third_party")
         self.activated_module = self._activate_device_module()
 
-    def non_tileir_default_backends(self):
-        return tuple(backend for backend in self.default_backends if backend != "tileir")
-
     def _activate_device_module(self, suffix=".py"):
         backend = self.flagtree_backend or "default"
         module_path = Path(os.path.dirname(__file__)) / backend
diff --git a/python/test/tle/integration/test_tle_distributed.py b/python/test/tle/integration/test_tle_distributed.py
index ca0a53b81..fc9e01b79 100644
--- a/python/test/tle/integration/test_tle_distributed.py
+++ b/python/test/tle/integration/test_tle_distributed.py
@@ -995,7 +995,7 @@ def test_remote_const_shard_load_high_block_encoding_no_regression(self):
         assert re.search(r"tensor<256x!tt\.ptr<f16,\s*3>,\s*#blocked[0-9]*>", ttgir) is not None
         assert re.search(
             r"\"tle\.remote_pointers\"\(%[^,]+,\s*%[^)]+\)\s*"
-            r"<\{space\s*=\s*\"cluster\"\}>\s*:\s*"
+            r"<\{[^}]*space\s*=\s*\"cluster\"[^}]*\}>\s*:\s*"
             r"\(tensor<256x!tt\.ptr<f16,\s*3>,\s*#blocked[0-9]*>,\s*i32\)\s*->\s*"
             r"tensor<256x!tt\.ptr<f16,\s*7>,\s*#blocked[0-9]*>",
             ttgir,
diff --git a/python/test/tle/unit/test_tle_d2d_barrier.py b/python/test/tle/unit/test_tle_d2d_barrier.py
new file mode 100644
index 000000000..cb2a78d72
--- /dev/null
+++ b/python/test/tle/unit/test_tle_d2d_barrier.py
@@ -0,0 +1,87 @@
+import triton.experimental.tle.language as tle
+import torch
+import triton
+import triton.language as tl
+import torch.distributed as dist
+
+DEVICE_MESH = tle.device_mesh(tle.MeshConfig(device=2))
+N = 8
+
+
+@triton.jit()
+def _barrier_d2d_kernel(out_ptr, device_dptr: tl.constexpr, mesh: tl.constexpr):
+    pid = tl.program_id(0)
+    local_rank = tle.shard_id(mesh, 'device', device_dptr=device_dptr)
+    n_rank = mesh.shape[0]
+    peer = (local_rank + 1) % n_rank  # noqa: F841
+
+    remote_mem = tle.remote(
+        device_dptr,
+        space="device",
+        dtype=tl.float32,
+        shard_id=peer,
+        offset=pid,
+    )
+    val = tl.load(remote_mem)
+    tl.store(out_ptr + pid, val)
+    tle.distributed_barrier(device_dptr=device_dptr, space="device")
+
+
+def _ir_verify(output, device_dptr, grid):
+    compiled = _barrier_d2d_kernel.warmup(
+        device_dptr=device_dptr,
+        out_ptr=output,
+        mesh=DEVICE_MESH,
+        grid=(grid, ),
+        num_ctas=1,
+        num_warps=4,
+    )
+    assert "distributed_barrier" in compiled.asm["ttgir"]
+    assert "remote_pointer" in compiled.asm["ttgir"]
+    assert "flagcxIntraBarrier" in compiled.asm['ptx']
+    assert "flagcxGetIntraPointerC" in compiled.asm['ptx']
+    assert "flagcxDevCommGetIntraRank" in compiled.asm['ptx']
+
+
+def _runtime_verify(output, device_dptr, grid, rank, world_size):
+    dist.barrier()
+    _barrier_d2d_kernel[grid](device_dptr=device_dptr, out_ptr=output, mesh=DEVICE_MESH)
+
+    torch.cuda.synchronize()
+
+    import sys
+    peer_rank = (rank + 1) % world_size
+    expected = torch.arange(N, dtype=torch.float32, device="cuda") + peer_rank * 1000
+    if torch.allclose(output, expected):
+        print(f"[Rank {rank}] [PASSED] read peer rank {peer_rank}")
+        print(f"[Rank {rank}] sample output[:4] = {output[:4].tolist()}")
+    else:
+        print(f"[Rank {rank}] [FAILED] read peer rank {peer_rank}")
+        print(f"[Rank {rank}] expected[:4] = {expected[:4].tolist()}")
+        print(f"[Rank {rank}] output[:4] = {output[:4].tolist()}")
+        sys.exit(1)
+
+    tle.cleanup_communicator()
+
+
+class TestD2DBarrier:
+
+    def test_tle_d2d_barrier(self):
+        grid = (N, )
+
+        mem_pool = tle.get_mem_pool()
+        world_size = dist.get_world_size()
+        rank = dist.get_rank()
+        with torch.cuda.use_mem_pool(mem_pool):
+            x = (torch.arange(N, dtype=torch.float32, device="cuda") + rank * 1000).clone()
+
+        device_dptr = tle.create_dist_tensor(x)
+        output = torch.zeros(N, dtype=torch.float32, device="cuda")
+
+        _ir_verify(output, device_dptr, grid)
+
+        _runtime_verify(output, device_dptr, grid, rank, world_size)
+
+
+if __name__ == "__main__":
+    TestD2DBarrier().test_tle_d2d_barrier()
diff --git a/python/test/tle/unit/test_tle_d2d_barrier.sh b/python/test/tle/unit/test_tle_d2d_barrier.sh
new file mode 100644
index 000000000..bb0a89a95
--- /dev/null
+++ b/python/test/tle/unit/test_tle_d2d_barrier.sh
@@ -0,0 +1,53 @@
+rm -rf ~/.triton/cache
+#!/bin/bash
+
+# Check if the debug flag is provided as an argument
+if [ "$1" == "debug" ]; then
+    export NCCL_DEBUG=INFO
+    export NCCL_DEBUG_SUBSYS=all
+    echo "NCCL debug information enabled."
+else
+    unset NCCL_DEBUG
+    unset NCCL_DEBUG_SUBSYS
+    echo "NCCL debug information disabled."
+fi
+#export FLAGCX_SOCKET_IFNAME=eth0
+export FLAGCX_IB_HCA=mlx5_0,mlx5_1,mlx5_2,mlx5_3,mlx5_6,mlx5_7,mlx5_8,mlx5_9
+export FLAGCX_USE_HETERO_COMM=1
+export FLAGCX_MEM_ENABLE=1
+export FLAGCX_VMM_ENABLE=0
+export FLAGCX_P2P_DISABLE=1
+#export CUDA_VISIBLE_DEVICES=0,1,2,3,4,5,6,7
+# Need to preload customized gloo library specified for FlagCX linkage
+#export LD_PRELOAD=/usr/local/lib/libgloo.so
+#export LD_PRELOAD=/usr/local/nccl/build/lib/libnccl.so
+#export TORCH_DISTRIBUTED_DETAIL=DEBUG
+#pytest  FlagTree/python/test/tle/unit/test_tle_distributed_d2d.py
+#export NCCL_DEBUG=INFO
+#export CUDA_LAUNCH_BLOCKING=1
+#export TORCH_USE_CUDA_DSA=1
+#export NCCL_NVLS_ENABLE=0
+#export FLAGCX_DMABUF_ENABLE=1
+#export FLAGCX_DEBUG=TRACE
+#export FLAGCX_DEBUG_SUBSY
+
+run_test() {
+    local script_dir
+    script_dir="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
+
+    torchrun \
+        --nproc_per_node=2 \
+        --nnodes=1 \
+        --node_rank=0 \
+        --master_addr=localhost \
+        --master_port=8333 \
+        "${script_dir}/test_tle_d2d_barrier.py"
+}
+
+run_test
+
+
+if [ $? -ne 0 ]; then
+    echo "ERROR: $CMD failed"
+    exit 1
+fi
diff --git a/python/test/tle/unit/test_tle_distributed_d2d.py b/python/test/tle/unit/test_tle_distributed_d2d.py
index 252abaa2e..4d04096de 100644
--- a/python/test/tle/unit/test_tle_distributed_d2d.py
+++ b/python/test/tle/unit/test_tle_distributed_d2d.py
@@ -32,8 +32,8 @@
 
 @triton.jit
 def lsa_read_kernel(
-    dev_mem_ptr,
     output_ptr,
+    device_dptr: tl.constexpr,
     N: tl.constexpr,
     MY_RANK: tl.constexpr,
     N_RANKS: tl.constexpr,
@@ -42,7 +42,7 @@ def lsa_read_kernel(
     peer = (MY_RANK + 1) % N_RANKS
 
     remote_mem = tle.remote(
-        dev_mem_ptr,
+        device_dptr,
         space="device",
         dtype=tl.float32,
         shard_id=peer,
@@ -77,9 +77,7 @@ def main():
           f"stride={buf_tensor.stride()}, "
           f"sample={buf_tensor[:4].tolist()}")
 
-    dev_comm_ptr, dev_mem_ptr = tle.create_comm_tensor(buf_tensor)
-
-    print(f"[Rank {rank}] dev_comm_ptr={dev_comm_ptr:#x}, dev_mem_ptr={dev_mem_ptr:#x}")
+    device_dptr = tle.create_dist_tensor(buf_tensor)
 
     output = torch.zeros(N, dtype=torch.float32, device="cuda")
 
@@ -87,8 +85,8 @@ def main():
 
     grid = (N, )
     lsa_read_kernel[grid](
-        dev_mem_ptr,
         output,
+        device_dptr=device_dptr,
         N=N,
         MY_RANK=rank,
         N_RANKS=world_size,
diff --git a/python/test/tle/unit/test_tle_get_local_pe.py b/python/test/tle/unit/test_tle_get_local_pe.py
index 566861c85..e87441f0d 100644
--- a/python/test/tle/unit/test_tle_get_local_pe.py
+++ b/python/test/tle/unit/test_tle_get_local_pe.py
@@ -27,9 +27,9 @@
 
 
 @triton.jit
-def _tle_local_pe_kernel(dev_comm_dptr, dev_mem_dptr, out_ptr, mesh: tl.constexpr, BLOCK: tl.constexpr):
+def _tle_local_pe_kernel(out_ptr, device_dptr: tl.constexpr, mesh: tl.constexpr, BLOCK: tl.constexpr):
     pid = tl.program_id(0)  # noqa: F841
-    local_rank = tle.shard_id(mesh, 'device', comm_ptr=dev_comm_dptr)
+    local_rank = tle.shard_id(mesh, 'device', device_dptr=device_dptr)
     n_rank = mesh.shape[0]
     peer = (local_rank + 1) % n_rank  # noqa: F841
 
@@ -43,12 +43,11 @@ def test_tle_local_pe_kernel(self):
         with torch.cuda.use_mem_pool(tle.get_mem_pool()):
             x = torch.randn((N, N), dtype=torch.float32, device="cuda")
         y = torch.empty_like(x)
-        dev_comm_dptr, dev_mem_dptr = tle.create_comm_tensor(x)
+        device_dptr = tle.create_dist_tensor(x)
 
         compiled = _tle_local_pe_kernel.warmup(
-            dev_comm_dptr=dev_comm_dptr,
-            dev_mem_dptr=dev_mem_dptr,
             out_ptr=y,
+            device_dptr=device_dptr,
             mesh=DEVICE_MESH,
             BLOCK=block,
             grid=(grid, ),
@@ -57,8 +56,7 @@ def test_tle_local_pe_kernel(self):
         )
         assert "get_device_id" in compiled.asm["ttgir"]
 
-        _tle_local_pe_kernel[(grid, )](dev_comm_dptr=dev_comm_dptr, dev_mem_dptr=dev_mem_dptr, out_ptr=y,
-                                       mesh=DEVICE_MESH, BLOCK=block)
+        _tle_local_pe_kernel[(grid, )](out_ptr=y, device_dptr=device_dptr, mesh=DEVICE_MESH, BLOCK=block)
 
         tle.cleanup_communicator()
 
diff --git a/python/triton/experimental/tle/language/__init__.py b/python/triton/experimental/tle/language/__init__.py
index 2aa235291..e0f4bdb48 100644
--- a/python/triton/experimental/tle/language/__init__.py
+++ b/python/triton/experimental/tle/language/__init__.py
@@ -43,6 +43,9 @@
     ShardingSpec,
     device_mesh,
     MeshConfig,
+    BarrierKind,
+    MemoryOrder,
+    GroupKind,
     distributed_barrier,
     distributed_dot,
     _infer_submesh_barrier_group,
@@ -56,7 +59,7 @@
     sharding,
 )
 from . import communication
-from .communication import get_mem_pool, create_comm_tensor, cleanup_communicator
+from .communication import get_mem_pool, create_dist_tensor, cleanup_communicator
 
 _EXTENSION_APIS = frozenset({
     "make_tensor_view",
@@ -110,8 +113,11 @@ def __getattr__(name):
     "raw",
     "mem_pool",
     "get_mem_pool",
-    "create_comm_tensor",
+    "create_dist_tensor",
     "cleanup_communicator",
+    "BarrierKind",
+    "MemoryOrder",
+    "GroupKind",
 ]
 
 from . import distributed, gpu, raw
diff --git a/python/triton/experimental/tle/language/communication.py b/python/triton/experimental/tle/language/communication.py
index 5467c3664..ebe2ae941 100644
--- a/python/triton/experimental/tle/language/communication.py
+++ b/python/triton/experimental/tle/language/communication.py
@@ -188,7 +188,7 @@ def init_communicator():
     _init_communicator_ = True
 
 
-def create_comm_tensor(buf_tensor):
+def create_dist_tensor(buf_tensor):
     global comm, rank, dev_mem, dev_comm, win
     buf_ptr = buf_tensor.data_ptr()
     buf_size = buf_tensor.numel() * buf_tensor.element_size()
@@ -225,4 +225,7 @@ def create_comm_tensor(buf_tensor):
 
     # Synchronize all ranks before kernel launch
     dist.barrier()
-    return dev_comm_dptr.value, dev_mem_dptr.value
+
+    from triton.runtime import DistributedRtContext
+    ctx = DistributedRtContext(dev_mem_dptr.value, dev_comm_dptr.value)
+    return ctx
diff --git a/python/triton/experimental/tle/language/distributed.py b/python/triton/experimental/tle/language/distributed.py
index b1754d40e..ac56d4ee2 100644
--- a/python/triton/experimental/tle/language/distributed.py
+++ b/python/triton/experimental/tle/language/distributed.py
@@ -25,7 +25,7 @@
 from dataclasses import dataclass, asdict
 from itertools import product
 from typing import Any, Iterable, Mapping, Sequence, List, Tuple, Union, Optional, Dict
-
+from enum import Enum
 import triton.language.core as tl
 
 Axis = Tuple[str, int]
@@ -47,12 +47,27 @@ def _as_positive_int(value: Any, label: str) -> int:
     return value
 
 
+def _parse_src_arg(builder, src, index=0):
+    try:
+        from triton.runtime import DistributedRtContext
+        src = tl._unwrap_if_constexpr(src)
+        if isinstance(src, DistributedRtContext):
+            return builder.get_int64(src[index])
+        elif src:
+            return src.handle
+        else:
+            return None
+    except Exception:
+        return src.handle
+
+
 # Get the current device id
 @tl.builtin
-def _get_local_rank(dev_mem_ptr, _semantic=None, ret_dtype=tl.int32):
+def _get_local_rank(device_dptr, _semantic=None, ret_dtype=tl.int32):
     builder = _semantic.builder
     ret_ir_ty = ret_dtype.to_ir(builder)
-    result = builder.get_device_id(ret_ir_ty, dev_mem_ptr.handle)
+    ptr = _parse_src_arg(builder, device_dptr, 1)
+    result = builder.get_device_id(ret_ir_ty, ptr)
     return tl.tensor(result, ret_dtype)
 
 
@@ -65,6 +80,28 @@ def n_pes(dev_mem_ptr, _semantic=None, ret_dtype=tl.int32):
     return tl.tensor(result, ret_dtype)
 
 
+class BarrierKind(str, Enum):
+    ARRIVE = "arrive"
+    WAIT = "wait"
+    SYNC = "sync"
+
+
+class MemoryOrder(str, Enum):
+    RELAXED = "relaxed"
+    ACQUIRE = "acquire"
+    RELEASE = "release"
+    ACQ_REL = "acqrel"
+
+
+class GroupKind(str, Enum):
+    THREAD = "thread"
+    WARP = "warp"
+    BLOCK = "block"
+    TILE_SPAN = "tile_span"
+    LANES = "lanes"
+    GRID = "grid"
+
+
 @dataclass
 class MeshConfig:
     """
@@ -595,7 +632,7 @@ def _resolve_launch_axis(mesh: device_mesh, axis: str | int) -> int:
 def shard_id(
     mesh: device_mesh,
     axis: str | int,
-    comm_ptr=None,
+    device_dptr=None,
     _semantic=None,
 ):
     """
@@ -607,8 +644,8 @@ def shard_id(
     mesh = tl._unwrap_if_constexpr(mesh)
     axis = tl._unwrap_if_constexpr(axis)
 
-    if comm_ptr is not None:
-        return _get_local_rank(comm_ptr, _semantic=_semantic, ret_dtype=tl.int32)
+    if axis in ("device", "node"):
+        return _get_local_rank(device_dptr, _semantic=_semantic, ret_dtype=tl.int32)
 
     if not isinstance(mesh, device_mesh):
         raise TypeError(f"mesh must be device_mesh, got {type(mesh).__name__}")
@@ -633,8 +670,40 @@ def shard_id(
     return coord
 
 
+def _parse_device_barrier_args(argType) -> str:
+    argType = tl._unwrap_if_constexpr(argType)
+    argTypes = (BarrierKind, GroupKind, MemoryOrder)
+    if isinstance(argType, argTypes):
+        return argType.value
+    else:
+        return str(argType).lower()
+
+
+def check_and_handle_device_intra_barrier(space: str = None, device_dptr=None,
+                                          barrier_kind: BarrierKind | str = BarrierKind.SYNC,
+                                          group_kind: str | GroupKind = GroupKind.BLOCK, index: int | None = 0,
+                                          order: MemoryOrder | str | int | None = MemoryOrder.ACQ_REL, _semantic=None):
+    if space and space in ("device", "node"):
+        builder = _semantic.builder
+        ptr = _parse_src_arg(builder, device_dptr, 1)
+        builder.create_distributed_barrier(
+            src=ptr,
+            barrier_index=index or 0,
+            space=_parse_device_barrier_args("device"),
+            group_kind=_parse_device_barrier_args(group_kind),
+            order=_parse_device_barrier_args(order),
+            barrier_kind=_parse_device_barrier_args(barrier_kind),
+        )
+        return True
+    return False
+
+
 @tl.builtin
-def distributed_barrier(mesh: device_mesh | None = None, _semantic=None):
+def distributed_barrier(mesh: device_mesh | None = None, device_dptr=None, space: str = None,
+                        group_kind: str | GroupKind = GroupKind.BLOCK,
+                        barrier_kind: BarrierKind | str = BarrierKind.SYNC,
+                        order: MemoryOrder | str | int | None = MemoryOrder.ACQ_REL, _semantic=None,
+                        index: int | None = 0):
     """
     M3 entrypoint: distributed synchronization primitive.
 
@@ -647,6 +716,10 @@ def distributed_barrier(mesh: device_mesh | None = None, _semantic=None):
     subgroup = None
     use_grid = mesh is not None and _mesh_uses_grid_barrier(mesh)
 
+    if check_and_handle_device_intra_barrier(space=space, device_dptr=device_dptr, barrier_kind=barrier_kind,
+                                             group_kind=group_kind, index=index, order=order, _semantic=_semantic):
+        return None
+
     if use_grid:
         if mesh is not None:
             _apply_mesh_grid_launch(mesh, _semantic)
@@ -773,7 +846,7 @@ def _create_remote_pointers_tensor(
         "cluster": (dtype, 7),
         "device": (dtype, 1),
     }.get(space))
-    if tensor.type.is_block():
+    if space == 'cluster' and tensor and tensor.type.is_block():
         remote_type = tl.block_type(remote_ptr_dtype, list(tensor.shape)).to_ir(builder)
     else:
         remote_type = remote_ptr_dtype.to_ir(builder)
@@ -798,11 +871,12 @@ def _create_remote_pointers_tensor(
         # automatic injection (e.g. inside this helper).
         if offset_tensor.dtype != tl.int64:
             offset_tensor = tl.cast(offset_tensor, tl.int64, _semantic=_semantic)
-        remote_op = builder.create_remote_pointers(remote_type, tensor.handle, shard_id_tensor.handle, space,
+        ptr = _parse_src_arg(builder, tensor, 0)
+        remote_op = builder.create_remote_pointers(remote_type, ptr, shard_id_tensor.handle, space,
                                                    offset_tensor.handle)
     else:
         remote_op = builder.create_remote_pointers(remote_type, tensor.handle, shard_id_tensor.handle, space)
-    if tensor.type.is_block():
+    if space == "cluster" and tensor and tensor.type.is_block():
         return tl.tensor(remote_op.get_result(0), tl.block_type(remote_ptr_dtype, list(tensor.shape)))
     return tl.tensor(remote_op.get_result(0), remote_ptr_dtype)
 
@@ -848,7 +922,7 @@ def _remote_pointer(
     _semantic=None,
 ) -> tl.tensor:
 
-    if not isinstance(tensor, tl.tensor):
+    if not isinstance(tensor, tl.tensor) and space not in ("device", "node"):
         raise TypeError(f"tensor must be tl.tensor, got {type(tensor).__name__}")
 
     space = tl._unwrap_if_constexpr(space)
@@ -877,8 +951,8 @@ def _remote_pointer(
 
 @tl.builtin
 def remote(
-    tensor,
-    shard_id,
+    tensor=tl.tensor | None,
+    shard_id=None,
     scope: device_mesh | None = None,
     space: str = "cluster",
     dtype: tl.dtype = None,
@@ -913,7 +987,7 @@ def remote(
 
     # Direct pointer path: support local_ptr scalar/tensor values and return
     # remote pointer with preserved shape.
-    if isinstance(tensor, tl.tensor):
+    if isinstance(tensor, tl.tensor) or (space in ("device", "node")):
         return _remote_pointer(tensor, shard_id, scope=scope, space=space, _semantic=_semantic, dtype=dtype,
                                offset=offset)
 
diff --git a/python/triton/runtime/__init__.py b/python/triton/runtime/__init__.py
index 5f6f81a09..8f1060508 100644
--- a/python/triton/runtime/__init__.py
+++ b/python/triton/runtime/__init__.py
@@ -28,6 +28,7 @@
 from .driver import driver
 from .jit import JITFunction, KernelInterface, MockTensor, TensorWrapper, reinterpret
 from .errors import OutOfResources, InterpreterError
+from ._distributed import DistributedRtContext
 
 __all__ = [
     "autotune",
@@ -45,4 +46,6 @@
     "reinterpret",
     "RemoteCacheBackend",
     "TensorWrapper",
+    # flagtree tle distributed
+    "DistributedRtContext",
 ]
diff --git a/python/triton/runtime/_distributed.py b/python/triton/runtime/_distributed.py
new file mode 100644
index 000000000..70fb3fa75
--- /dev/null
+++ b/python/triton/runtime/_distributed.py
@@ -0,0 +1,76 @@
+class DistributedRtContext:
+    _instance = None
+    _initialized = False
+    _init_count = 0
+
+    def __new__(cls, *args, **kwargs):
+        if cls._instance is None:
+            cls._instance = super().__new__(cls)
+            cls._instance._comm_ptr = None
+            cls._instance._mem_ptr = None
+        return cls._instance
+
+    def __init__(self, _comm_ptr=None, _mem_ptr=None):
+        if _comm_ptr and _mem_ptr:
+            type(self)._init_count += 1
+
+        if self._initialized:
+            return
+        self._comm_ptr = _comm_ptr
+        self._mem_ptr = _mem_ptr
+        self._initialized = True
+
+    def get_packed_data(self):
+        return int(self._mem_ptr), int(self._comm_ptr)
+
+    @property
+    def comm_ptr(self) -> int | None:
+        """Communication runtime pointer."""
+        return self._comm_ptr
+
+    def _get_needed_params(self):
+        return {"device_comm_ptr": self._comm_ptr, "device_mem_ptr": self._mem_ptr}
+
+    @property
+    def mem_ptr(self) -> int | None:
+        """Distributed memory runtime pointer."""
+        return self._mem_ptr
+
+    @property
+    def is_lite_mode(self) -> bool:
+        import os
+        user_action = os.getenv("FLAGTREE_LITE_DIST", "").strip().upper() in {
+            "1",
+            "ON",
+            "TRUE",
+        }
+        inner_action = self._init_count == 1
+        return inner_action and user_action
+
+    def __getitem__(self, index=0):
+        return list(self._get_needed_params().values())[index]
+
+    def add_args_to_jitfunction(self, **kwargs):
+        params = kwargs['params']  # list
+        _kwargs = kwargs['kwargs']
+        if (isinstance(params, list) and len(params) > 0):
+            needed_params = self._get_needed_params()
+            needed_params_size = len(needed_params)
+            template_ele = params[0]
+            KernelParam = type(template_ele)  # KernelParam type
+            Parameter = type(template_ele._param)  # inspect.Parameter type
+
+            dist_params = []
+            for i, (name, val) in enumerate(needed_params.items()):
+                _kwargs[name] = val
+                param = Parameter(name, kind=template_ele._param.kind)
+                dist_params.append(KernelParam(i, param, False, False))
+
+            new_params = []
+            for param in params:
+                new_loc = param.num + needed_params_size
+                #num: int, param: inspect.Parameter, do_not_specialize: bool, do_not_specialize_on_alignment: bool
+                new_params.append(
+                    KernelParam(new_loc, param._param, param.do_not_specialize, param.do_not_specialize_on_alignment))
+            new_params = dist_params + new_params
+            params[:] = new_params
diff --git a/python/triton/runtime/jit.py b/python/triton/runtime/jit.py
index b45a3943c..d2e92b2a5 100644
--- a/python/triton/runtime/jit.py
+++ b/python/triton/runtime/jit.py
@@ -40,6 +40,7 @@
 from .._utils import find_paths_if, get_iterable_path, type_canonicalisation_dict, is_namedtuple
 from .cache import get_cache_key
 from triton._C.libtriton import get_cache_invalidating_env_vars, native_specialize_impl
+from ._distributed import DistributedRtContext
 
 TRITON_MODULE = "triton.language"
 GLUON_MODULE = "triton.experimental.gluon.language"
@@ -781,8 +782,15 @@ def run(self, *args, grid, warmup, **kwargs):
                 kernel = kernel.result()
             # launch kernel
             launch_metadata = kernel.launch_metadata(grid, stream, *bound_args.values())
-            kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,
-                       knobs.runtime.launch_enter_hook, knobs.runtime.launch_exit_hook, *bound_args.values())
+            # flagtree tle distributed: Add dist_param to kernel.run
+            dist_param = []
+            ctx = DistributedRtContext()
+            if ctx.is_lite_mode:
+                dist_param += [ctx.comm_ptr, ctx.mem_ptr]
+            kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata,
+                       launch_metadata, knobs.runtime.launch_enter_hook, knobs.runtime.launch_exit_hook, *dist_param,
+                       *bound_args.values())
+
         return kernel
 
     def repr(self, _):
diff --git a/setup.py b/setup.py
index ad3655f81..18a294658 100644
--- a/setup.py
+++ b/setup.py
@@ -698,9 +698,8 @@ def download_and_copy_dependencies():
 
 if helper.flagtree_backend:
     if helper.flagtree_backend in ("aipu", "tsingmicro", "enflame", "rpu", "thrive", "sunrise", "tileir"):
-        default_backends = helper.configs.non_tileir_default_backends()
         backends = [
-            *BackendInstaller.copy(default_backends + tuple(helper.configs.extend_backends)),
+            *BackendInstaller.copy(helper.configs.default_backends + tuple(helper.configs.extend_backends)),
             *BackendInstaller.copy_externals(),
         ]
     else:
diff --git a/third_party/nvidia/backend/compiler.py b/third_party/nvidia/backend/compiler.py
index 8fee79a6e..64ae02436 100644
--- a/third_party/nvidia/backend/compiler.py
+++ b/third_party/nvidia/backend/compiler.py
@@ -3,6 +3,7 @@
 from triton._C.libtriton import tle
 from triton import knobs
 from triton.runtime.errors import PTXASError
+from triton.runtime._distributed import DistributedRtContext
 
 from dataclasses import dataclass
 import functools
@@ -279,6 +280,9 @@ def make_ttgir(mod, metadata, opt, capability):
         pm = ir.pass_manager(mod.context)
         dump_enabled = pm.enable_debug()
         emuTF32 = (capability // 10 >= 8)
+        # flagtree tle distributed
+        if DistributedRtContext().is_lite_mode:
+            tle.passes.add_params_for_distribution(pm)
         passes.ttir.add_convert_to_ttgpuir(pm, f"cuda:{capability}", opt.num_warps, 32, opt.num_ctas)
         # flagtree tle raw
         tle.raw_passes.add_tle_convert_arg_to_memdesc(pm)
diff --git a/third_party/nvidia/backend/distributed.py b/third_party/nvidia/backend/distributed.py
index 798500573..159a4d902 100644
--- a/third_party/nvidia/backend/distributed.py
+++ b/third_party/nvidia/backend/distributed.py
@@ -69,8 +69,8 @@ def __init__(self, path_order=0):
             self.include_path = self._get_include_paths()[path_order]
 
     def _check_path_available(self, paths):
-        available_paths = [Path(p) for p in paths if p and p.exists()]
-        if len(paths) == 0:
+        available_paths = [Path(p) for p in paths if p and Path(p).exists()]
+        if len(available_paths) == 0:
             raise RuntimeError(f"There are no available {self.bt_name} path in this {available_paths}")
         return available_paths
 
diff --git a/third_party/nvidia/backend/driver.py b/third_party/nvidia/backend/driver.py
index 8e9c97acf..0c5e98f56 100644
--- a/third_party/nvidia/backend/driver.py
+++ b/third_party/nvidia/backend/driver.py
@@ -9,6 +9,7 @@
 from triton.runtime import _allocation
 from triton.backends.compiler import GPUTarget
 from triton.backends.driver import GPUDriver
+from triton.runtime._distributed import DistributedRtContext
 
 dirname = os.path.dirname(os.path.realpath(__file__))
 include_dirs = [os.path.join(dirname, "include")]
@@ -215,6 +216,10 @@ def format_of(ty):
     flat_signature = []
     for sig in signature.values():
         _flatten_signature(sig, flat_signature)
+    # flagtree tle distributed
+    if DistributedRtContext().is_lite_mode:
+        flat_signature.insert(0, "*i64")  # flagcx_dev_comm_ptr
+        flat_signature.insert(1, "*i64")  # flagcx_dev_comm_ptr
     signature = {i: s for i, s in enumerate(flat_signature)}
     args_list = ', ' + ', '.join(f"&_arg{i}" for i, ty in signature.items()) if len(signature) > 0 else ''
     # Record the end of regular arguments;
diff --git a/third_party/nvidia/lib/TritonNVIDIAGPUToLLVM/CMakeLists.txt b/third_party/nvidia/lib/TritonNVIDIAGPUToLLVM/CMakeLists.txt
index 23137749c..3bc3be2bf 100644
--- a/third_party/nvidia/lib/TritonNVIDIAGPUToLLVM/CMakeLists.txt
+++ b/third_party/nvidia/lib/TritonNVIDIAGPUToLLVM/CMakeLists.txt
@@ -4,6 +4,9 @@ if(FLAGTREE_TLE)
         TleToLLVM # flagtree tle raw
         TritonTLETransforms
     )
+    if(FLAGCX_ENABLED)
+        list(APPEND _TLE_LIBS FlagCxToLLVM)
+    endif()
 else()
     set(_TLE_LIBS "")
 endif()
diff --git a/third_party/nvidia/lib/TritonNVIDIAGPUToLLVM/TritonGPUToLLVM.cpp b/third_party/nvidia/lib/TritonNVIDIAGPUToLLVM/TritonGPUToLLVM.cpp
index b3c9b193c..552dc79ed 100644
--- a/third_party/nvidia/lib/TritonNVIDIAGPUToLLVM/TritonGPUToLLVM.cpp
+++ b/third_party/nvidia/lib/TritonNVIDIAGPUToLLVM/TritonGPUToLLVM.cpp
@@ -25,8 +25,8 @@
 #include "tle/dialect/include/Conversion/TleToLLVM/DistributedBarrierOpToLLVM.h"
 #include "tle/dialect/include/Conversion/TleToLLVM/ExclusiveCumsumOpToLLVM.h"
 #include "tle/dialect/include/Conversion/TleToLLVM/ExtractOpToLLVM.h"
+#include "tle/dialect/include/Conversion/TleToLLVM/FlagCxOpToLLVM/FlagCxOpToLLVM.h"
 #include "tle/dialect/include/Conversion/TleToLLVM/GetDeviceIdToFlagCX.h"
-#include "tle/dialect/include/Conversion/TleToLLVM/GetLocalRankOpToLLVM.h"
 #include "tle/dialect/include/Conversion/TleToLLVM/LocalPointersOpToLLVM.h"
 #include "tle/dialect/include/Conversion/TleToLLVM/PackOpToLLVM.h"
 #include "tle/dialect/include/IR/Dialect.h"
@@ -112,10 +112,6 @@ class TleLLVMConversionTarget : public ConversionTarget {
           return hasLegalRegions && typeConverter.isLegal(op);
         });
     addLegalOp<tle::RemotePointersOp>();
-    addLegalOp<tle::GetDeviceIdOp>();
-    // addIllegalOp<tle::GetLocalRankOp>();
-    // // addLegalOp<tle::GetDeviceIdOp>();
-    // addLegalOp<tle::GetNumPesOp>();
     // Allow non-TLE ops to remain during this partial conversion.
     markUnknownOpDynamicallyLegal([](Operation *) -> bool { return true; });
   }
@@ -180,8 +176,6 @@ struct ConvertTritonGPUToLLVM
                                                       benefit);
       mlir::triton::tle::populateDistributedBarrierOpToLLVMPatterns(
           typeConverter, patterns, benefit);
-      // mlir::triton::tle::populateGetNumPesOpToLLVMPatterns(
-      //     typeConverter, patterns, benefit + 1);
       mlir::triton::tle::populateLocalPointersOpToLLVMPatterns(
           typeConverter, targetInfo, patterns, benefit);
       mlir::triton::tle::populateExtractTileOpToLLVMPatterns(
@@ -196,21 +190,15 @@ struct ConvertTritonGPUToLLVM
           typeConverter, patterns, benefit);
       mlir::triton::tle::populateTMAStoreCommitGroupOpToLLVMPatterns(
           typeConverter, patterns, benefit);
+      // FlagCX ops are lowered to LLVM.
+#ifdef FLAGCX_ENABLED
+      mlir::triton::tle::populateFlagCxOpToLLVMPatterns(typeConverter, patterns,
+                                                        benefit);
+#endif
       if (failed(applyPartialConversion(mod, target, std::move(patterns)))) {
         return signalPassFailure();
       }
     }
-#ifdef FLAGCX_ENABLED
-    {
-      mlir::triton::tle::populateGetDeviceIdOpToFlagCxPatterns(
-          typeConverter, patterns, benefit);
-      mlir::triton::tle::populateGetLocalRankOpToLLVMPatterns(
-          typeConverter, patterns, benefit);
-      mlir::triton::tle::populateGetNumPesOpToLLVMPatterns(typeConverter,
-                                                           patterns, benefit);
-    }
-#endif
-
 #endif
 
     mlir::triton::NVIDIA::populateConvertLayoutOpToLLVMPatterns(
diff --git a/third_party/tileir/CMakeLists.txt b/third_party/tileir/CMakeLists.txt
index 96ce81423..de8277ecf 100644
--- a/third_party/tileir/CMakeLists.txt
+++ b/third_party/tileir/CMakeLists.txt
@@ -40,8 +40,7 @@ else()
   set(CUDA_TILE_SUBMODULE_DIR ${CMAKE_CURRENT_SOURCE_DIR}/third_party/cuda-tile)
   if(NOT EXISTS ${CUDA_TILE_SUBMODULE_DIR}/CMakeLists.txt)
     message(FATAL_ERROR
-      "cuda-tile submodule not initialized at ${CUDA_TILE_SUBMODULE_DIR}. "
-      "Run: git submodule update --init --recursive")
+      "cuda-tile submodule not initialized at ${CUDA_TILE_SUBMODULE_DIR}. ")
   endif()
 
   execute_process(
diff --git a/third_party/tileir/README.md b/third_party/tileir/README.md
index e6ed4d907..72fdf89a9 100644
--- a/third_party/tileir/README.md
+++ b/third_party/tileir/README.md
@@ -35,13 +35,7 @@ FlagTree's normal NVIDIA assemblers, managed by `setup.py` at:
 - `third_party/nvidia/backend/bin/ptxas-blackwell`
 
 `pip install` and `pip install -e` initialize the pinned cuda-tile submodule
-automatically through `setup.py`; no manual submodule command is needed. If you
-invoke CMake directly without `setup.py`, initialize the submodule first:
-
-```bash
-git submodule update --init --recursive \
-  third_party/tileir/third_party/cuda-tile
-```
+automatically through `setup.py`; no manual submodule command is needed.
 
 At runtime, TileIR looks for `tileiras` in this order:
 
diff --git a/third_party/tileir/third_party/cuda-tile b/third_party/tileir/third_party/cuda-tile
deleted file mode 160000
index 2e5ccba66..000000000
--- a/third_party/tileir/third_party/cuda-tile
+++ /dev/null
@@ -1 +0,0 @@
-Subproject commit 2e5ccba66fb3afdba34b26cf358418283027c248
diff --git a/third_party/tle/dialect/include/Conversion/TleToLLVM/FlagCxOpToLLVM/DeviceIntraBarrierOpToLLVM.h b/third_party/tle/dialect/include/Conversion/TleToLLVM/FlagCxOpToLLVM/DeviceIntraBarrierOpToLLVM.h
new file mode 100644
index 000000000..fe002abed
--- /dev/null
+++ b/third_party/tle/dialect/include/Conversion/TleToLLVM/FlagCxOpToLLVM/DeviceIntraBarrierOpToLLVM.h
@@ -0,0 +1,15 @@
+#ifndef TLE_CONVERSION_TLE_TO_LLVM_DEVICE_INTRA_BARRIER_OP_TO_LLVM_H
+#define TLE_CONVERSION_TLE_TO_LLVM_DEVICE_INTRA_BARRIER_OP_TO_LLVM_H
+
+#include "mlir/Conversion/LLVMCommon/TypeConverter.h"
+#include "mlir/IR/PatternMatch.h"
+
+namespace mlir::triton::tle {
+
+void populateDeviceIntraBarrierOpToLLVMPatterns(
+    LLVMTypeConverter &typeConverter, RewritePatternSet &patterns,
+    PatternBenefit benefit);
+
+} // namespace mlir::triton::tle
+
+#endif // TLE_CONVERSION_TLE_TO_LLVM_DEVICE_INTRA_BARRIER_OP_TO_LLVM_H
diff --git a/third_party/tle/dialect/include/Conversion/TleToLLVM/FlagCxOpToLLVM/FlagCxOpToLLVM.h b/third_party/tle/dialect/include/Conversion/TleToLLVM/FlagCxOpToLLVM/FlagCxOpToLLVM.h
new file mode 100644
index 000000000..e2704e8e5
--- /dev/null
+++ b/third_party/tle/dialect/include/Conversion/TleToLLVM/FlagCxOpToLLVM/FlagCxOpToLLVM.h
@@ -0,0 +1,8 @@
+#include "tle/dialect/include/Conversion/TleToLLVM/FlagCxOpToLLVM/DeviceIntraBarrierOpToLLVM.h"
+#include "tle/dialect/include/Conversion/TleToLLVM/FlagCxOpToLLVM/GetLocalRankOpToLLVM.h"
+
+namespace mlir::triton::tle {
+void populateFlagCxOpToLLVMPatterns(LLVMTypeConverter &typeConverter,
+                                    RewritePatternSet &patterns,
+                                    PatternBenefit benefit);
+}
diff --git a/third_party/tle/dialect/include/Conversion/TleToLLVM/GetLocalRankOpToLLVM.h b/third_party/tle/dialect/include/Conversion/TleToLLVM/FlagCxOpToLLVM/GetLocalRankOpToLLVM.h
similarity index 100%
rename from third_party/tle/dialect/include/Conversion/TleToLLVM/GetLocalRankOpToLLVM.h
rename to third_party/tle/dialect/include/Conversion/TleToLLVM/FlagCxOpToLLVM/GetLocalRankOpToLLVM.h
diff --git a/third_party/tle/dialect/include/IR/FlagCxOps.td b/third_party/tle/dialect/include/IR/FlagCxOps.td
index 7c283d78b..e882c8d7f 100644
--- a/third_party/tle/dialect/include/IR/FlagCxOps.td
+++ b/third_party/tle/dialect/include/IR/FlagCxOps.td
@@ -23,3 +23,36 @@ def Tle_GetLocalRankOp : Tle_Op<"get local rank"> {
   let results = (outs I32:$result);
   let hasVerifier = 1;
 }
+
+
+def Tle_DeviceIntraBarrierOp : Tle_Op<"device intra barrier"> {
+  let summary = "Device-local synchronization barrier";
+  let description = [{
+    Performs a device-local synchronization barrier by lowering to one of the
+    following runtime calls:
+
+      * flagcxIntraBarrierArriveS(comm, coopKind, index, multimem, order)
+      * flagcxIntraBarrierWaitS(comm, coopKind, index, multimem, order)
+      * flagcxIntraBarrierSyncS(comm, coopKind, index, multimem, order)
+
+    The specific runtime routine is selected according to the barrier operation
+    (arrive, wait, or sync).
+
+    The barrier is identified by `index`, scoped by `coopKind`, and obeys the
+    specified memory ordering semantics (`order`). `multimem` indicates whether
+    the barrier operates on multimem storage.
+
+    Returns the runtime status code.
+  }];
+
+  let arguments = (ins
+    Tle_LocalPointerResultType:$comm,
+    OptionalAttr<StrAttr>:$barrier_type,
+    OptionalAttr<I32Attr>:$coop_kind,
+    OptionalAttr<I32Attr>:$index,
+    OptionalAttr<I1Attr>:$multimem,
+    OptionalAttr<I32Attr>:$order
+  );
+
+  let hasVerifier = 1;
+}
diff --git a/third_party/tle/dialect/include/IR/TleOps.td b/third_party/tle/dialect/include/IR/TleOps.td
index 3fffa4d47..c72884494 100644
--- a/third_party/tle/dialect/include/IR/TleOps.td
+++ b/third_party/tle/dialect/include/IR/TleOps.td
@@ -319,19 +319,23 @@ def Tle_TMAStoreCommitGroupOp
 def Tle_DistributedBarrierOp : Tle_Op<"distributed_barrier",
                                       [MemoryEffects<[MemRead, MemWrite]>]> {
   let arguments = (ins
+    Optional<AnyType>:$src,
+    OptionalAttr<StrAttr>:$space,
+    OptionalAttr<StrAttr>:$barrier_type,
+    OptionalAttr<StrAttr>:$order,
     OptionalAttr<StrAttr>:$group_kind,
+    OptionalAttr<I32Attr>:$barrier_index,
     OptionalAttr<I32Attr>:$group_rank,
     OptionalAttr<DenseI32ArrayAttr>:$group_shape,
     OptionalAttr<DenseI32ArrayAttr>:$group_axes,
     OptionalAttr<DenseI32ArrayAttr>:$group_mask
   );
-  let assemblyFormat = "attr-dict";
   let hasVerifier = 1;
 }
 
-def Tle_RemotePointersOp : Tle_Op<"remote_pointers", [Pure]> {
+def Tle_RemotePointersOp : Tle_Op<"remote_pointers", [Pure, AttrSizedOperandSegments]> {
   let arguments = (ins
-    Tle_LocalPointerResultType:$src,
+    Optional<Tle_LocalPointerResultType>:$src,
     TT_Int:$shard_id,
     StrAttr:$space,
     Optional<TT_IntLike>:$offset
@@ -352,12 +356,13 @@ def Tle_GetNumPesOp : Tle_Op<"get_num_pes"> {
 def Tle_GetDeviceIdOp : Tle_Op<"get_device_id"> {
   let summary = "Get local PE id";
   let arguments = (ins
-     AnyType:$input
+     Optional<AnyType>:$input
   );
   let results = (outs I32:$result);
   let hasVerifier = 1;
 }
 
+
 def Tle_DSLRegionOp : Tle_Op<"dsl_region", [IsolatedFromAbove, MemDescViewTrait,
                                             DeclareOpInterfaceMethods<MemoryEffectsOpInterface>]> {
   let arguments = (ins
diff --git a/third_party/tle/dialect/include/IR/VerfiyUtils.h b/third_party/tle/dialect/include/IR/VerfiyUtils.h
index 24d2e45e8..bdbf2fdf2 100644
--- a/third_party/tle/dialect/include/IR/VerfiyUtils.h
+++ b/third_party/tle/dialect/include/IR/VerfiyUtils.h
@@ -17,4 +17,8 @@ namespace RemotePointers {
 llvm::LogicalResult verifyDeviceSpace(mlir::Value src, mlir::Value result);
 }
 
+namespace DistributedBarrier {
+llvm::LogicalResult verifyDeviceSpace(mlir::Operation *op, mlir::Value src);
+}
+
 } // namespace mlir::triton::tle
diff --git a/third_party/tle/dialect/include/Tools/FlagcxUtils.h b/third_party/tle/dialect/include/Tools/FlagcxUtils.h
index 84b5ecbfa..144f0add8 100644
--- a/third_party/tle/dialect/include/Tools/FlagcxUtils.h
+++ b/third_party/tle/dialect/include/Tools/FlagcxUtils.h
@@ -16,4 +16,9 @@ LLVM::CallOp getNumPesFunCall(mlir::Location loc,
                               ConversionPatternRewriter &rewriter,
                               Value memPtrInt);
 
+LLVM::CallOp getBarrierFuncCall(mlir::Location loc,
+                                ConversionPatternRewriter &rewriter, Value comm,
+                                size_t barrier_index, size_t coopKind,
+                                size_t order, llvm::StringRef barrierType);
+
 } // namespace mlir::triton::tle
diff --git a/third_party/tle/dialect/include/Transforms/Passes.td b/third_party/tle/dialect/include/Transforms/Passes.td
index 2ddd3b1bf..50523383b 100644
--- a/third_party/tle/dialect/include/Transforms/Passes.td
+++ b/third_party/tle/dialect/include/Transforms/Passes.td
@@ -397,4 +397,28 @@ def TleRemoveRedundantCopy
 
 def TleDSLRegionInline : Pass<"tle-dslregion-inline", "mlir::ModuleOp"> {}
 
+def TritonTleAddDistributedParams
+    : Pass<"triton-tle-add-distributed-params", "mlir::ModuleOp"> {
+
+  let summary = "Add distributed runtime parameters to Triton kernel functions";
+
+  let description = [{
+    This pass updates Triton kernel function signatures by inserting two
+    runtime parameters at the beginning of each kernel function:
+
+      * device_comm_ptr
+      * device_mem_ptr
+
+    These parameters provide device-side communication and runtime state
+    required by distributed TLE operations.
+
+    All call sites and function argument uses are updated accordingly.
+  }];
+
+  let dependentDialects = [
+    "mlir::triton::TritonDialect"
+  ];
+}
+
+
 #endif // TRITON_TLE_PASSES
diff --git a/third_party/tle/dialect/lib/Conversion/TleToLLVM/CMakeLists.txt b/third_party/tle/dialect/lib/Conversion/TleToLLVM/CMakeLists.txt
index 0a634494a..d59aef3a5 100644
--- a/third_party/tle/dialect/lib/Conversion/TleToLLVM/CMakeLists.txt
+++ b/third_party/tle/dialect/lib/Conversion/TleToLLVM/CMakeLists.txt
@@ -11,10 +11,11 @@ set(TLE_TO_LLVM_SRC
     WGMMAFenceOpToLLVM.cpp
 )
 
+
 if(FLAGCX_ENABLED)
- list(APPEND TLE_TO_LLVM_SRC
+add_subdirectory(FlagCxOpToLLVM)
+list(APPEND TLE_TO_LLVM_SRC
   GetDeviceIdToFlagCX.cpp
-  GetLocalRankOpToLLVM.cpp
 )
 endif()
 
diff --git a/third_party/tle/dialect/lib/Conversion/TleToLLVM/DistributedBarrierOpToLLVM.cpp b/third_party/tle/dialect/lib/Conversion/TleToLLVM/DistributedBarrierOpToLLVM.cpp
index 14a9950e1..5bd5ec7c9 100644
--- a/third_party/tle/dialect/lib/Conversion/TleToLLVM/DistributedBarrierOpToLLVM.cpp
+++ b/third_party/tle/dialect/lib/Conversion/TleToLLVM/DistributedBarrierOpToLLVM.cpp
@@ -1,4 +1,5 @@
 #include "tle/dialect/include/Conversion/TleToLLVM/DistributedBarrierOpToLLVM.h"
+#include "tle/dialect/include/Tools/FlagcxUtils.h"
 
 #include "mlir/Conversion/LLVMCommon/Pattern.h"
 #include "mlir/Dialect/GPU/IR/GPUDialect.h"
@@ -18,6 +19,9 @@ namespace {
 using namespace mlir;
 namespace tle = mlir::triton::tle;
 
+constexpr llvm::StringLiteral kSpaceAttr = "space";
+constexpr llvm::StringLiteral kOrderAttr = "order";
+constexpr llvm::StringLiteral kIndexAttr = "barrier_index";
 constexpr llvm::StringLiteral kGroupKindAttr = "group_kind";
 constexpr llvm::StringLiteral kGroupShapeAttr = "group_shape";
 constexpr llvm::StringLiteral kGroupMaskAttr = "group_mask";
@@ -38,6 +42,16 @@ constexpr int32_t kGridScratchAlignment = 4;
 constexpr int32_t kGridScratchBytes = 4;
 constexpr int32_t kGridArrivedOffsetBytes = 0;
 
+Value getDistDevicePtr(tle::DistributedBarrierOp op,
+                       SmallVector<Value> &srcElems) {
+  if (!srcElems.empty())
+    return srcElems[0];
+  else {
+    auto func = op->getParentOfType<LLVM::LLVMFuncOp>();
+    return func.getArgument(1);
+  }
+}
+
 FailureOr<int32_t> getOrCreateSubmeshScratchOffset(ModuleOp mod) {
   if (auto existing =
           mod->getAttrOfType<IntegerAttr>(kSubmeshScratchOffsetAttr)) {
@@ -422,9 +436,61 @@ struct DistributedBarrierOpConversion
     return success();
   }
 
+  LogicalResult
+  lowerDeviceSpaceBarrier(tle::DistributedBarrierOp op, OpAdaptor adaptor,
+                          ConversionPatternRewriter &rewriter) const {
+    auto kindAttr = op->getAttrOfType<StringAttr>(kGroupKindAttr);
+    auto orderAttr = op->getAttrOfType<StringAttr>(kOrderAttr);
+    auto indexAttr = op->getAttrOfType<IntegerAttr>(kIndexAttr);
+    auto loc = op.getLoc();
+    SmallVector<Value> srcElems;
+    auto getCoopKindValue = [](StringRef kind) -> int32_t {
+      return llvm::StringSwitch<int32_t>(kind)
+          .Case("thread", 0)
+          .Case("warp", 1)
+          .Case("block", 2)
+          .Case("grid", 3)
+          .Default(-1);
+    };
+    auto getOrderValue = [](StringRef order) -> int32_t {
+      return llvm::StringSwitch<int32_t>(order)
+          .Case("relaxed", 0)
+          .Case("acquire", 1)
+          .Case("release", 2)
+          .Case("acqrel", 3)
+          .Default(-1);
+    };
+
+    int32_t coopKind = getCoopKindValue(kindAttr.getValue());
+    int32_t order = getOrderValue(orderAttr.getValue());
+    if (coopKind < 0)
+      return rewriter.notifyMatchFailure(op, "invalid coop_kind");
+
+    if (order < 0)
+      return rewriter.notifyMatchFailure(op, "invalid order");
+
+    if (auto src = adaptor.getSrc())
+      srcElems = unpackLLElements(loc, src, rewriter);
+
+    auto comm = getDistDevicePtr(op, srcElems);
+    auto coopKindAttr = rewriter.getI32IntegerAttr(coopKind);
+    auto newOrderAttr = rewriter.getI32IntegerAttr(order);
+    auto barrierTypeAttr = op.getBarrierTypeAttr();
+    auto multimemAttr = rewriter.getBoolAttr(false);
+#ifdef FLAGCX_ENABLED
+    rewriter.replaceOpWithNewOp<tle::DeviceIntraBarrierOp>(
+        op, comm, barrierTypeAttr, coopKindAttr, indexAttr, multimemAttr,
+        newOrderAttr);
+#endif
+    return success();
+  }
   LogicalResult
   matchAndRewrite(tle::DistributedBarrierOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
+    if (auto spaceAttr = op->getAttrOfType<StringAttr>(kSpaceAttr))
+      if (spaceAttr.getValue() == "device")
+        return lowerDeviceSpaceBarrier(op, adaptor, rewriter);
+
     if (auto kindAttr = op->getAttrOfType<StringAttr>(kGroupKindAttr)) {
       if (kindAttr.getValue() == "grid")
         return lowerGridBarrier(op, rewriter);
diff --git a/third_party/tle/dialect/lib/Conversion/TleToLLVM/FlagCxOpToLLVM/CMakeLists.txt b/third_party/tle/dialect/lib/Conversion/TleToLLVM/FlagCxOpToLLVM/CMakeLists.txt
new file mode 100644
index 000000000..124815cb3
--- /dev/null
+++ b/third_party/tle/dialect/lib/Conversion/TleToLLVM/FlagCxOpToLLVM/CMakeLists.txt
@@ -0,0 +1,17 @@
+set(FLAGCX_TO_LLVM_SRC
+    DeviceIntraBarrierOpToLLVM.cpp
+    GetLocalRankOpToLLVM.cpp
+    FlagCxOpToLLVM.cpp
+)
+
+add_triton_library(FlagCxToLLVM
+    ${FLAGCX_TO_LLVM_SRC}
+
+    DEPENDS
+    TleTableGen
+    TritonGPUTypeInterfacesIncGen
+    TleTools
+
+    LINK_LIBS PUBLIC
+    TleTools
+)
diff --git a/third_party/tle/dialect/lib/Conversion/TleToLLVM/FlagCxOpToLLVM/DeviceIntraBarrierOpToLLVM.cpp b/third_party/tle/dialect/lib/Conversion/TleToLLVM/FlagCxOpToLLVM/DeviceIntraBarrierOpToLLVM.cpp
new file mode 100644
index 000000000..972dd07c6
--- /dev/null
+++ b/third_party/tle/dialect/lib/Conversion/TleToLLVM/FlagCxOpToLLVM/DeviceIntraBarrierOpToLLVM.cpp
@@ -0,0 +1,57 @@
+#include "tle/dialect/include/Conversion/TleToLLVM/FlagCxOpToLLVM/DeviceIntraBarrierOpToLLVM.h"
+#include "tle/dialect/include/Tools/FlagcxUtils.h"
+
+#include "mlir/Conversion/LLVMCommon/Pattern.h"
+#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
+#include "mlir/Dialect/LLVMIR/LLVMTypes.h"
+#include "mlir/Dialect/LLVMIR/NVVMDialect.h"
+#include "mlir/IR/BuiltinTypes.h"
+#include "mlir/Transforms/DialectConversion.h"
+#include "tle/dialect/include/IR/Dialect.h"
+#include "triton/Conversion/TritonGPUToLLVM/Utility.h"
+#include "triton/Dialect/Triton/IR/Types.h"
+#include "triton/Dialect/TritonGPU/IR/Dialect.h"
+#include "triton/Dialect/TritonGPU/IR/LinearLayoutConversions.h"
+#include "triton/Tools/LayoutUtils.h"
+#include "llvm/Support/raw_ostream.h"
+
+namespace {
+using namespace mlir;
+namespace ttg = mlir::triton::gpu;
+namespace tle = mlir::triton::tle;
+
+struct DeviceIntraBarrierOpConversion
+    : public ConvertOpToLLVMPattern<tle::DeviceIntraBarrierOp> {
+  DeviceIntraBarrierOpConversion(LLVMTypeConverter &typeConverter,
+                                 PatternBenefit benefit)
+      : ConvertOpToLLVMPattern(typeConverter, benefit) {}
+
+  LogicalResult
+  matchAndRewrite(tle::DeviceIntraBarrierOp op, OpAdaptor adaptor,
+                  ConversionPatternRewriter &rewriter) const override {
+    auto loc = op.getLoc();
+
+    auto indexValue = op.getIndexAttr().getInt();
+    auto coopValue = op.getCoopKindAttr().getInt();
+    auto orderValue = op.getOrderAttr().getInt();
+    auto barrierType = op.getBarrierTypeAttr().getValue();
+    if (!llvm::is_contained(std::array<size_t, 5>{0, 1, 2, 3, 4}, coopValue))
+      return rewriter.notifyMatchFailure(op, "invalid coop_kind");
+
+    if (!llvm::is_contained(std::array<size_t, 4>{0, 1, 2, 3}, orderValue))
+      return rewriter.notifyMatchFailure(op, "invalid coop_kind");
+
+    tle::getBarrierFuncCall(loc, rewriter, adaptor.getComm(), indexValue,
+                            coopValue, orderValue, barrierType);
+    rewriter.eraseOp(op);
+    return success();
+  }
+};
+
+} // namespace
+
+void tle::populateDeviceIntraBarrierOpToLLVMPatterns(
+    LLVMTypeConverter &typeConverter, RewritePatternSet &patterns,
+    PatternBenefit benefit) {
+  patterns.add<DeviceIntraBarrierOpConversion>(typeConverter, benefit);
+}
diff --git a/third_party/tle/dialect/lib/Conversion/TleToLLVM/FlagCxOpToLLVM/FlagCxOpToLLVM.cpp b/third_party/tle/dialect/lib/Conversion/TleToLLVM/FlagCxOpToLLVM/FlagCxOpToLLVM.cpp
new file mode 100644
index 000000000..214bcc590
--- /dev/null
+++ b/third_party/tle/dialect/lib/Conversion/TleToLLVM/FlagCxOpToLLVM/FlagCxOpToLLVM.cpp
@@ -0,0 +1,21 @@
+#include "tle/dialect/include/Conversion/TleToLLVM/FlagCxOpToLLVM/DeviceIntraBarrierOpToLLVM.h"
+#include "tle/dialect/include/Conversion/TleToLLVM/FlagCxOpToLLVM/GetLocalRankOpToLLVM.h"
+#include "tle/dialect/include/Conversion/TleToLLVM/GetDeviceIdToFlagCX.h"
+
+namespace mlir::triton::tle {
+void populateFlagCxOpToLLVMPatterns(LLVMTypeConverter &typeConverter,
+                                    RewritePatternSet &patterns,
+                                    PatternBenefit benefit) {
+#ifdef FLAGCX_ENABLED
+  mlir::triton::tle::populateGetDeviceIdOpToFlagCxPatterns(typeConverter,
+                                                           patterns, benefit);
+  mlir::triton::tle::populateGetLocalRankOpToLLVMPatterns(typeConverter,
+                                                          patterns, benefit);
+  mlir::triton::tle::populateGetNumPesOpToLLVMPatterns(typeConverter, patterns,
+                                                       benefit);
+  mlir::triton::tle::populateDeviceIntraBarrierOpToLLVMPatterns(
+      typeConverter, patterns, benefit);
+#endif
+}
+
+} // namespace mlir::triton::tle
diff --git a/third_party/tle/dialect/lib/Conversion/TleToLLVM/GetLocalRankOpToLLVM.cpp b/third_party/tle/dialect/lib/Conversion/TleToLLVM/FlagCxOpToLLVM/GetLocalRankOpToLLVM.cpp
similarity index 95%
rename from third_party/tle/dialect/lib/Conversion/TleToLLVM/GetLocalRankOpToLLVM.cpp
rename to third_party/tle/dialect/lib/Conversion/TleToLLVM/FlagCxOpToLLVM/GetLocalRankOpToLLVM.cpp
index 9a63b05f9..ae793adb7 100644
--- a/third_party/tle/dialect/lib/Conversion/TleToLLVM/GetLocalRankOpToLLVM.cpp
+++ b/third_party/tle/dialect/lib/Conversion/TleToLLVM/FlagCxOpToLLVM/GetLocalRankOpToLLVM.cpp
@@ -1,4 +1,4 @@
-#include "tle/dialect/include/Conversion/TleToLLVM/GetLocalRankOpToLLVM.h"
+#include "tle/dialect/include/Conversion/TleToLLVM/FlagCxOpToLLVM/GetLocalRankOpToLLVM.h"
 #include "tle/dialect/include/Tools/FlagcxUtils.h"
 
 #include "mlir/Conversion/LLVMCommon/Pattern.h"
@@ -53,20 +53,18 @@ struct GetLocalRankOpConversion
   LogicalResult
   matchAndRewrite(tle::GetLocalRankOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
-
     auto reportFailure = [&](StringRef msg) -> LogicalResult {
       llvm::errs() << "[GetLocalRankOpConversion] " << msg << "\n";
       return failure();
     };
     auto loc = op.getLoc();
-    auto srcElems = unpackLLElements(loc, adaptor.getSrc(), rewriter);
-    auto getLocalPeCall = tle::getLocalPeFuncCall(loc, rewriter, srcElems[0]);
+    auto comm = op.getSrc();
+    auto getLocalPeCall = tle::getLocalPeFuncCall(loc, rewriter, comm);
 
     Value localPe = getLocalPeCall.getResult();
     if (!localPe.getType().isInteger(32))
       return reportFailure("expected i32 result");
     rewriter.replaceOp(op, localPe);
-
     return success();
   }
 };
diff --git a/third_party/tle/dialect/lib/Conversion/TleToLLVM/GetDeviceIdToFlagCX.cpp b/third_party/tle/dialect/lib/Conversion/TleToLLVM/GetDeviceIdToFlagCX.cpp
index 79061205e..dff002aea 100644
--- a/third_party/tle/dialect/lib/Conversion/TleToLLVM/GetDeviceIdToFlagCX.cpp
+++ b/third_party/tle/dialect/lib/Conversion/TleToLLVM/GetDeviceIdToFlagCX.cpp
@@ -20,6 +20,15 @@ using namespace mlir;
 namespace ttg = mlir::triton::gpu;
 namespace tle = mlir::triton::tle;
 
+Value getDistDevicePtr(tle::GetDeviceIdOp op, SmallVector<Value> &srcElems) {
+  if (!srcElems.empty())
+    return srcElems[0];
+  else {
+    auto func = op->getParentOfType<LLVM::LLVMFuncOp>();
+    return func.getArgument(1);
+  }
+}
+
 struct GetDeviceIdOpConversion
     : public ConvertOpToLLVMPattern<tle::GetDeviceIdOp> {
   GetDeviceIdOpConversion(LLVMTypeConverter &typeConverter,
@@ -29,11 +38,19 @@ struct GetDeviceIdOpConversion
   LogicalResult
   matchAndRewrite(tle::GetDeviceIdOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
-    Value src = adaptor.getInput();
-
+    auto loc = op.getLoc();
+    SmallVector<Value> srcElems;
+    if (auto src = adaptor.getInput())
+      srcElems = unpackLLElements(loc, src, rewriter);
+    auto func = op->getParentOfType<LLVM::LLVMFuncOp>();
+    if (!func) {
+      return rewriter.notifyMatchFailure(
+          op, "expected parent LLVM::LLVMFuncOp, but none was found. ");
+    }
+    auto comm = getDistDevicePtr(op, srcElems);
+    rewriter.modifyOpInPlace(op, [&]() { op->insertOperands(0, comm); });
     auto localRank = rewriter.create<tle::GetLocalRankOp>(
-        op.getLoc(), rewriter.getI32Type(), src);
-
+        op.getLoc(), rewriter.getI32Type(), comm);
     rewriter.replaceOp(op, localRank.getResult());
 
     return success();
diff --git a/third_party/tle/dialect/lib/Conversion/TleToLLVM/LocalPointersOpToLLVM.cpp b/third_party/tle/dialect/lib/Conversion/TleToLLVM/LocalPointersOpToLLVM.cpp
index 4dd033c45..28ef347da 100644
--- a/third_party/tle/dialect/lib/Conversion/TleToLLVM/LocalPointersOpToLLVM.cpp
+++ b/third_party/tle/dialect/lib/Conversion/TleToLLVM/LocalPointersOpToLLVM.cpp
@@ -374,7 +374,7 @@ LogicalResult lowerClusterSpace(Location loc, ValueRange srcElems,
   return success();
 }
 
-LogicalResult lowerDeviceSpace(Location loc, ValueRange srcElems,
+LogicalResult lowerDeviceSpace(Location loc, Value mem_ptr,
                                ValueRange shardElems, Value offsetVal,
                                int elemBytes,
                                ConversionPatternRewriter &rewriter,
@@ -387,7 +387,7 @@ LogicalResult lowerDeviceSpace(Location loc, ValueRange srcElems,
   }
   auto func = getOrInsertGetPeerPointer(module, rewriter.getContext());
 
-  Value memPtr = srcElems[0];
+  Value memPtr = mem_ptr;
   Value peer = shardElems[0];
 
   auto i64Ty = rewriter.getI64Type();
@@ -452,6 +452,15 @@ LogicalResult lowerNodeSpace(Location loc, ValueRange srcElems,
   return failure(); // Not implemented yet
 }
 
+Value getDistDevicePtr(tle::RemotePointersOp op, SmallVector<Value> &srcElems) {
+  if (!srcElems.empty())
+    return srcElems[0];
+  else {
+    auto func = op->getParentOfType<LLVM::LLVMFuncOp>();
+    return func.getArgument(1);
+  }
+}
+
 struct RemotePointersOpConversion
     : public ConvertOpToLLVMPattern<tle::RemotePointersOp> {
   RemotePointersOpConversion(LLVMTypeConverter &typeConverter,
@@ -467,19 +476,24 @@ struct RemotePointersOpConversion
       llvm::errs() << "[RemotePointersOpConversion] " << msg << "\n";
       return rewriter.notifyMatchFailure(op, msg);
     };
-    auto srcElems = unpackLLElements(loc, adaptor.getSrc(), rewriter);
-    if (srcElems.empty())
+
+    SmallVector<Value> srcElems;
+    auto space = adaptor.getSpace();
+
+    if (auto src = adaptor.getSrc())
+      srcElems = unpackLLElements(loc, adaptor.getSrc(), rewriter);
+
+    if (space != "device" && srcElems.empty())
       return reportFailure("expected non-empty source pointer elements");
 
     auto shardElems = unpackLLElements(loc, adaptor.getShardId(), rewriter);
     if (shardElems.empty())
       return reportFailure("expected non-empty shard_id elements");
-    if (shardElems.size() != 1 && shardElems.size() != srcElems.size())
+    if (space != "device" && shardElems.size() != 1 &&
+        shardElems.size() != srcElems.size())
       return reportFailure(
           "shard_id must be scalar or match source pointer element count");
 
-    auto space = adaptor.getSpace();
-
     Value offsetVal;
     if (adaptor.getOffset()) {
       auto offsetElems = unpackLLElements(loc, adaptor.getOffset(), rewriter);
@@ -489,6 +503,8 @@ struct RemotePointersOpConversion
     }
 
     SmallVector<Value> mappedPtrs;
+    auto mem = getDistDevicePtr(op, srcElems);
+
     if (space == "cluster") {
       if (failed(lowerClusterSpace(loc, srcElems, shardElems, rewriter,
                                    mappedPtrs))) {
@@ -506,19 +522,17 @@ struct RemotePointersOpConversion
               "result pointee type must be scalar int or float");
         elemBytes = elemBits.value() / 8;
       }
-      if (failed(lowerDeviceSpace(loc, srcElems, shardElems, offsetVal,
-                                  elemBytes, rewriter, mappedPtrs))) {
+      if (failed(lowerDeviceSpace(loc, mem, shardElems, offsetVal, elemBytes,
+                                  rewriter, mappedPtrs))) {
         return rewriter.notifyMatchFailure(op, "device lowering failed");
       }
     } else if (space == "node") {
-      if (failed(lowerNodeSpace(loc, srcElems, shardElems, rewriter,
-                                mappedPtrs))) {
+      if (failed(lowerNodeSpace(loc, mem, shardElems, rewriter, mappedPtrs))) {
         return rewriter.notifyMatchFailure(op, "node lowering failed");
       }
     } else {
       return reportFailure("unsupported remote space: " + space.str());
     }
-
     Value packed =
         packLLElements(loc, typeConverter, mappedPtrs, rewriter, op.getType());
     rewriter.replaceOp(op, packed);
diff --git a/third_party/tle/dialect/lib/IR/FlagCxOps.cpp b/third_party/tle/dialect/lib/IR/FlagCxOps.cpp
index 1af15b524..1cca185a5 100644
--- a/third_party/tle/dialect/lib/IR/FlagCxOps.cpp
+++ b/third_party/tle/dialect/lib/IR/FlagCxOps.cpp
@@ -14,6 +14,21 @@
 #include "triton/Dialect/TritonGPU/IR/Dialect.h"
 #include "triton/Dialect/TritonGPU/IR/LinearLayoutConversions.h"
 
+enum class CoopKind : int32_t {
+  Thread = 0,
+  Warp = 1,
+  Block = 2,
+  TileSpan = 3,
+  Lanes = 4,
+};
+
+enum class MemoryOrder : int32_t {
+  Relaxed = 0,
+  Acquire = 1,
+  Release = 2,
+  AcqRel = 3,
+};
+
 namespace mlir::triton::tle {
 
 LogicalResult GetLocalRankOp::verify() {
@@ -25,4 +40,75 @@ LogicalResult GetLocalRankOp::verify() {
   return success();
 }
 
+LogicalResult DeviceIntraBarrierOp::verify() {
+  auto *op = getOperation();
+
+  auto barrierTypeAttr = getBarrierTypeAttr();
+  auto coopKindAttr = getCoopKindAttr();
+  auto orderAttr = getOrderAttr();
+
+  auto emitInvalidIntAttr = [&](StringRef attrName, int64_t value,
+                                StringRef expected) -> LogicalResult {
+    return op->emitOpError() << "invalid " << attrName << " (" << value
+                             << "), expected one of: " << expected;
+  };
+
+  auto emitInvalidStrAttr = [&](StringRef attrName, StringRef value,
+                                StringRef expected) -> LogicalResult {
+    return op->emitOpError() << "invalid " << attrName << " '" << value
+                             << "', expected one of: " << expected;
+  };
+
+  // barrier_type
+  if (barrierTypeAttr) {
+    StringRef barrierType = barrierTypeAttr.getValue();
+
+    bool valid = llvm::StringSwitch<bool>(barrierType)
+                     .Case("arrive", true)
+                     .Case("wait", true)
+                     .Case("sync", true)
+                     .Default(false);
+
+    if (!valid)
+      return emitInvalidStrAttr("barrier_type", barrierType,
+                                "arrive, wait, sync");
+  }
+
+  // coop_kind
+  if (coopKindAttr) {
+    auto coopKind = static_cast<CoopKind>(coopKindAttr.getInt());
+
+    switch (coopKind) {
+    case CoopKind::Thread:
+    case CoopKind::Warp:
+    case CoopKind::Block:
+    case CoopKind::TileSpan:
+    case CoopKind::Lanes:
+      break;
+    default:
+      return emitInvalidIntAttr(
+          "coop_kind", coopKindAttr.getInt(),
+          "Thread(0), Warp(1), Block(2), TileSpan(3), Lanes(4)");
+    }
+  }
+
+  // order
+  if (orderAttr) {
+    auto order = static_cast<MemoryOrder>(orderAttr.getInt());
+
+    switch (order) {
+    case MemoryOrder::Relaxed:
+    case MemoryOrder::Acquire:
+    case MemoryOrder::Release:
+    case MemoryOrder::AcqRel:
+      break;
+    default:
+      return emitInvalidIntAttr(
+          "order", orderAttr.getInt(),
+          "Relaxed(0), Acquire(1), Release(2), AcqRel(3)");
+    }
+  }
+
+  return success();
+}
 } // namespace mlir::triton::tle
diff --git a/third_party/tle/dialect/lib/IR/Ops.cpp b/third_party/tle/dialect/lib/IR/Ops.cpp
index e3a868b9b..2856d82bc 100644
--- a/third_party/tle/dialect/lib/IR/Ops.cpp
+++ b/third_party/tle/dialect/lib/IR/Ops.cpp
@@ -812,6 +812,11 @@ LogicalResult ExclusiveCumsumOp::verify() {
 
 LogicalResult DistributedBarrierOp::verify() {
   auto *op = getOperation();
+  auto spaceAttr = op->getAttrOfType<StringAttr>("space");
+
+  if (spaceAttr && spaceAttr.getValue() == "device")
+    return DistributedBarrier::verifyDeviceSpace(op, getSrc());
+
   auto kindAttr = op->getAttrOfType<StringAttr>("group_kind");
   auto rankAttr = op->getAttrOfType<IntegerAttr>("group_rank");
   auto shapeAttr = op->getAttrOfType<DenseI32ArrayAttr>("group_shape");
@@ -892,13 +897,13 @@ LogicalResult DistributedBarrierOp::verify() {
 }
 
 LogicalResult RemotePointersOp::verify() {
-  Type srcTy = getSrc().getType();
-  Type resultTy = getResult().getType();
   auto spaceAttr = getSpace();
   if (spaceAttr == "device") {
     if (failed(RemotePointers::verifyDeviceSpace(getSrc(), getResult())))
       return failure();
   } else {
+    Type srcTy = getSrc().getType();
+    Type resultTy = getResult().getType();
     auto getPtrInfo = [&](Type ty, triton::PointerType &ptr, bool &isTensor,
                           ArrayRef<int64_t> &shape,
                           Attribute &encoding) -> LogicalResult {
diff --git a/third_party/tle/dialect/lib/IR/VerfiyUtils.cpp b/third_party/tle/dialect/lib/IR/VerfiyUtils.cpp
index 753faccae..5452c7b19 100644
--- a/third_party/tle/dialect/lib/IR/VerfiyUtils.cpp
+++ b/third_party/tle/dialect/lib/IR/VerfiyUtils.cpp
@@ -17,10 +17,8 @@
 namespace mlir::triton::tle {
 namespace RemotePointers {
 llvm::LogicalResult verifyDeviceSpace(mlir::Value src, mlir::Value result) {
-  // flagcxGetIntraPointerC accept raw device pointers represented as signless
-  // i64 values.
-  if (!src.getType().isSignlessInteger(64))
-    return failure();
+  if (!src)
+    return success();
 
   if (auto tensorTy = dyn_cast<RankedTensorType>(result.getType())) {
     auto ptr = dyn_cast<triton::PointerType>(tensorTy.getElementType());
@@ -32,4 +30,21 @@ llvm::LogicalResult verifyDeviceSpace(mlir::Value src, mlir::Value result) {
 }
 } // namespace RemotePointers
 
+namespace DistributedBarrier {
+llvm::LogicalResult verifyDeviceSpace(mlir::Operation *op, mlir::Value src) {
+
+  auto kindAttr = op->getAttrOfType<StringAttr>("group_kind");
+  auto barrierTypeAttr = op->getAttrOfType<StringAttr>("barrier_type");
+  auto orderAttr = op->getAttrOfType<StringAttr>("order");
+
+  if (kindAttr && barrierTypeAttr && orderAttr)
+    return success();
+  else
+    return op->emitOpError()
+           << "expects src, group_kind, barrier_type and order attributes to "
+              "be present for device space distributed barrier";
+}
+
+} // namespace DistributedBarrier
+
 } // namespace mlir::triton::tle
diff --git a/third_party/tle/dialect/lib/Tools/FlagcxUtils.cpp b/third_party/tle/dialect/lib/Tools/FlagcxUtils.cpp
index bc7bdd3ea..e43034b63 100644
--- a/third_party/tle/dialect/lib/Tools/FlagcxUtils.cpp
+++ b/third_party/tle/dialect/lib/Tools/FlagcxUtils.cpp
@@ -13,7 +13,9 @@ using namespace mlir;
 static const llvm::StringMap<StringRef> runtimeNames = {
     {"getLocalPeFunction", "flagcxDevCommGetIntraRank"},
     {"getNumPesFunction", "flagcxDevCommGetIntraSize"},
-    {"getReadSignalFunction", "flagcxDevNetReadSignal"}};
+    {"getIntraBarrierArriveSignalFunction", "flagcxIntraBarrierArriveS"},
+    {"getIntraBarrierWaitSignalFunction", "flagcxIntraBarrierWaitS"},
+    {"getIntraBarrierSyncSignalFunction", "flagcxIntraBarrierSyncS"}};
 
 static inline LLVM::LLVMFuncOp createFuncInstance(const char *funcName,
                                                   ModuleOp module,
@@ -31,9 +33,12 @@ static inline LLVM::LLVMFuncOp createFuncInstance(const char *funcName,
   return func;
 }
 
-static inline Value getFlagcxMemPtr(mlir::Location loc,
-                                    ConversionPatternRewriter &rewriter,
-                                    Value memPtrInt) {
+// The frontend passes the FlagCX global memory/communication pointer as an
+// integer. Convert it back to an LLVM pointer in global address space (AS1)
+// before passing it to device/runtime functions.
+static inline Value getFlagcxMemOrCommPtr(mlir::Location loc,
+                                          ConversionPatternRewriter &rewriter,
+                                          Value memPtrInt) {
   auto ctx = rewriter.getContext();
   auto ptrTy = LLVM::LLVMPointerType::get(ctx, 1);
   return rewriter.create<LLVM::IntToPtrOp>(loc, ptrTy, memPtrInt);
@@ -51,28 +56,49 @@ LLVM::CallOp getNumPesFunCall(mlir::Location loc,
   auto func = createFuncInstance(
       runtimeNames.lookup("getNumPesFunction").data(), module, {PtrTy}, i32Ty);
 
-  auto comm_dev_ptr = getFlagcxMemPtr(loc, rewriter, memPtrInt);
+  auto comm_dev_ptr = getFlagcxMemOrCommPtr(loc, rewriter, memPtrInt);
   return rewriter.create<LLVM::CallOp>(
       loc, TypeRange{func.getFunctionType().getReturnType()},
       FlatSymbolRefAttr::get(func), ValueRange{comm_dev_ptr});
 }
 
-LLVM::CallOp getReadSignalFunCall(mlir::Location loc,
-                                  ConversionPatternRewriter &rewriter,
-                                  Value memPtrInt) {
+LLVM::CallOp getBarrierFuncCall(mlir::Location loc,
+                                ConversionPatternRewriter &rewriter, Value comm,
+                                size_t barrier_index, size_t coopKind,
+                                size_t order, llvm::StringRef barrierType) {
   auto ctx = rewriter.getContext();
   ModuleOp module =
       rewriter.getInsertionPoint()->getParentOp()->getParentOfType<ModuleOp>();
 
   auto PtrTy = LLVM::LLVMPointerType::get(ctx, 1);
   auto i32Ty = IntegerType::get(ctx, 32);
-  auto func = createFuncInstance(
-      runtimeNames.lookup("getNumPesFunction").data(), module, {PtrTy}, i32Ty);
-
-  auto comm_dev_ptr = getFlagcxMemPtr(loc, rewriter, memPtrInt);
+  auto i1Ty = IntegerType::get(ctx, 1);
+  auto funcName = "";
+  if (barrierType == "arrive") {
+    funcName = "getIntraBarrierArriveSignalFunction";
+  } else if (barrierType == "wait") {
+    funcName = "getIntraBarrierWaitSignalFunction";
+  } else if (barrierType == "sync") {
+    funcName = "getIntraBarrierSyncSignalFunction";
+  } else {
+    llvm_unreachable("Unknown barrier type");
+  }
+
+  auto func = createFuncInstance(runtimeNames.lookup(funcName).data(), module,
+                                 {PtrTy, i32Ty, i32Ty, i1Ty, i32Ty}, i32Ty);
+
+  auto comm_dev_ptr = getFlagcxMemOrCommPtr(loc, rewriter, comm);
+  auto falseVal =
+      rewriter.create<LLVM::ConstantOp>(loc, i1Ty, rewriter.getBoolAttr(false));
+  auto barrierIndexVal =
+      rewriter.create<LLVM::ConstantOp>(loc, i32Ty, barrier_index);
+  auto coopKindVal = rewriter.create<LLVM::ConstantOp>(loc, i32Ty, coopKind);
+  auto orderVal = rewriter.create<LLVM::ConstantOp>(loc, i32Ty, order);
   return rewriter.create<LLVM::CallOp>(
       loc, TypeRange{func.getFunctionType().getReturnType()},
-      FlatSymbolRefAttr::get(func), ValueRange{comm_dev_ptr});
+      FlatSymbolRefAttr::get(func),
+      ValueRange{comm_dev_ptr, coopKindVal, barrierIndexVal, falseVal,
+                 orderVal});
 }
 
 LLVM::CallOp getLocalPeFuncCall(mlir::Location loc,
@@ -87,7 +113,7 @@ LLVM::CallOp getLocalPeFuncCall(mlir::Location loc,
   auto func = createFuncInstance(
       runtimeNames.lookup("getLocalPeFunction").data(), module, {PtrTy}, i32Ty);
 
-  auto comm_dev_ptr = getFlagcxMemPtr(loc, rewriter, memPtrInt);
+  auto comm_dev_ptr = getFlagcxMemOrCommPtr(loc, rewriter, memPtrInt);
   return rewriter.create<LLVM::CallOp>(
       loc, TypeRange{func.getFunctionType().getReturnType()},
       FlatSymbolRefAttr::get(func), ValueRange{comm_dev_ptr});
diff --git a/third_party/tle/dialect/lib/Transforms/CMakeLists.txt b/third_party/tle/dialect/lib/Transforms/CMakeLists.txt
index f77a3b937..f6ee4f946 100644
--- a/third_party/tle/dialect/lib/Transforms/CMakeLists.txt
+++ b/third_party/tle/dialect/lib/Transforms/CMakeLists.txt
@@ -28,6 +28,7 @@ add_triton_library(TritonTLETransforms
   RemoveRedundantCopy.cpp
   DSLRegionInline.cpp
   TleUtility.cpp
+  TleAddDistributedParams.cpp
 
   DEPENDS
   TritonTLETransformsIncGen
diff --git a/third_party/tle/dialect/lib/Transforms/TleAddDistributedParams.cpp b/third_party/tle/dialect/lib/Transforms/TleAddDistributedParams.cpp
new file mode 100644
index 000000000..49fd83f50
--- /dev/null
+++ b/third_party/tle/dialect/lib/Transforms/TleAddDistributedParams.cpp
@@ -0,0 +1,74 @@
+#include "mlir/Dialect/Arith/IR/Arith.h"
+#include "mlir/IR/BuiltinOps.h"
+#include "mlir/IR/PatternMatch.h"
+#include "mlir/Pass/Pass.h"
+#include "tle/dialect/include/IR/Dialect.h"
+#include "tle/dialect/include/Transforms/Passes.h"
+#include "triton/Dialect/TritonGPU/IR/Dialect.h"
+#include "triton/Dialect/TritonNvidiaGPU/IR/Dialect.h"
+
+namespace mlir::triton::tle {
+
+namespace ttg = mlir::triton::gpu;
+namespace ttng = mlir::triton::nvidia_gpu;
+
+#define GEN_PASS_DEF_TRITONTLEADDDISTRIBUTEDPARAMS
+#include "tle/dialect/include/Transforms/Passes.h.inc"
+
+namespace {
+
+void addDistributedParams(MLIRContext *ctx, mlir::triton::FuncOp func) {
+  auto i64Ty = IntegerType::get(ctx, 64);
+
+  //
+  // 1. Update function type.
+  //
+  auto oldType = func.getFunctionType();
+
+  SmallVector<Type> inputTypes;
+  inputTypes.push_back(i64Ty); // device_comm_ptr
+  inputTypes.push_back(i64Ty); // device_mem_ptr
+  llvm::append_range(inputTypes, oldType.getInputs());
+
+  func.setFunctionType(
+      FunctionType::get(ctx, inputTypes, oldType.getResults()));
+
+  //
+  // 2. Update argument attributes.
+  //
+  auto makeNameAttr = [&](StringRef name) {
+    return DictionaryAttr::get(ctx,
+                               {NamedAttribute(StringAttr::get(ctx, "tt.name"),
+                                               StringAttr::get(ctx, name))});
+  };
+
+  SmallVector<Attribute> argAttrs;
+  argAttrs.push_back(makeNameAttr("device_comm_ptr"));
+  argAttrs.push_back(makeNameAttr("device_mem_ptr"));
+
+  if (auto oldAttrs = func.getArgAttrsAttr())
+    llvm::append_range(argAttrs, oldAttrs);
+
+  func.setArgAttrsAttr(ArrayAttr::get(ctx, argAttrs));
+
+  //
+  // 3. Insert block arguments.
+  //
+  Block &entry = func.front();
+  entry.insertArgument(static_cast<unsigned>(0), i64Ty, func.getLoc());
+  entry.insertArgument(static_cast<unsigned>(1), i64Ty, func.getLoc());
+}
+
+struct TritonTleAddDistributedParams
+    : public impl::TritonTleAddDistributedParamsBase<
+          TritonTleAddDistributedParams> {
+  void runOnOperation() override {
+    ModuleOp m = getOperation();
+    MLIRContext *ctx = m.getContext();
+
+    m.walk([&](mlir::triton::FuncOp func) { addDistributedParams(ctx, func); });
+  }
+};
+
+} // namespace
+} // namespace mlir::triton::tle
diff --git a/third_party/tle/triton_tle.cc b/third_party/tle/triton_tle.cc
index 7709e90b2..8d595947a 100644
--- a/third_party/tle/triton_tle.cc
+++ b/third_party/tle/triton_tle.cc
@@ -525,9 +525,37 @@ void init_triton_tle_ir(py::module &&m) {
       .def("create_distributed_barrier",
            [](TritonOpBuilder &self) -> void {
              self.create<tle::DistributedBarrierOp>(
-                 StringAttr(), IntegerAttr(), DenseI32ArrayAttr(),
-                 DenseI32ArrayAttr(), DenseI32ArrayAttr());
+                 Value(), StringAttr(), StringAttr(), StringAttr(),
+                 StringAttr(), IntegerAttr(), IntegerAttr(),
+                 DenseI32ArrayAttr(), DenseI32ArrayAttr(), DenseI32ArrayAttr());
            })
+      .def(
+          "create_distributed_barrier",
+          [](TritonOpBuilder &self, std::optional<Value> src,
+             size_t barrier_index = 0, const std::string &space = "device",
+             const std::string &group_kind = "block",
+             const std::string &order = "acqrel",
+             const std::string &barrier_kind = "sync") -> void {
+            auto &builder = self.getBuilder();
+            auto *ctx = builder.getContext();
+            auto getOptStrAttr = [&](const std::string &s) -> StringAttr {
+              return s.empty() ? StringAttr() : builder.getStringAttr(s);
+            };
+            auto spaceAttr = getOptStrAttr(space);
+            auto kindAttr = getOptStrAttr(group_kind);
+            auto orderAttr = getOptStrAttr(order);
+            auto barrierTypeAttr = getOptStrAttr(barrier_kind);
+            auto barrierIndexAttr =
+                builder.getI32IntegerAttr(static_cast<int32_t>(barrier_index));
+
+            self.create<tle::DistributedBarrierOp>(
+                src.value_or(Value()), spaceAttr, barrierTypeAttr, orderAttr,
+                kindAttr, barrierIndexAttr, IntegerAttr(), DenseI32ArrayAttr(),
+                DenseI32ArrayAttr(), DenseI32ArrayAttr());
+          },
+          py::arg("src") = py::none(), py::arg("barrier_index"),
+          py::arg("space"), py::arg("group_kind"), py::arg("order"),
+          py::arg("barrier_kind"))
       .def(
           "create_distributed_barrier",
           [](TritonOpBuilder &self, const std::string &groupKind,
@@ -563,14 +591,15 @@ void init_triton_tle_ir(py::module &&m) {
             }
 
             self.create<tle::DistributedBarrierOp>(
-                kindAttr, rankAttr, shapeAttr, axesAttr, maskAttr);
+                Value(), StringAttr(), StringAttr(), StringAttr(), kindAttr,
+                IntegerAttr(), rankAttr, shapeAttr, axesAttr, maskAttr);
           },
           py::arg("group_kind"), py::arg("group_shape"), py::arg("group_axes"),
           py::arg("group_mask"))
       .def(
           "create_remote_pointers",
-          [](TritonOpBuilder &self, Type resultTy, Value src, Value shardId,
-             const std::string &space,
+          [](TritonOpBuilder &self, Type resultTy, std::optional<Value> &src,
+             Value shardId, const std::string &space,
              std::optional<Value> &offset) -> OpState {
             auto &builder = self.getBuilder();
             static const std::unordered_set<std::string> valid = {
@@ -581,15 +610,19 @@ void init_triton_tle_ir(py::module &&m) {
                   ". Expected one of: cluster, device, node.");
             }
             auto space_attr = builder.getStringAttr(space);
+
             return self.create<tle::RemotePointersOp>(
-                resultTy, src, shardId, space_attr, offset.value_or(Value()));
+                resultTy, src.value_or(Value()), shardId, space_attr,
+                offset.value_or(Value()));
           },
-          py::arg("resultTy"), py::arg("src"), py::arg("shardId"),
+          py::arg("resultTy"), py::arg("src") = py::none(), py::arg("shardId"),
           py::arg("space"), py::arg("offset") = py::none())
       .def("get_device_id",
-           [](TritonOpBuilder &self, Type resultTy, Value src) -> Value {
+           [](TritonOpBuilder &self, Type resultTy,
+              std::optional<Value> src) -> Value {
              auto &builder = self.getBuilder();
-             return self.create<tle::GetDeviceIdOp>(resultTy, src);
+             return self.create<tle::GetDeviceIdOp>(resultTy,
+                                                    src.value_or(Value()));
            })
       .def("get_n_pes",
            [](TritonOpBuilder &self, Type resultTy, Value src) -> Value {
@@ -632,6 +665,8 @@ void init_triton_tle_ir(py::module &&m) {
 }
 
 void init_triton_tle_passes(py::module &&m) {
+  ADD_PASS_WRAPPER_0("add_params_for_distribution",
+                     tle::createTritonTleAddDistributedParams);
   ADD_PASS_WRAPPER_0("add_early_assign_memory_space",
                      tle::createTritonTleEarlyAssignMemorySpace);
   ADD_PASS_WRAPPER_0("add_select_encodings",