From 9d1bd7cab009cdbd5f57a7c0fbb8fb5a4ad1ebd8 Mon Sep 17 00:00:00 2001
From: Chapaman <14204271+Chapaman@users.noreply.github.com>
Date: Fri, 30 Jan 2026 16:35:04 -0300
Subject: [PATCH 1/7] Basic gall_gather implementation

---
 exla/lib/exla/defn.ex       | 18 ++++++++++++++++++
 exla/lib/exla/mlir/value.ex | 27 +++++++++++++++++++++++++++
 nx/lib/nx/defn/evaluator.ex |  5 +++++
 nx/lib/nx/defn/expr.ex      | 27 +++++++++++++++++++++++++++
 nx/lib/nx/defn/kernel.ex    | 26 ++++++++++++++++++++++++++
 5 files changed, 103 insertions(+)

diff --git a/exla/lib/exla/defn.ex b/exla/lib/exla/defn.ex
index 592a8279b4..4cd674ab97 100644
--- a/exla/lib/exla/defn.ex
+++ b/exla/lib/exla/defn.ex
@@ -1471,6 +1471,24 @@ defmodule EXLA.Defn do
     EXLA.Lib.argsort(state.builder, tensor, dimension, stable, comp, ans.type)
   end
 
+## to_operator collective ops
+
+  defp to_operator(:all_gather, [%Value{} = tensor, opts], ans, _state) do
+    all_gather_dim = Keyword.fetch!(opts, :all_gather_dim)
+    replica_groups = Keyword.fetch!(opts, :replica_groups)
+    use_global_device_ids = Keyword.get(opts, :use_global_device_ids, false)
+
+    Value.all_gather(
+      [tensor],
+      expr_to_typespec(ans),
+      all_gather_dim,
+      replica_groups,
+      use_global_device_ids,
+      Keyword.take(opts, [:channel_id])
+    )
+    |> hd()
+  end
+
   defp fft(exla_op, [%Value{} = tensor, opts], %{type: type} = ans, state) do
     n = opts[:length]
     axis = opts[:axis]
diff --git a/exla/lib/exla/mlir/value.ex b/exla/lib/exla/mlir/value.ex
index 393b6d57a8..9b6822c6dd 100644
--- a/exla/lib/exla/mlir/value.ex
+++ b/exla/lib/exla/mlir/value.ex
@@ -64,6 +64,33 @@ defmodule EXLA.MLIR.Value do
     end
   end
 
+  def all_gather([%Value{function: func} | _] = operands, typespec, all_gather_dim, replica_groups, use_global_device_ids, opts \\ []) do
+    result_types = typespecs_to_mlir_types([typespec])
+
+    opts =
+      Keyword.validate!(opts, [
+        channel_id: nil,
+      ])
+
+      num_groups = length(replica_groups)
+      group_size = if num_groups > 0, do: length(hd(replica_groups)), else: 0
+      flat_groups = List.flatten(replica_groups)
+
+      attributes = [
+        all_gather_dim: attr_i64(all_gather_dim),
+        replica_groups: attr_dense_elements(flat_groups, {:s, 64}, {num_groups, group_size}),
+        use_global_device_ids: attr_boolean(use_global_device_ids)
+      ]
+
+    attributes =
+      if opts[:channel_id] do
+        attributes ++ [channel_id: attr_i64(opts[:channel_id])]
+      else
+        attributes end
+
+    op(func, "stablehlo.all_gather", operands, result_types, attributes: attributes)
+  end
+
   defp compare_and_return_bool(func, lhs, rhs, typespec, direction, total_order? \\ false) do
     %{type: lhs_type} = get_typespec(lhs)
     %{type: rhs_type} = get_typespec(rhs)
diff --git a/nx/lib/nx/defn/evaluator.ex b/nx/lib/nx/defn/evaluator.ex
index 601f750942..5e33925625 100644
--- a/nx/lib/nx/defn/evaluator.ex
+++ b/nx/lib/nx/defn/evaluator.ex
@@ -478,6 +478,11 @@ defmodule Nx.Defn.Evaluator do
           {Nx.Shared.list_impl!(args), [ans | args]}
       end
 
+    if op == :all_gather and not function_exported?(mod, :all_gather, 3) do
+      raise ArgumentError,
+        "all_gather/3 is not supported by backend #{inspect(mod)}."
+    end
+
     {apply(mod, op, args), caches}
   end
 
diff --git a/nx/lib/nx/defn/expr.ex b/nx/lib/nx/defn/expr.ex
index 899b430da4..ab11c46af1 100644
--- a/nx/lib/nx/defn/expr.ex
+++ b/nx/lib/nx/defn/expr.ex
@@ -1166,6 +1166,33 @@ defmodule Nx.Defn.Expr do
     expr(out, context, :gather, [tensor, indices, opts])
   end
 
+  def all_gather(tensor, opts) do
+    {[tensor], context} = to_exprs([tensor])
+
+    _all_gather_dim = opts[:all_gather_dim]
+    replica_groups = opts[:replica_groups]
+
+    # Calculate group size (number of replicas per group)
+    _group_size =
+      case replica_groups do
+        [first_group | _] -> length(first_group)
+        [] -> 1
+      end
+
+    # Calculate output shape by multiplying the gather dimension by group_size
+    input_shape = tensor.shape
+    output_shape =
+      input_shape
+#      |> Tuple.to_list()
+#      |> List.update_at(all_gather_dim, &(&1 * group_size))
+#      |> List.to_tuple()
+
+    # Create output tensor with the new shape
+    out = %{tensor | shape: output_shape}
+
+    expr(out, context, :all_gather, [tensor, opts])
+  end
+
   @impl true
   def reverse(out, tensor, axes) do
     tensor = to_expr(tensor)
diff --git a/nx/lib/nx/defn/kernel.ex b/nx/lib/nx/defn/kernel.ex
index ab913ab61f..a0cf4f4493 100644
--- a/nx/lib/nx/defn/kernel.ex
+++ b/nx/lib/nx/defn/kernel.ex
@@ -1669,6 +1669,32 @@ defmodule Nx.Defn.Kernel do
     end
   end
 
+  @doc """
+  Gathers tensors from all replicas along a specified dimension.
+
+  This operation concatenates tensors from multiple replicas/devices along
+  the specified dimension. Requires a backend that supports multi-device operations.
+
+  ## Parameters
+
+    * `tensor` - The input tensor to gather
+    * `all_gather_dim` - The dimension along which to gather
+    * `replica_groups` - 2D list defining how replicas are grouped (required)
+    * `opts` - Optional keyword list:
+      * `:use_global_device_ids` - Whether to use global device IDs (default: false)
+      * `:channel_id` - Channel ID for communication (optional)
+
+  ## Examples
+
+    all_gather(tensor, 0, [[0, 1, 2, 3]])
+    all_gather(tensor, 1, [[0, 1], [2, 3]], use_global_device_ids: true)
+  """
+  def all_gather(tensor, all_gather_dim, replica_groups, opts \\ []) do
+    opts = Keyword.put(opts, :all_gather_dim, all_gather_dim)
+    opts = Keyword.put(opts, :replica_groups, replica_groups)
+    Nx.Defn.Expr.all_gather(tensor, opts)
+  end
+
   @definitions (Module.definitions_in(__MODULE__, :def) ++
                   Module.definitions_in(__MODULE__, :defmacro)) --
                  [

From cc8761d4fa10441bfbe658dd30610a9d1d1c74c1 Mon Sep 17 00:00:00 2001
From: Chapaman <14204271+Chapaman@users.noreply.github.com>
Date: Mon, 2 Feb 2026 19:33:10 -0300
Subject: [PATCH 2/7] changes due to code review by @polvalente

---
 exla/lib/exla.ex                      | 15 ++++++
 exla/lib/exla/defn.ex                 | 21 ++++----
 exla/lib/exla/mlir/value.ex           | 30 +++++------
 exla/test/exla/defn/sharding_test.exs | 73 ++++++++++++++++++++++++++-
 nx/lib/nx/defn/evaluator.ex           |  5 --
 nx/lib/nx/defn/kernel.ex              | 14 ++---
 nx/test/nx/defn_test.exs              | 13 +++++
 7 files changed, 128 insertions(+), 43 deletions(-)

diff --git a/exla/lib/exla.ex b/exla/lib/exla.ex
index 78c9016361..403c6fbe76 100644
--- a/exla/lib/exla.ex
+++ b/exla/lib/exla.ex
@@ -215,6 +215,21 @@ defmodule EXLA do
   The metadata is:
 
     * `:key` - the compilation key for debugging
+
+  ## Sharding
+
+  EXLA supports sharding, which is a way to partition a computation across multiple devices.
+  There are a number of collective operations that are supported by sharding.
+
+  ### [`all_gather`](https://openxla.org/stablehlo/spec#all_gather)
+
+  #### Options
+
+    * `:all_gather_dim` - the dimension along which to gather
+    * `:replica_groups` - 2D list defining how replicas are grouped
+    * `:use_global_device_ids` - Whether to use global device IDs (default: `false`)
+    * `:channel_id` - Channel ID for communication (optional)
+
   """
 
   @behaviour Nx.Defn.Compiler
diff --git a/exla/lib/exla/defn.ex b/exla/lib/exla/defn.ex
index 4cd674ab97..65239373ba 100644
--- a/exla/lib/exla/defn.ex
+++ b/exla/lib/exla/defn.ex
@@ -1478,15 +1478,18 @@ defmodule EXLA.Defn do
     replica_groups = Keyword.fetch!(opts, :replica_groups)
     use_global_device_ids = Keyword.get(opts, :use_global_device_ids, false)
 
-    Value.all_gather(
-      [tensor],
-      expr_to_typespec(ans),
-      all_gather_dim,
-      replica_groups,
-      use_global_device_ids,
-      Keyword.take(opts, [:channel_id])
-    )
-    |> hd()
+    # We might want to surface all_gather as an operation that takes a container of operands instead of a single one.
+    [result] =
+      Value.all_gather(
+        [tensor],
+        expr_to_typespec(ans),
+        all_gather_dim,
+        replica_groups,
+        use_global_device_ids,
+        opts[:channel_id]
+      )
+
+    result
   end
 
   defp fft(exla_op, [%Value{} = tensor, opts], %{type: type} = ans, state) do
diff --git a/exla/lib/exla/mlir/value.ex b/exla/lib/exla/mlir/value.ex
index 9b6822c6dd..e548693497 100644
--- a/exla/lib/exla/mlir/value.ex
+++ b/exla/lib/exla/mlir/value.ex
@@ -64,29 +64,25 @@ defmodule EXLA.MLIR.Value do
     end
   end
 
-  def all_gather([%Value{function: func} | _] = operands, typespec, all_gather_dim, replica_groups, use_global_device_ids, opts \\ []) do
+  def all_gather([%Value{function: func} | _] = operands, typespec, all_gather_dim, replica_groups, use_global_device_ids, channel_id \\ nil) do
     result_types = typespecs_to_mlir_types([typespec])
 
-    opts =
-      Keyword.validate!(opts, [
-        channel_id: nil,
-      ])
+    num_groups = length(replica_groups)
+    group_size = if num_groups > 0, do: length(hd(replica_groups)), else: 0
+    flat_groups = List.flatten(replica_groups)
 
-      num_groups = length(replica_groups)
-      group_size = if num_groups > 0, do: length(hd(replica_groups)), else: 0
-      flat_groups = List.flatten(replica_groups)
-
-      attributes = [
-        all_gather_dim: attr_i64(all_gather_dim),
-        replica_groups: attr_dense_elements(flat_groups, {:s, 64}, {num_groups, group_size}),
-        use_global_device_ids: attr_boolean(use_global_device_ids)
-      ]
+    attributes = [
+      all_gather_dim: attr_i64(all_gather_dim),
+      replica_groups: attr_dense_elements(flat_groups, {:s, 64}, {num_groups, group_size}),
+      use_global_device_ids: attr_boolean(use_global_device_ids)
+    ]
 
     attributes =
-      if opts[:channel_id] do
-        attributes ++ [channel_id: attr_i64(opts[:channel_id])]
+      if channel_id do
+        Keyword.put(attributes, :channel_id, attr_i64(channel_id))
       else
-        attributes end
+        attributes
+      end
 
     op(func, "stablehlo.all_gather", operands, result_types, attributes: attributes)
   end
diff --git a/exla/test/exla/defn/sharding_test.exs b/exla/test/exla/defn/sharding_test.exs
index ed46f76b6a..058e2683b6 100644
--- a/exla/test/exla/defn/sharding_test.exs
+++ b/exla/test/exla/defn/sharding_test.exs
@@ -6,7 +6,8 @@ defmodule EXLA.Defn.ShardingTest do
   describe "MLIR module generation with sharding" do
     @moduletag :multi_device
     test "generates correct MLIR with simple 2D mesh and sharding" do
-      fun = fn x, y -> Nx.add(x, y) end
+      fun = fn x, y -> Nx.add(x, y)
+    end
 
       mesh = %Mesh{name: "mesh", shape: {2, 2}}
       # First arg: shard dim 0 on mesh axis 0, dim 1 on mesh axis 1
@@ -737,5 +738,75 @@ defmodule EXLA.Defn.ShardingTest do
       assert result.mlir_module =~ ~r/"axis_0"/
       assert result.mlir_module =~ ~r/"axis_1"/
     end
+
+    @moduletag :multi_device
+    test "generates correct MLIR with all_gather" do
+      fun = fn x, y -> Nx.add(x, y)
+      |> Nx.Defn.Kernel.all_gather(all_gather_dim: 0, replica_groups: [[0]])
+      |> Nx.Defn.Kernel.all_gather(all_gather_dim: 1, replica_groups: [[0]])
+    end
+
+      mesh = %Mesh{name: "mesh", shape: {2, 2}}
+      # First arg: shard dim 0 on mesh axis 0, dim 1 on mesh axis 1
+      # Second arg: shard dim 0 on mesh axis 0, dim 1 not sharded
+      input_shardings = [%{0 => [0], 1 => [1]}, %{0 => [0]}]
+
+      # For mesh {2, 2}, we have 4 partitions
+      # Each partition gets a shard of the inputs
+      # First input: shape {8, 2} sharded as [[0], [1]] -> each partition gets {4, 1}
+      # Second input: shape {8, 1} sharded as [[0], []] -> each partition gets {4, 1}
+      args = [
+        # partition 0
+        [Nx.iota({4, 1}), Nx.iota({4, 1})],
+        # partition 1
+        [Nx.iota({4, 1}), Nx.iota({4, 1})],
+        # partition 2
+        [Nx.iota({4, 1}), Nx.iota({4, 1})],
+        # partition 3
+        [Nx.iota({4, 1}), Nx.iota({4, 1})]
+      ]
+
+      result = EXLA.to_mlir_module(fun, args, mesh: mesh, input_shardings: input_shardings)
+
+      expected_mlir = """
+      module {
+        sdy.mesh @mesh = <["axis_0"=2, "axis_1"=2]>
+        func.func public @main(%arg0: tensor<8x2xi32> {sdy.sharding = #sdy.sharding<@mesh, [{"axis_0", ?}p0, {"axis_1", ?}p0]>}, %arg1: tensor<8x1xi32> {sdy.sharding = #sdy.sharding<@mesh, [{"axis_0", ?}p0, {?}p0]>}) -> tensor<8x2xi32> {
+          %0 = stablehlo.broadcast_in_dim %arg1, dims = [0, 1] : (tensor<8x1xi32>) -> tensor<8x2xi32>
+          %1 = stablehlo.add %arg0, %0 : tensor<8x2xi32>
+          %2 = "stablehlo.all_gather"(%1) <{all_gather_dim = 0 : i64, replica_groups = dense<0> : tensor<1x1xi64>}> : (tensor<8x2xi32>) -> tensor<8x2xi32>
+          %3 = "stablehlo.all_gather"(%2) <{all_gather_dim = 1 : i64, replica_groups = dense<0> : tensor<1x1xi64>}> : (tensor<8x2xi32>) -> tensor<8x2xi32>
+          return %3 : tensor<8x2xi32>
+        }
+      }
+      """
+
+      assert expected_mlir == result.mlir_module
+
+      results = EXLA.shard_jit(fun, mesh, input_shardings: input_shardings).(args)
+
+      assert length(results) == 4
+
+      # After all_gather on both dims, each partition has the full tensor: add(iota, iota) -> 2*iota
+      # Each shard had iota({4,1}) = [[0],[1],[2],[3]], so add gives [[0],[2],[4],[6]]
+      # After gathering: replicated 8x2 with pattern [[0,0],[2,2],[4,4],[6,6],[0,0],[2,2],[4,4],[6,6]]
+      expected_result =
+        Nx.tensor([
+          [0, 0],
+          [2, 2],
+          [4, 4],
+          [6, 6],
+          [0, 0],
+          [2, 2],
+          [4, 4],
+          [6, 6]
+        ])
+
+      for r <- results do
+        assert_equal(r, expected_result)
+      end
+    end
+
+
   end
 end
diff --git a/nx/lib/nx/defn/evaluator.ex b/nx/lib/nx/defn/evaluator.ex
index 5e33925625..601f750942 100644
--- a/nx/lib/nx/defn/evaluator.ex
+++ b/nx/lib/nx/defn/evaluator.ex
@@ -478,11 +478,6 @@ defmodule Nx.Defn.Evaluator do
           {Nx.Shared.list_impl!(args), [ans | args]}
       end
 
-    if op == :all_gather and not function_exported?(mod, :all_gather, 3) do
-      raise ArgumentError,
-        "all_gather/3 is not supported by backend #{inspect(mod)}."
-    end
-
     {apply(mod, op, args), caches}
   end
 
diff --git a/nx/lib/nx/defn/kernel.ex b/nx/lib/nx/defn/kernel.ex
index a0cf4f4493..809a66b480 100644
--- a/nx/lib/nx/defn/kernel.ex
+++ b/nx/lib/nx/defn/kernel.ex
@@ -1678,20 +1678,12 @@ defmodule Nx.Defn.Kernel do
   ## Parameters
 
     * `tensor` - The input tensor to gather
-    * `all_gather_dim` - The dimension along which to gather
-    * `replica_groups` - 2D list defining how replicas are grouped (required)
-    * `opts` - Optional keyword list:
-      * `:use_global_device_ids` - Whether to use global device IDs (default: false)
-      * `:channel_id` - Channel ID for communication (optional)
 
-  ## Examples
+    * `opts` - Optional keyword list. These are backend- and compiler-specific;
+    see your backend or compiler docs for supported options.
 
-    all_gather(tensor, 0, [[0, 1, 2, 3]])
-    all_gather(tensor, 1, [[0, 1], [2, 3]], use_global_device_ids: true)
   """
-  def all_gather(tensor, all_gather_dim, replica_groups, opts \\ []) do
-    opts = Keyword.put(opts, :all_gather_dim, all_gather_dim)
-    opts = Keyword.put(opts, :replica_groups, replica_groups)
+  def all_gather(tensor, opts \\ []) do
     Nx.Defn.Expr.all_gather(tensor, opts)
   end
 
diff --git a/nx/test/nx/defn_test.exs b/nx/test/nx/defn_test.exs
index 62993b07a3..621b4f4e77 100644
--- a/nx/test/nx/defn_test.exs
+++ b/nx/test/nx/defn_test.exs
@@ -2952,4 +2952,17 @@ defmodule Nx.DefnTest do
       assert vectorized_metadata_tuple(x, z) == vec_nonvec_result
     end
   end
+
+  describe "sharding" do
+    defn all_gather_test(tensor) do
+      Nx.Defn.Kernel.all_gather(tensor, all_gather_dim: 0, replica_groups: [[0]])
+    end
+
+    @tag compiler: Evaluator
+    test "all_gather works" do
+      assert_raise UndefinedFunctionError, fn ->
+        all_gather_test(Nx.tensor([1, 2, 3, 4]))
+      end
+    end
+  end
 end

From 0a50a3e5fb457b8fb84a60d553b1e1ebccb325e9 Mon Sep 17 00:00:00 2001
From: Chapaman <14204271+Chapaman@users.noreply.github.com>
Date: Mon, 9 Feb 2026 18:26:32 -0300
Subject: [PATCH 3/7] added test in defn to guarantee output format

---
 nx/test/nx/defn_test.exs | 13 +++++++++++++
 1 file changed, 13 insertions(+)

diff --git a/nx/test/nx/defn_test.exs b/nx/test/nx/defn_test.exs
index 621b4f4e77..d15951ce71 100644
--- a/nx/test/nx/defn_test.exs
+++ b/nx/test/nx/defn_test.exs
@@ -2958,6 +2958,19 @@ defmodule Nx.DefnTest do
       Nx.Defn.Kernel.all_gather(tensor, all_gather_dim: 0, replica_groups: [[0]])
     end
 
+    test "all_gather produces correct expr format for compiler" do
+      # Uses debug_expr to inspect the expression without compiling.
+      # Guarantees the format passed to compilers (e.g. EXLA) stays stable.
+      assert %T{data: %Expr{op: :all_gather, args: [tensor, opts]}} =
+               Nx.Defn.debug_expr(&all_gather_test/1).(Nx.tensor([1, 2, 3, 4]))
+
+      assert %T{data: %Expr{op: :parameter, args: [0]}} = tensor
+
+      # Compilers expect opts with :all_gather_dim and :replica_groups
+      assert opts[:all_gather_dim] == 0
+      assert opts[:replica_groups] == [[0]]
+    end
+
     @tag compiler: Evaluator
     test "all_gather works" do
       assert_raise UndefinedFunctionError, fn ->

From 2daaad8dc3f64918cae50f040a031e35a373be78 Mon Sep 17 00:00:00 2001
From: Chapaman <14204271+Chapaman@users.noreply.github.com>
Date: Mon, 9 Feb 2026 18:37:28 -0300
Subject: [PATCH 4/7] added sharding confirmation to test

---
 exla/test/exla/defn/sharding_test.exs | 12 +++++++-----
 1 file changed, 7 insertions(+), 5 deletions(-)

diff --git a/exla/test/exla/defn/sharding_test.exs b/exla/test/exla/defn/sharding_test.exs
index 058e2683b6..7d1eb993e4 100644
--- a/exla/test/exla/defn/sharding_test.exs
+++ b/exla/test/exla/defn/sharding_test.exs
@@ -802,11 +802,13 @@ defmodule EXLA.Defn.ShardingTest do
           [6, 6]
         ])
 
-      for r <- results do
-        assert_equal(r, expected_result)
-      end
-    end
-
+      device_ids =
+        for r <- results do
+          assert_equal(r, expected_result)
+          r.data.buffer.device_id
+        end
 
+      assert Enum.sort(device_ids) == [0, 1, 2, 3]
+    end
   end
 end

From 1be394a1e7dec11666e14728e5cd4f734485f74e Mon Sep 17 00:00:00 2001
From: Chapaman <14204271+Chapaman@users.noreply.github.com>
Date: Wed, 11 Feb 2026 18:47:15 -0300
Subject: [PATCH 5/7] updated the test so the result on the test is clearer

---
 exla/test/exla/defn/sharding_test.exs | 60 +++++++++++++--------------
 1 file changed, 28 insertions(+), 32 deletions(-)

diff --git a/exla/test/exla/defn/sharding_test.exs b/exla/test/exla/defn/sharding_test.exs
index 7d1eb993e4..ccf3ce2583 100644
--- a/exla/test/exla/defn/sharding_test.exs
+++ b/exla/test/exla/defn/sharding_test.exs
@@ -747,23 +747,22 @@ defmodule EXLA.Defn.ShardingTest do
     end
 
       mesh = %Mesh{name: "mesh", shape: {2, 2}}
-      # First arg: shard dim 0 on mesh axis 0, dim 1 on mesh axis 1
-      # Second arg: shard dim 0 on mesh axis 0, dim 1 not sharded
-      input_shardings = [%{0 => [0], 1 => [1]}, %{0 => [0]}]
+      # First arg: 0..15 (8x2), shard dim 0 on mesh axis 0, dim 1 on mesh axis 1
+      # Second arg: 100..115 (8x2), same sharding — makes sharded results easy to read
+      input_shardings = [%{0 => [0], 1 => [1]}, %{0 => [0], 1 => [1]}]
 
-      # For mesh {2, 2}, we have 4 partitions
-      # Each partition gets a shard of the inputs
-      # First input: shape {8, 2} sharded as [[0], [1]] -> each partition gets {4, 1}
-      # Second input: shape {8, 1} sharded as [[0], []] -> each partition gets {4, 1}
+      # For mesh {2, 2}, 4 partitions. Each gets {4, 1}. Full 8x2 row-major: [[0,1],[2,3],...,[14,15]].
+      # Partition (axis_0, axis_1): (0,0)=rows 0-3 col 0, (0,1)=rows 0-3 col 1, (1,0)=rows 4-7 col 0, (1,1)=rows 4-7 col 1.
+      # So partition 0 gets (0,0),(1,0),(2,0),(3,0) = 0,2,4,6; partition 1 gets (0,1),(1,1),... = 1,3,5,7; etc.
       args = [
-        # partition 0
-        [Nx.iota({4, 1}), Nx.iota({4, 1})],
-        # partition 1
-        [Nx.iota({4, 1}), Nx.iota({4, 1})],
-        # partition 2
-        [Nx.iota({4, 1}), Nx.iota({4, 1})],
-        # partition 3
-        [Nx.iota({4, 1}), Nx.iota({4, 1})]
+        # partition 0: rows 0–3 col 0 -> 0,2,4,6 and 100,102,104,106
+        [Nx.tensor([[0], [2], [4], [6]]), Nx.tensor([[100], [102], [104], [106]])],
+        # partition 1: rows 0–3 col 1 -> 1,3,5,7 and 101,103,105,107
+        [Nx.tensor([[1], [3], [5], [7]]), Nx.tensor([[101], [103], [105], [107]])],
+        # partition 2: rows 4–7 col 0 -> 8,10,12,14 and 108,110,112,114
+        [Nx.tensor([[8], [10], [12], [14]]), Nx.tensor([[108], [110], [112], [114]])],
+        # partition 3: rows 4–7 col 1 -> 9,11,13,15 and 109,111,113,115
+        [Nx.tensor([[9], [11], [13], [15]]), Nx.tensor([[109], [111], [113], [115]])]
       ]
 
       result = EXLA.to_mlir_module(fun, args, mesh: mesh, input_shardings: input_shardings)
@@ -771,12 +770,11 @@ defmodule EXLA.Defn.ShardingTest do
       expected_mlir = """
       module {
         sdy.mesh @mesh = <["axis_0"=2, "axis_1"=2]>
-        func.func public @main(%arg0: tensor<8x2xi32> {sdy.sharding = #sdy.sharding<@mesh, [{"axis_0", ?}p0, {"axis_1", ?}p0]>}, %arg1: tensor<8x1xi32> {sdy.sharding = #sdy.sharding<@mesh, [{"axis_0", ?}p0, {?}p0]>}) -> tensor<8x2xi32> {
-          %0 = stablehlo.broadcast_in_dim %arg1, dims = [0, 1] : (tensor<8x1xi32>) -> tensor<8x2xi32>
-          %1 = stablehlo.add %arg0, %0 : tensor<8x2xi32>
-          %2 = "stablehlo.all_gather"(%1) <{all_gather_dim = 0 : i64, replica_groups = dense<0> : tensor<1x1xi64>}> : (tensor<8x2xi32>) -> tensor<8x2xi32>
-          %3 = "stablehlo.all_gather"(%2) <{all_gather_dim = 1 : i64, replica_groups = dense<0> : tensor<1x1xi64>}> : (tensor<8x2xi32>) -> tensor<8x2xi32>
-          return %3 : tensor<8x2xi32>
+        func.func public @main(%arg0: tensor<8x2xi32> {sdy.sharding = #sdy.sharding<@mesh, [{"axis_0", ?}p0, {"axis_1", ?}p0]>}, %arg1: tensor<8x2xi32> {sdy.sharding = #sdy.sharding<@mesh, [{"axis_0", ?}p0, {"axis_1", ?}p0]>}) -> tensor<8x2xi32> {
+          %0 = stablehlo.add %arg0, %arg1 : tensor<8x2xi32>
+          %1 = "stablehlo.all_gather"(%0) <{all_gather_dim = 0 : i64, replica_groups = dense<0> : tensor<1x1xi64>}> : (tensor<8x2xi32>) -> tensor<8x2xi32>
+          %2 = "stablehlo.all_gather"(%1) <{all_gather_dim = 1 : i64, replica_groups = dense<0> : tensor<1x1xi64>}> : (tensor<8x2xi32>) -> tensor<8x2xi32>
+          return %2 : tensor<8x2xi32>
         }
       }
       """
@@ -787,19 +785,17 @@ defmodule EXLA.Defn.ShardingTest do
 
       assert length(results) == 4
 
-      # After all_gather on both dims, each partition has the full tensor: add(iota, iota) -> 2*iota
-      # Each shard had iota({4,1}) = [[0],[1],[2],[3]], so add gives [[0],[2],[4],[6]]
-      # After gathering: replicated 8x2 with pattern [[0,0],[2,2],[4,4],[6,6],[0,0],[2,2],[4,4],[6,6]]
+      # After all_gather: full first arg 0..15 + full second 100..115 -> 100,102,...,130
       expected_result =
         Nx.tensor([
-          [0, 0],
-          [2, 2],
-          [4, 4],
-          [6, 6],
-          [0, 0],
-          [2, 2],
-          [4, 4],
-          [6, 6]
+          [100, 102],
+          [104, 106],
+          [108, 110],
+          [112, 114],
+          [116, 118],
+          [120, 122],
+          [124, 126],
+          [128, 130]
         ])
 
       device_ids =

From eaef13671e3be66c6c7945e0204aa6992eea234c Mon Sep 17 00:00:00 2001
From: Chapaman <14204271+Chapaman@users.noreply.github.com>
Date: Wed, 11 Feb 2026 19:46:48 -0300
Subject: [PATCH 6/7] intermediate commit

---
 exla/test/exla/defn/sharding_test.exs | 11 ++++-------
 1 file changed, 4 insertions(+), 7 deletions(-)

diff --git a/exla/test/exla/defn/sharding_test.exs b/exla/test/exla/defn/sharding_test.exs
index ccf3ce2583..3052fcc30d 100644
--- a/exla/test/exla/defn/sharding_test.exs
+++ b/exla/test/exla/defn/sharding_test.exs
@@ -798,13 +798,10 @@ defmodule EXLA.Defn.ShardingTest do
           [128, 130]
         ])
 
-      device_ids =
-        for r <- results do
-          assert_equal(r, expected_result)
-          r.data.buffer.device_id
-        end
-
-      assert Enum.sort(device_ids) == [0, 1, 2, 3]
+      for {r, partition_idx} <- Enum.with_index(results) do
+        assert_equal(r, expected_result)
+        assert r.data.buffer.device_id == partition_idx
+      end
     end
   end
 end

From d1c683869e8964e4f36012af8c3d4ee4080bd106 Mon Sep 17 00:00:00 2001
From: Chapaman <14204271+Chapaman@users.noreply.github.com>
Date: Wed, 11 Feb 2026 20:10:16 -0300
Subject: [PATCH 7/7] added partially sharded test + Enum.zip_with on previous
 test

---
 exla/test/exla/defn/sharding_test.exs | 50 ++++++++++++++++++++++++---
 1 file changed, 46 insertions(+), 4 deletions(-)

diff --git a/exla/test/exla/defn/sharding_test.exs b/exla/test/exla/defn/sharding_test.exs
index 4154d71808..b22e20a97a 100644
--- a/exla/test/exla/defn/sharding_test.exs
+++ b/exla/test/exla/defn/sharding_test.exs
@@ -950,10 +950,52 @@ defmodule EXLA.Defn.ShardingTest do
           [128, 130]
         ])
 
-      for {r, partition_idx} <- Enum.with_index(results) do
-        assert_equal(r, expected_result)
-        assert r.data.buffer.device_id == partition_idx
-      end
+      Enum.zip_with([results, 0..3], fn [result, i] ->
+        assert_equal(result, expected_result)
+        assert result.data.buffer.device_id == i
+      end)
+    end
+
+    @moduletag :multi_device
+    test "can return partially sharded results" do
+      fun = fn x, y -> Nx.add(x, y) end
+
+      mesh = %Mesh{name: "mesh", shape: {2, 2}}
+      # Inputs sharded on both axes
+      input_shardings = [%{0 => [0], 1 => [1]}, %{0 => [0], 1 => [1]}]
+      # Output: sharded only on axis 0 (dim 1 replicated) -> each partition gets {4, 2}
+      output_shardings = [%{0 => [0]}]
+
+      # Logical x: 8x2, y: 8x2. Each partition gets {4, 1} of each
+      args = [
+        [Nx.tensor([[0], [1], [2], [3]]), Nx.tensor([[100], [101], [102], [103]])],
+        [Nx.tensor([[10], [11], [12], [13]]), Nx.tensor([[110], [111], [112], [113]])],
+        [Nx.tensor([[4], [5], [6], [7]]), Nx.tensor([[104], [105], [106], [107]])],
+        [Nx.tensor([[14], [15], [16], [17]]), Nx.tensor([[114], [115], [116], [117]])]
+      ]
+
+      results =
+        EXLA.shard_jit(fun, mesh,
+          input_shardings: input_shardings,
+          output_shardings: output_shardings
+        ).(args)
+
+      assert length(results) == 4
+
+      # Partially sharded output: dim 0 sharded on axis 0, dim 1 not in output spec
+      # Each device returns its local shard {4, 1} (x+y computed locally)
+      # Dev0: col0 rows 0-3, Dev1: col1 rows 0-3, Dev2: col0 rows 4-7, Dev3: col1 rows 4-7
+      expected_results = [
+        Nx.tensor([[100], [102], [104], [106]]),
+        Nx.tensor([[120], [122], [124], [126]]),
+        Nx.tensor([[108], [110], [112], [114]]),
+        Nx.tensor([[128], [130], [132], [134]])
+      ]
+
+      Enum.zip_with([results, expected_results, 0..3], fn [result, expected, i] ->
+        assert_equal(result, expected)
+        assert result.data.buffer.device_id == i
+      end)
     end
   end
 end