Kernels:

---

jburtoft

/

gte-encoder-neuron-kernels

Like 0

"""On-device accuracy for NewBERT embedding models (gte-*-en-v1.5, gte-multilingual-base): CLS embeddings of the
patched model vs the stock Hugging Face model on CPU in fp32 (exact GELU), and vs stock PyTorch Native.

Gates: mean per-text cosine >= min(0.9999, stock bf16 mean) (patched vs CPU fp32; stock itself is below 0.9999 at L=1024); min per-text cosine within COS_MIN_SLACK of stock
bf16 PyTorch Native's min (an absolute 0.9999 floor on the minimum is not met by stock bf16 itself);
retrieval top-1 per query equal to the CPU fp32 ranking at least as often as stock Native; padding isolation
bit-exact; fused ops actually traced.

  GTE_TEST_MODEL=Alibaba-NLP/gte-large-en-v1.5 NEURON_LOGICAL_NC_CONFIG=2 python -m pytest -m neuron -s tests/test_device_embed.py
"""
import json
import os
import sys

import pytest

from conftest import load_package
from test_device import TOPICS, _pairs

torch = pytest.importorskip("torch")
pytest.importorskip("torch_neuronx")
transformers = pytest.importorskip("transformers")
pytestmark = pytest.mark.neuron

MODEL = os.environ.get("GTE_TEST_MODEL", "Alibaba-NLP/gte-large-en-v1.5")
RESULTS = os.environ.get("GTE_RESULTS", os.path.join(os.path.dirname(os.path.abspath(__file__)),
                                                     "device_results.jsonl"))
SHAPES = [(256, 32), (128, 32), (512, 3), (1024, 8), (768, 3)]
COS_MEAN_MIN = 0.9999
ATTN_LENGTHS = tuple(int(v) for v in os.environ.get("GTE_ATTN_LENGTHS", "128,256,384,512,768,1024").split(","))
COS_MIN_SLACK = 1e-3     # worst single text may trail stock's worst by this much (bf16 noise over 24 layers)


def _load(dtype):
    m = transformers.AutoModel.from_pretrained(MODEL, trust_remote_code=True, attn_implementation="eager").eval()
    if dtype is not None:
        for mod in m.modules():
            for cn, c in mod.named_children():
                if "GELU" in type(c).__name__:
                    setattr(mod, cn, torch.nn.GELU(approximate="tanh"))
        m = m.to(dtype)
    return m


class _CLS(torch.nn.Module):
    def __init__(self, m):
        super().__init__()
        self.m = m

    def forward(self, ids, msk):
        return self.m(input_ids=ids, attention_mask=msk).last_hidden_state[:, 0]


@pytest.fixture(scope="module")
def setup():
    import torch._dynamo as dyn
    dyn.config.cache_size_limit = 64
    pkg = load_package()
    to = sys.modules["gte_encoder_neuron_kernels.torch_ops"]
    calls = {}
    for name in ("attn_block", "gated_mlp_ln", "res_ln", "ln"):
        def wrap(f, name=name):
            def g(*a, **k):
                calls[name] = calls.get(name, 0) + 1
                return f(*a, **k)
            return g
        setattr(to, name, wrap(getattr(to, name)))
    tok = transformers.AutoTokenizer.from_pretrained(MODEL, trust_remote_code=True)
    cpu = _CLS(_load(None))
    dev = torch.device("neuron")
    stock = _CLS(_load(torch.bfloat16)).to(dev)
    p = _load(torch.bfloat16)
    n = pkg.patch_model(p, attn_lengths=ATTN_LENGTHS)
    patched = _CLS(p).to(dev)
    return dict(pkg=pkg, tok=tok, cpu=cpu, dev=dev, stock=stock, patched=patched, n_layers=n, calls=calls)


def _emb(model, ids, msk, bs, dev):
    cm = torch.compile(model, backend="neuron", dynamic=False)
    with torch.no_grad():
        out = [cm(ids[i:i + bs].to(dev), msk[i:i + bs].to(dev)).float().cpu() for i in range(0, ids.shape[0], bs)]
    return torch.cat(out), cm


def _cos(a, b):
    return torch.nn.functional.cosine_similarity(a, b, dim=-1)


@pytest.mark.parametrize("L,bs", SHAPES)
def test_embed_accuracy_and_isolation(setup, L, bs):
    tok, cpu, dev = setup["tok"], setup["cpu"], setup["dev"]
    per = 8
    Q, Ps = _pairs(L, per)
    queries = list(TOPICS)
    texts = queries + Ps                          # 8 queries + 64 passages
    n = len(texts) - len(texts) % bs
    texts = texts[:n]
    e = tok(texts, return_tensors="pt", max_length=L, padding="max_length", truncation=True)
    ids, msk = e["input_ids"], e["attention_mask"]
    with torch.no_grad():
        ref = torch.cat([cpu(ids[i:i + 8], msk[i:i + 8]).float() for i in range(0, n, 8)])
    e_stock, _ = _emb(setup["stock"], ids, msk, bs, dev)
    setup["calls"].clear()
    e_nki, cm = _emb(setup["patched"], ids, msk, bs, dev)
    traced = dict(setup["calls"])

    cos_nki = _cos(e_nki, ref)
    cos_stock = _cos(e_stock, ref)

    def top1(E):
        En = torch.nn.functional.normalize(E, dim=-1)
        q = En[:len(queries)]
        p = En[len(queries):]
        return (q @ p.T).argmax(-1)

    t_ref, t_stock, t_nki = top1(ref), top1(e_stock), top1(e_nki)
    agree_stock = float((t_stock == t_ref).float().mean())
    agree_nki = float((t_nki == t_ref).float().mean())

    b1 = ids[:bs].clone(), msk[:bs].clone()
    if bs > 1:
        perm = torch.arange(n - 1, n - bs, -1)
        b1[0][1:], b1[1][1:] = ids[perm], msk[perm]
    with torch.no_grad():
        y0 = cm(ids[:bs].to(dev), msk[:bs].to(dev))[0].float().cpu()
        y1 = cm(b1[0].to(dev), b1[1].to(dev))[0].float().cpu()
    isolated = bool(torch.equal(y0, y1))
    lens = msk.sum(-1).numpy()
    worst = torch.argsort(cos_nki)[:4].tolist()
    rec_worst = [dict(i=int(i), len=int(lens[i]), cos_nki=float(cos_nki[i]), cos_stock=float(cos_stock[i]))
                 for i in worst]
    rec = dict(model=MODEL, kind="embed", lnc=setup["pkg"].LNC, L=L, bs=bs, texts=n,
               real_len=[int(lens.min()), int(lens.max())], layers_patched=setup["n_layers"],
               cos_min_stock=float(cos_stock.min()), cos_mean_stock=float(cos_stock.mean()),
               cos_min_nki=float(cos_nki.min()), cos_mean_nki=float(cos_nki.mean()),
               top1_agree_stock=agree_stock, top1_agree_nki=agree_nki, padding_isolated=isolated,
               traced_ops=traced, worst=rec_worst)
    print(json.dumps(rec), flush=True)
    with open(RESULTS, "a") as f:
        f.write(json.dumps(rec) + "\n")
    nl = setup["n_layers"]
    if L in ATTN_LENGTHS:
        assert traced.get("attn_block", 0) >= nl, rec
    assert traced.get("gated_mlp_ln", 0) >= nl, rec
    assert float(cos_nki.mean()) >= min(COS_MEAN_MIN, float(cos_stock.mean())), rec
    assert float(cos_nki.min()) >= float(cos_stock.min()) - COS_MIN_SLACK, rec
    assert agree_nki >= agree_stock, rec
    assert isolated, rec