@@ -153,3 +153,59 @@ def _lower_compare(lhs: ir.Value, rhs: ir.Value) -> ir.Value:
         arith.select(rhs_longer, arith.constant(i32, -1), zero_i32),
     )
     return arith.select(found, byte_diff, len_cmp)
+
+
+def _bytes_equal(
+    a_data: ir.Value, a_off: ir.Value, b_data: ir.Value, n: ir.Value
+) -> ir.Value:
+    """i1: whether ``a_data[a_off:a_off+n] == b_data[0:n]`` byte-for-byte.
+
+    Early-exits (``scf.while``) at the first differing byte; callers must pass
+    ``n == 0`` when the range would read out of bounds.
+    """
+    i1 = ir.IntegerType.get_signless(1)
+    zero_i64 = arith.constant(T.i64(), 0)
+    loop = scf.WhileOp([T.i64(), i1], [zero_i64, arith.constant(i1, 1)])
+    before = loop.before.blocks.append(T.i64(), i1)
+    with ir.InsertionPoint(before):
+        idx, eq = before.arguments
+        in_bounds = arith.cmpi(arith.CmpIPredicate.ult, idx, n)
+        scf.condition(arith.andi(in_bounds, eq), [idx, eq])
+    after = loop.after.blocks.append(T.i64(), i1)
+    with ir.InsertionPoint(after):
+        idx, _eq = after.arguments
+        av = _byte_at(a_data, arith.addi(a_off, idx))
+        bv = _byte_at(b_data, idx)
+        next_idx = arith.addi(idx, arith.constant(T.i64(), 1))
+        scf.YieldOp([next_idx, arith.cmpi(arith.CmpIPredicate.eq, av, bv)])
+    return loop.results[1]
+
+
+def _lower_startswith(str_view: ir.Value, prefix: ir.Value) -> ir.Value:
+    """``startswith(prefix)`` -> ``i1`` (byte prefix match)."""
+    i1 = ir.IntegerType.get_signless(1)
+    src_nb = llvm.zext(T.i64(), _view_nbytes(str_view))
+    pfx_nb = llvm.zext(T.i64(), _view_nbytes(prefix))
+    too_long = arith.cmpi(arith.CmpIPredicate.ugt, pfx_nb, src_nb)
+    # Compare 0 bytes when the prefix cannot fit, so the scan never over-reads
+    # the source; the result is false anyway.
+    n = arith.select(too_long, arith.constant(T.i64(), 0), pfx_nb)
+    eq = _bytes_equal(
+        _view_data(str_view), arith.constant(T.i64(), 0), _view_data(prefix), n
+    )
+    return arith.andi(arith.xori(too_long, arith.constant(i1, 1)), eq)
+
+
+def _lower_endswith(str_view: ir.Value, suffix: ir.Value) -> ir.Value:
+    """``endswith(suffix)`` -> ``i1`` (byte suffix match)."""
+    i1 = ir.IntegerType.get_signless(1)
+    src_nb = llvm.zext(T.i64(), _view_nbytes(str_view))
+    sfx_nb = llvm.zext(T.i64(), _view_nbytes(suffix))
+    too_long = arith.cmpi(arith.CmpIPredicate.ugt, sfx_nb, src_nb)
+    zero_i64 = arith.constant(T.i64(), 0)
+    # Clamp both the start offset and the length when the suffix cannot fit so
+    # the scan stays in bounds; the result is false anyway.
+    n = arith.select(too_long, zero_i64, sfx_nb)
+    offset = arith.select(too_long, zero_i64, arith.subi(src_nb, sfx_nb))
+    eq = _bytes_equal(_view_data(str_view), offset, _view_data(suffix), n)
+    return arith.andi(arith.xori(too_long, arith.constant(i1, 1)), eq)
@@ -2,11 +2,11 @@
 # SPDX-License-Identifier: Apache-2.0
 """Typing and lowering for ``mlir_string`` operations.
 
-Covers ``len`` (UTF-8 character count) and the string comparison operators
-(``==``/``!=``/``<``/``<=``/``>``/``>=``) over ``mlir_string``,
-``Masked(mlir_string)`` and string-literal operands; masked operands yield a
-``Masked`` result whose validity is the AND of the operand validity bits. All
-lowerings are pure MLIR (see
+Covers ``len`` (UTF-8 character count), the string comparison operators
+(``==``/``!=``/``<``/``<=``/``>``/``>=``), and the affix methods
+``startswith``/``endswith`` over ``mlir_string``, ``Masked(mlir_string)`` and
+string-literal operands; masked operands yield a ``Masked`` result whose validity
+is the AND of the operand validity bits. All lowerings are pure MLIR (see
 :mod:`cudf.core.udf.mlir_backend.string_lowering_impl`). Registered with
 ``numba_cuda_mlir`` at import time via :func:`_register`.
 """
@@ -20,8 +20,10 @@
 from numba_cuda_mlir._mlir import ir
 from numba_cuda_mlir._mlir.dialects import arith, llvm
 from numba_cuda_mlir.extending import lowering_registry, typing_registry
+from numba_cuda_mlir.lowering_utilities import DeferredMethodCall
 from numba_cuda_mlir.numba_cuda.typing.templates import (
     AbstractTemplate,
+    AttributeTemplate,
     Signature,
 )
 from numba_cuda_mlir.typing import signature as nb_signature
@@ -71,6 +73,12 @@
     (MLIRStringType, _masked_string),
 )
 
+# affix method name -> its pure-MLIR impl (one string arg -> boolean).
+_BOOL_METHODS = {
+    "startswith": _impl._lower_startswith,
+    "endswith": _impl._lower_endswith,
+}
+
 
 def _is_plain_string(ty: types.Type) -> bool:
     """Whether ``ty`` is an ``mlir_string`` or a compile-time string literal."""
@@ -161,6 +169,70 @@ def generic(self, args, kws):
     return _CmpOpTemplate
 
 
+def _make_method_attr(
+    attrname: str, retty: types.Type, masked: bool
+) -> object:
+    """Build a ``resolve_<method>`` for a one-string-arg method.
+
+    Parameters
+    ----------
+    attrname : str
+        Method name (e.g. ``"startswith"``).
+    retty : types.Type
+        The unmasked return type.
+    masked : bool
+        Whether the receiver is ``Masked(mlir_string)`` (wraps the result).
+
+    Returns
+    -------
+    callable
+        A ``resolve_`` function for the attribute template.
+    """
+    result_ty = MaskedType(retty) if masked else retty
+    recvr = _masked_string if masked else mlir_string
+
+    class _MethodTemplate(AbstractTemplate):
+        key = f"MLIRString.{attrname}.{'m' if masked else 's'}"
+
+        def generic(self, args, kws):
+            """Resolve ``s.<method>(other)`` for a string ``other``."""
+            if len(args) == 1 and not kws and _is_string_arg(args[0]):
+                return nb_signature(result_ty, args[0], recvr=self.this)
+            return None
+
+    def resolve(self, mod):
+        return types.BoundFunction(_MethodTemplate, recvr)
+
+    return resolve
+
+
+@typing_registry.register_attr
+class MLIRStringAttrs(AttributeTemplate):
+    """Attribute typing for ``mlir_string`` methods."""
+
+    key = mlir_string
+
+
+@typing_registry.register_attr
+class MaskedMLIRStringAttrs(AttributeTemplate):
+    """Attribute typing for ``Masked(mlir_string)`` methods."""
+
+    key = _masked_string
+
+
+for _name in _BOOL_METHODS:
+    setattr(
+        MLIRStringAttrs,
+        f"resolve_{_name}",
+        _make_method_attr(_name, types.boolean, masked=False),
+    )
+    setattr(
+        MaskedMLIRStringAttrs,
+        f"resolve_{_name}",
+        _make_method_attr(_name, types.boolean, masked=True),
+    )
+
+
 def _lower_len(
     builder: MLIRLower, target: Var, args: list[Var], kwargs: list
 ) -> None:
@@ -286,8 +358,29 @@ def _lower(builder, target, args, kwargs):
     return _lower
 
 
+def _make_lower_binary(impl_fn: object) -> object:
+    """Build the ``lower_getattr`` for a one-string-arg method (startswith/endswith).
+
+    The returned getattr defers to a method lowering that handles plain, literal
+    and masked operands uniformly and packs a ``Masked`` result when required.
+    """
+
+    def _impl_lower(builder, target, args, kwargs):
+        valid_ty = builder.get_mlir_type(types.boolean)
+        true_val = arith.constant(valid_ty, 1)
+        recv, vr = _view_and_valid(builder, args[0], true_val)
+        other, vo = _view_and_valid(builder, args[1], true_val)
+        result = impl_fn(recv, other)
+        _store_maybe_masked(builder, target, result, arith.andi(vr, vo))
+
+    def _getattr(context, builder, target, value, attr=None):
+        builder.store_var(target, DeferredMethodCall(value, _impl_lower))
+
+    return _getattr
+
+
 def _register() -> None:
-    """Register ``len`` + comparison typing and lowering with ``numba_cuda_mlir``."""
+    """Register ``len`` + comparison + affix typing/lowering with numba_cuda_mlir."""
     typing_registry.register_global(len)(LenMLIRStringTemplate)
     lowering_registry.lower(len, mlir_string)(_lower_len)
     lowering_registry.lower(len, MaskedType)(_lower_masked_len)
@@ -298,5 +391,11 @@ def _register() -> None:
         for lhs_ty, rhs_ty in _STR_COMBOS:
             lowering_registry.lower(cmp_op, lhs_ty, rhs_ty)(impl)
 
+    # startswith / endswith methods on mlir_string and Masked(mlir_string)
+    for name, impl_fn in _BOOL_METHODS.items():
+        getattr_fn = _make_lower_binary(impl_fn)
+        lowering_registry.lower_getattr(mlir_string, name)(getattr_fn)
+        lowering_registry.lower_getattr(_masked_string, name)(getattr_fn)
+
 
 _register()
@@ -367,3 +367,66 @@ def k(o, ov, sa, sav, sb, sbv):
     for i in range(n):
         if got_valid[i]:
             assert bool(out.get()[i]) is (a[i] == b[i])
+
+
+def test_string_affix():
+    """``startswith`` / ``endswith`` against a literal, over an array."""
+    strings = ["abc", "abcd", "xabc", "ab", "", "cab", "h\u00e9llo"]
+    arr, _k = _make_mlir_string_array(strings)
+    n = len(strings)
+    starts = cp.zeros(n, dtype=np.bool_)
+    ends = cp.zeros(n, dtype=np.bool_)
+
+    @cuda.jit(
+        types.void(
+            types.boolean[::1],
+            types.boolean[::1],
+            types.CPointer(mlir_string),
+        ),
+        extensions=[mlir_string_arg_handler],
+    )
+    def k(so, eo, s):
+        i = cuda.grid(1)
+        if i < so.size:
+            so[i] = s[i].startswith("ab")
+            eo[i] = s[i].endswith("bc")
+
+    with MLIRNumbaCudaConfig():
+        k[1, n](starts, ends, arr)
+    cuda.synchronize()
+    assert starts.get().tolist() == [s.startswith("ab") for s in strings]
+    assert ends.get().tolist() == [s.endswith("bc") for s in strings]
+
+
+def test_masked_string_affix_propagates_validity():
+    """``Masked(str).startswith(literal)`` -> ``Masked(bool)``; validity carried."""
+    strings = ["abc", "abd", "xyz"]
+    valid = [True, False, True]
+    arr, _k = _make_mlir_string_array(strings)
+    n = len(strings)
+    out = cp.zeros(n, dtype=np.bool_)
+    out_valid = cp.zeros(n, dtype=np.bool_)
+
+    @cuda.jit(
+        types.void(
+            types.boolean[::1],
+            types.boolean[::1],
+            types.CPointer(mlir_string),
+            types.boolean[::1],
+        ),
+        extensions=[mlir_string_arg_handler],
+    )
+    def k(o, ov, s, sv):
+        i = cuda.grid(1)
+        if i < o.size:
+            r = Masked(s[i], sv[i]).startswith("ab")
+            o[i] = r.value
+            ov[i] = r.valid
+
+    with MLIRNumbaCudaConfig():
+        k[1, n](out, out_valid, arr, cp.array(valid, dtype=np.bool_))
+    cuda.synchronize()
+    assert out_valid.get().tolist() == valid
+    for i in range(n):
+        if valid[i]:
+            assert bool(out.get()[i]) is strings[i].startswith("ab")