@@ -200,29 +200,49 @@ def _lower_masked_na_compare(builder, target, args, kwargs, *, is_null):
     builder.store_var(target, valid)
 
 
-# datetime64 / timedelta64 ``+``/``-`` need numba_cuda_mlir's ``datetime``
-# lowering (which scales by unit), not a raw i64 op.
-def _needs_datetimelike_delegate(op, ty1, ty2):
-    if op not in (operator.add, operator.sub):
+def _needs_datetimelike_delegate(
+    op: Callable, ty1: types.Type, ty2: types.Type
+) -> bool:
+    """Whether a masked binary ``op`` over datetime64/timedelta64 operands must
+    delegate to numba_cuda_mlir's unit-aware scalar lowering.
+
+    A raw i64 op on the payloads is wrong whenever units matter:
+
+    * ``+``/``-`` must scale the operands to a common unit before combining.
+    * comparisons must likewise compare on a common unit (``1 s < 2 ns`` is
+      ``1_000_000_000 < 2``, not ``1 < 2``).
+
+    So delegate for add/sub and all comparison operators when either operand is
+    temporal.
+    """
+    if op not in (operator.add, operator.sub, *comparison_ops):
         return False
     return isinstance(
         ty1, (types.NPDatetime, types.NPTimedelta)
     ) or isinstance(ty2, (types.NPDatetime, types.NPTimedelta))
 
 
 def _apply_masked_datetimelike_binary(
-    builder,
-    target,
-    target_type,
-    v1,
-    v2,
-    result_valid,
-    op,
-    ty1,
-    ty2,
-    ref_var,
-):
-    """TODO: write docstring."""
+    builder: MLIRLower,
+    target: Var,
+    target_type: MaskedType,
+    v1: mlir_ir.Value,
+    v2: mlir_ir.Value,
+    result_valid: mlir_ir.Value,
+    op: Callable,
+    ty1: types.Type,
+    ty2: types.Type,
+) -> None:
+    """Lower a temporal masked binary op by delegating to the registered
+    unit-aware scalar builder, then repacking with the operand validity.
+
+    The two payloads are staged into fresh typed IR vars (with their true
+    datetime/timedelta types), handed to numba_cuda_mlir's scalar lowering for
+    ``op`` (which handles unit scaling), and the result is converted to the
+    target Masked's value type and packed with ``result_valid``. Staging vars
+    are keyed off the unique ``target`` so distinct temporal ops in one
+    expression never share a typemap slot.
+    """
     ret_ty = target_type.value_type
     nb_sig = nb_typing.signature(ret_ty, ty1, ty2)
     cg = builder.get_registered_builder(op, nb_sig)
@@ -231,9 +251,10 @@ def _apply_masked_datetimelike_binary(
             f"No MLIR lowering for masked {op!r} with {ty1}, {ty2}; "
             f"signature {nb_sig}"
         )
-    in1 = _make_temp_var(builder, ref_var, "mdt_l", ty1)
-    in2 = _make_temp_var(builder, ref_var, "mdt_r", ty2)
-    outv = _make_temp_var(builder, ref_var, "mdt_o", ret_ty)
+    tag = getattr(op, "__name__", "op")
+    in1 = _make_temp_var(builder, target, f"dt_{tag}_l", ty1)
+    in2 = _make_temp_var(builder, target, f"dt_{tag}_r", ty2)
+    outv = _make_temp_var(builder, target, f"dt_{tag}_o", ret_ty)
     builder.store_var(in1, convert(v1, builder.get_mlir_type(ty1)))
     builder.store_var(in2, convert(v2, builder.get_mlir_type(ty2)))
     cg(builder, outv, [in1, in2], ())
@@ -253,18 +274,16 @@ def _apply_masked_binary_op(
     *,
     inner_ty1: types.Type | None = None,
     inner_ty2: types.Type | None = None,
-    ref_var: Var | None = None,
 ) -> None:
     """Apply ``op(v1, v2)`` to two scalar MLIR values, convert the result to
     the target Masked's value type, and pack it with the given validity bit.
     Numeric/boolean only at this layer.
     """
-    # datetime/timedelta add/sub: delegate to the unit-aware scalar
-    # lowering when we know the operand inner types.
+    # datetime/timedelta add/sub/comparisons: delegate to the unit-aware
+    # scalar lowering when we know the operand inner types.
     if (
         inner_ty1 is not None
         and inner_ty2 is not None
-        and ref_var is not None
         and _needs_datetimelike_delegate(op, inner_ty1, inner_ty2)
     ):
         _apply_masked_datetimelike_binary(
@@ -277,7 +296,6 @@ def _apply_masked_binary_op(
             op,
             inner_ty1,
             inner_ty2,
-            ref_var,
         )
         return
 
@@ -320,7 +338,6 @@ def _lower(
             op,
             inner_ty1=ty1,
             inner_ty2=ty2,
-            ref_var=args[0],
         )
 
     return _lower
@@ -383,7 +400,6 @@ def _lower(
                 op,
                 inner_ty1=m_inner_ty,
                 inner_ty2=s_inner_ty,
-                ref_var=m_var,
             )
         else:
             _apply_masked_binary_op(
@@ -396,7 +412,6 @@ def _lower(
                 op,
                 inner_ty1=s_inner_ty,
                 inner_ty2=m_inner_ty,
-                ref_var=m_var,
             )
 
     return _lower
@@ -424,14 +439,15 @@ def _make_temp_var(
 ) -> Var:
     """Create a fresh typed IR ``Var`` for staging an intermediate value.
 
-    Used to feed the masked payload into a registered scalar lowering (which
+    Used to feed a masked payload into a registered scalar lowering (which
     operates on plain IR vars) and to receive its result. The name is derived
-    from ``base_var`` plus ``name_suffix`` so typemap keys stay unique when one
-    operand feeds several unary calls in a single expression.
+    from ``base_var`` plus ``name_suffix``; callers pass a unique ``base_var``
+    (e.g. the result ``target``) so typemap keys stay distinct when one operand
+    feeds several op lowerings in a single expression.
     """
     scope = getattr(base_var, "scope", None)
     loc = getattr(base_var, "loc", None)
-    name = f"$masked_uop_{base_var.name}_{name_suffix}"
+    name = f"$masked_tmp_{base_var.name}_{name_suffix}"
     temp = numba_ir.Var(scope=scope, name=name, loc=loc)
     builder.fndesc.typemap[temp.name] = numba_type
     return temp
 
@@ -817,7 +817,7 @@ def _bool(v):
 
 
 def test_masked_datetime_minus_datetime_is_timedelta():
-    """TODO: write docstring."""
+    """Masked ``datetime - datetime`` yields a Masked timedelta."""
 
     @cuda.jit(
         types.void(
@@ -850,7 +850,7 @@ def k(out_v, out_valid, a, av, b, bv):
 
 
 def test_masked_datetime_plus_timedelta_is_datetime():
-    """TODO: write docstring."""
+    """Masked ``datetime + timedelta`` yields a Masked datetime."""
 
     @cuda.jit(
         types.void(
@@ -870,7 +870,7 @@ def k(out_v, a, av, t, tv):
 
 
 def test_masked_timedelta_plus_timedelta_is_timedelta():
-    """TODO: write docstring."""
+    """Masked ``timedelta + timedelta`` yields a Masked timedelta."""
 
     @cuda.jit(
         types.void(
@@ -889,12 +889,9 @@ def k(out_v, a, av, b, bv):
     assert int(out_v.get().view("int64")[0]) == 420
 
 
-@pytest.mark.parametrize(
-    "op,ref",
-    [(operator.lt, lambda a, b: a < b), (operator.gt, lambda a, b: a > b)],
-)
-def test_masked_datetime_comparison(op, ref):
-    """TODO: write docstring."""
+@pytest.mark.parametrize("op", [operator.lt, operator.gt])
+def test_masked_datetime_comparison(op):
+    """Masked datetime comparisons evaluate on the operand payloads."""
 
     @cuda.jit(
         types.void(
@@ -910,11 +907,41 @@ def k(out, a, av, b, bv):
 
     out = cp.zeros(1, dtype=np.bool_)
     _launch(k, out, _dt_in([400]), _bool(True), _dt_in([1000]), _bool(True))
-    assert bool(out.get()[0]) == ref(400, 1000)
+    assert bool(out.get()[0]) == op(400, 1000)
+
+
+def test_masked_datetime_comparison_mixed_units():
+    """Mixed-unit comparisons scale to a common unit rather than comparing
+    raw payloads: ``1 s`` vs ``2 ns`` compares as ``1_000_000_000`` vs ``2``
+    (True for ``>``), not the raw ``1 < 2``.
+    """
+    dt_s = types.NPDatetime("s")
+    dt_ns = types.NPDatetime("ns")
+
+    @cuda.jit(
+        types.void(
+            types.boolean[::1],
+            dt_s[::1],
+            types.boolean[::1],
+            dt_ns[::1],
+            types.boolean[::1],
+        )
+    )
+    def k(out, a, av, b, bv):
+        out[0] = (Masked(a[0], av[0]) > Masked(b[0], bv[0])).value
+
+    a = cp.asarray(np.array([1], dtype="int64")).view("datetime64[s]")
+    b = cp.asarray(np.array([2], dtype="int64")).view("datetime64[ns]")
+    out = cp.zeros(1, dtype=np.bool_)
+    _launch(k, out, a, _bool(True), b, _bool(True))
+    # 1 second (1e9 ns) > 2 ns is True; a raw i64 compare would give 1 > 2 = False.
+    assert bool(out.get()[0]) is True
 
 
 def test_masked_datetime_arith_validity_propagates():
-    """TODO: write docstring."""
+    """Temporal arithmetic ANDs operand validity: an NA operand poisons the
+    result.
+    """
 
     @cuda.jit(
         types.void(
 
@@ -55,7 +55,9 @@ def test_masked_type_unsupported_value_becomes_poison():
 
 @pytest.mark.parametrize("unit", ["ns", "us", "ms", "s"])
 def test_masked_datetime_timedelta_not_poisoned(unit):
-    """TODO: write docstring."""
+    """datetime64/timedelta64 payloads are supported MaskedType value types
+    (not poisoned) across the units cuDF columns use.
+    """
     dt = MaskedType(types.NPDatetime(unit))
     td = MaskedType(types.NPTimedelta(unit))
     assert dt.value_type == types.NPDatetime(unit)