SYCL: --batch slots honour the request's repeat / frequency / presence penalties: each penalised slot's reply tail is staged before its window (set_slot_history), sample_rows builds the per-row history with slot_penalty_rows (the CUDA/HIP/Metal logic of 2678c9732), pen_from names where a continued request's reply began; a plain greedy slot still keeps the window's argmax token. logit_bias stays refused on SYCL. slot_history_test (CPU) added to the SYCL targets

sycl/CMakeLists.txt

@@ -463,6 +463,9 @@ if(STRATA_SYCL_PARITY)
   add_executable(pipeline_auto_test ${CMAKE_CURRENT_SOURCE_DIR}/src/core/pipeline_auto_test.cpp)   # the default of --pipeline-windows (no GPU)
   target_include_directories(pipeline_auto_test PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/include)
   add_test(NAME pipeline_auto_test COMMAND pipeline_auto_test)
+  add_executable(slot_history_test ${CMAKE_CURRENT_SOURCE_DIR}/src/kernels/slot_history_test.cpp)   # the penalty-history rows of a batch window, per slot (CPU, no model)
+  target_include_directories(slot_history_test PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/include)
+  add_test(NAME slot_history_test COMMAND slot_history_test)
   if(_i8_esimd)   # DPAS s8 8x16x32 (Xe2): not for an Alchemist / Xe-HPC AOT build, nor with STRATA_SYCL_I8_ESIMD=OFF
     add_executable(esimd_dpas_probe ${CMAKE_CURRENT_SOURCE_DIR}/src/kernels/esimd_dpas_probe.cpp)   # ESIMD DPAS peak (registers only)
     add_executable(esimd_i8_bench ${CMAKE_CURRENT_SOURCE_DIR}/src/kernels/esimd_i8_bench.cpp)   # ESIMD int8 DPAS GEMM on pre-decoded weights

sycl/include/strata/core/verify.hpp

@@ -178,13 +178,19 @@ public:
     /// -1 = an error (err).  Serves every layer that has rung so far.
     int batch_poll(PoolMultiFn pool, void* user, std::string& err);
     bool batch_busy() const { return b_running_; }
-    /// A slot's sampling (temperature / top_p / top_k / min_p / seed; penalties are not applied in batch windows):
-    /// its row is drawn again on the last stage with Philox(seed, position), as a solo window draws it.  Greedy by
-    /// default.  Set on the first stage, it reaches the last.
+    /// A slot's sampling (temperature / top_p / top_k / min_p / seed and the repeat / frequency / presence penalties,
+    /// whose window comes from set_slot_history): its row is drawn again on the last stage with Philox(seed, position),
+    /// as a solo window draws it.  Greedy by default.  Set on the first stage, it reaches the last.
     void set_slot_sampling(int slot, const strata::kernels::SamplerParams& sp) {
         if (slot >= 0 && slot < (int) slot_sp_.size()) slot_sp_[(size_t) slot] = sp;
         if (next_) next_->set_slot_sampling(slot, sp);
     }
+    /// The slot's reply tail for its penalty window: the `n` tokens its session consumed since the reply began (the last penalty_last_n are used), the
+    /// fed-back head not included.  Set before every batch window that holds the slot while it has a penalty; reaches the last stage.
+    void set_slot_history(int slot, const int32_t* tail, int64_t n) {
+        if (slot >= 0 && slot < (int) slot_tail_.size()) slot_tail_[(size_t) slot].assign(tail, tail + (n > 0 ? n : 0));
+        if (next_) next_->set_slot_history(slot, tail, n);
+    }
     const int32_t* batch_out() const { return b_out_; }
     bool last_stage() const { return g_ != nullptr && le_ == g_->n_layers; }
     /// commit() returns without waiting for its graph (a single-GPU session sets it): the next window follows it on
@@ -312,6 +318,9 @@ private:
     std::chrono::steady_clock::time_point b_last_;
     int32_t b_out_[8] = {};
     std::vector<strata::kernels::SamplerParams> slot_sp_;
+    std::vector<std::vector<int32_t>> slot_tail_;        ///< per slot: its reply tail for the penalty window (set_slot_history)
+    int32_t* slot_hist_d_ = nullptr;                     ///< the penalty-history rows of the current batch window (device)
+    size_t slot_hist_cap_ = 0;
     bool sample_rows(int S, std::string& err);   ///< the sampled slots' rows of the last batch window
     int32_t* h_commitb_ = nullptr; int32_t* m_commitb_ = nullptr;   // per slot [1, 0, pos, -1 ..], stride 2 + max_t
     int32_t* commitb_ = nullptr;

sycl/include/strata/kernels/sampler.hpp

@@ -96,6 +96,31 @@ inline void penalty_rows(const int32_t* tail, int64_t n_tail, const int32_t* win
     }
 }
 
+/// The penalty-history rows of a BATCH window, on the host.  Row t of the window feeds `tokens[t]` into slot `rows[t]`; its row of `out` (h slots, most
+/// recent LAST, -1 in unused front slots) is the last `h` tokens of that slot's reply tail (`tails[slot]`, `n_tails[slot]` tokens: what the slot's
+/// session consumed since the reply began, the fed-back head not yet included) followed by the tokens fed so far by THIS window's rows of the same slot
+/// up to and including row t (a slot may hold several rows: MTP drafts).  Exactly what `penalty_rows` gives a solo window, per slot.
+inline void slot_penalty_rows(const int* rows, const int32_t* tokens, int S, const int32_t* const* tails, const int64_t* n_tails, int h, int32_t* out) {
+    for (int t = 0; t < S; ++t) {
+        int32_t* row = out + (size_t) t * (size_t) h;
+        const int slot = rows[t];
+        const int32_t* tail = tails[slot];
+        const int64_t n_tail = tail != nullptr ? n_tails[slot] : 0;
+        int64_t in_window = 0;                                     // this window's rows of the slot up to t
+        for (int u = 0; u <= t; ++u) in_window += rows[u] == slot;
+        const int64_t avail = n_tail + in_window;
+        const int take = (int) std::min<int64_t>(h, avail);
+        std::fill(row, row + (h - take), -1);
+        for (int j = 0; j < take; ++j) {
+            const int64_t i = avail - take + j;                     // index into (tail, the slot's window tokens)
+            if (i < n_tail) { row[h - take + j] = tail[i]; continue; }
+            int64_t k = i - n_tail, seen = 0;
+            for (int u = 0; u <= t; ++u)
+                if (rows[u] == slot) { if (seen == k) { row[h - take + j] = tokens[u]; break; } ++seen; }
+        }
+    }
+}
+
 // Gated proposal verification is not implemented by the SYCL backend.
 bool sample_tokens_proposal(const float* logits, int n_tokens, int n_vocab, const int* history, int history_len,
                             const SamplerParams& p, const int32_t* dtok, const int32_t* qbuf, int n_q, int* out,

sycl/src/core/verify.cpp

@@ -423,6 +423,7 @@ Verifier::~Verifier() try {
         slot.compare_exchange_strong(me, nullptr);
     }
     if (cs_) cs_->wait();
+    if (slot_hist_d_ != nullptr && cs_) { try { sycl::free(slot_hist_d_, *cs_); } catch (...) {} slot_hist_d_ = nullptr; }
     if (sh_cs_) sh_cs_->wait();
     for (auto& e : exec_)
         if (e) delete (e);
@@ -2943,6 +2944,7 @@ bool Verifier::init_slots(const std::vector<SessionState *> &slots,
     tail_snap_b_ = (float*) ((uint8_t*) d + a);
     slots_ = slots;
     slot_sp_.assign(slots.size(), sampling_);   // greedy until set_slot_sampling
+    slot_tail_.assign(slots.size(), {});
     std::fprintf(stderr, "strata verify: batch windows of up to %lld sequences (layers [%lld, %lld))\n", (long long) S,
                  (long long) lb_, (long long) le_);
     return true;
@@ -3520,12 +3522,37 @@ catch (sycl::exception const &exc) {
 
 bool Verifier::sample_rows(int S, std::string &err) try {
     bool any = false;
+    constexpr int kSlotHistMax = 4096;   // serve's kPenaltyWindowCap
+    std::vector<int32_t> hrows;
+    std::vector<const int32_t*> tails(slots_.size(), nullptr);
+    std::vector<int64_t> n_tails(slots_.size(), 0);
+    for (size_t b = 0; b < slots_.size() && b < slot_tail_.size(); ++b) { tails[b] = slot_tail_[b].data(); n_tails[b] = (int64_t) slot_tail_[b].size(); }
     for (int t = 0; t < S; ++t) {
         strata::kernels::SamplerParams sp = slot_sp_[(size_t) last_rows_[t]];
-        if (sp.greedy || sp.temperature <= 0.0f) continue;
+        const bool pen = sp.penalty_last_n > 0 && (sp.penalty_repeat != 1.0f || sp.penalty_freq != 0.0f || sp.penalty_present != 0.0f);
+        if ((sp.greedy || sp.temperature <= 0.0f) && !pen) continue;   // a plain greedy slot keeps the window's argmax token untouched
         sp.counter = (uint64_t) last_pos_b_[t];   // Philox(seed, position): the solo window's draw for this position
-        sp.penalty_last_n = 0;
-        strata::kernels::sample_tokens(head_logits_ + (size_t) t * (size_t) n_vocab_, 1, (int) n_vocab_, nullptr, 0, sp,
+        const int32_t* hist = nullptr;
+        int hlen = 0;
+        if (pen) {
+            // the slot's reply tail plus its tokens in this window up to row t (slot_penalty_rows), staged on the device like a solo window's
+            hlen = std::min(sp.penalty_last_n, kSlotHistMax);
+            const size_t need = (size_t) S * (size_t) hlen;
+            hrows.assign(need, -1);
+            strata::kernels::slot_penalty_rows(last_rows_, last_tokens_, S, tails.data(), n_tails.data(), hlen, hrows.data());
+            const OnDevice on_hist(device_);
+            if (slot_hist_cap_ < need) {
+                if (slot_hist_d_ != nullptr) sycl::free(slot_hist_d_, *cs_);
+                slot_hist_d_ = sycl::malloc_device<int32_t>(need, *cs_);
+                if (slot_hist_d_ == nullptr) { slot_hist_cap_ = 0; err = "verify batch: the penalty-history allocation failed"; return false; }
+                slot_hist_cap_ = need;
+            }
+            cs_->memcpy(slot_hist_d_ + (size_t) t * (size_t) hlen, hrows.data() + (size_t) t * (size_t) hlen, (size_t) hlen * sizeof(int32_t)).wait();   // (waited: the host rows are rebuilt for the next row)
+            hist = slot_hist_d_ + (size_t) t * (size_t) hlen;
+        } else {
+            sp.penalty_last_n = 0;
+        }
+        strata::kernels::sample_tokens(head_logits_ + (size_t) t * (size_t) n_vocab_, 1, (int) n_vocab_, hist, hlen, sp,
                                        m_out_ + t, cs_);
         any = true;
     }

sycl/src/kernels/slot_history_test.cpp

@@ -0,0 +1,89 @@
+// Penalty-history rows of a BATCH window, per slot (strata::kernels::slot_penalty_rows).  Header-only, CPU, no model.
+//
+// Row t of a batch window feeds tokens[t] into slot rows[t]; its history is the slot's reply tail followed by that slot's tokens in this window up to
+// and including row t - what penalty_rows gives a solo window of that slot, whatever the other slots do.
+#include "strata/kernels/sampler.hpp"
+
+#include <cstdio>
+#include <algorithm>
+#include <cstdlib>
+#include <random>
+#include <vector>
+
+using strata::kernels::penalty_rows;
+using strata::kernels::slot_penalty_rows;
+
+namespace {
+int checks = 0;
+void check(bool ok, const char* label) {
+    ++checks;
+    if (!ok) { std::fprintf(stderr, "FAIL: %s\n", label); std::exit(1); }
+}
+}  // namespace
+
+int main() {
+    std::mt19937 rng(7);
+    constexpr int kSlots = 6, kH = 16;
+    for (int trial = 0; trial < 400; ++trial) {
+        // per slot: a reply tail of 0..40 tokens
+        std::vector<std::vector<int32_t>> tail(kSlots);
+        std::vector<const int32_t*> tails(kSlots, nullptr);
+        std::vector<int64_t> n_tails(kSlots, 0);
+        for (int b = 0; b < kSlots; ++b) {
+            const int n = (int) (rng() % 41);
+            for (int i = 0; i < n; ++i) tail[b].push_back((int32_t) (rng() % 500));
+            tails[b] = tail[b].empty() ? nullptr : tail[b].data();
+            n_tails[b] = (int64_t) tail[b].size();
+        }
+        // a window: 1..8 rows; each slot at most once, or as a contiguous group of 1..3 rows (MTP drafts)
+        std::vector<int> rows; std::vector<int32_t> toks;
+        std::vector<int> order = {0, 1, 2, 3, 4, 5};
+        std::shuffle(order.begin(), order.end(), rng);
+        for (int b : order) {
+            if (rows.size() >= 8) break;
+            if (rng() % 3 == 0) continue;
+            const int g = 1 + (int) (rng() % 3);
+            for (int j = 0; j < g && rows.size() < 8; ++j) { rows.push_back(b); toks.push_back((int32_t) (rng() % 500)); }
+        }
+        if (rows.empty()) { rows.push_back(2); toks.push_back(9); }
+        const int S = (int) rows.size();
+        std::vector<int32_t> out((size_t) S * kH, 12345);
+        slot_penalty_rows(rows.data(), toks.data(), S, tails.data(), n_tails.data(), kH, out.data());
+        // the definition: per slot, penalty_rows over (its tail, its rows of this window)
+        for (int b = 0; b < kSlots; ++b) {
+            std::vector<int> idx;
+            std::vector<int32_t> win;
+            for (int t = 0; t < S; ++t) if (rows[t] == b) { idx.push_back(t); win.push_back(toks[t]); }
+            if (idx.empty()) continue;
+            std::vector<int32_t> ref(idx.size() * kH);
+            penalty_rows(tail[b].data(), (int64_t) tail[b].size(), win.data(), (int) idx.size(), kH, ref.data());
+            for (size_t k = 0; k < idx.size(); ++k)
+                for (int j = 0; j < kH; ++j)
+                    check(out[(size_t) idx[k] * kH + j] == ref[k * kH + j], "slot row equals penalty_rows of the slot alone");
+        }
+    }
+    // ---- fixed cases: an empty tail, the fed-back head counted, other slots' tokens never in a row
+    {
+        const int rows[3] = {0, 1, 1};
+        const int32_t toks[3] = {10, 20, 21};
+        const int32_t t0[2] = {1, 2};
+        const int32_t t1[3] = {7, 8, 9};
+        const int32_t* tails[2] = {t0, t1};
+        const int64_t n_tails[2] = {2, 3};
+        std::vector<int32_t> out(3 * 4, 0);
+        slot_penalty_rows(rows, toks, 3, tails, n_tails, 4, out.data());
+        const int32_t e[12] = {-1, 1, 2, 10,    7, 8, 9, 20,    8, 9, 20, 21};
+        for (int i = 0; i < 12; ++i) check(out[i] == e[i], "fixed case");
+    }
+    {   // a slot with no tail (no penalty history): its row holds the fed token only
+        const int rows[1] = {0};
+        const int32_t toks[1] = {42};
+        const int32_t* tails[1] = {nullptr};
+        const int64_t n_tails[1] = {0};
+        int32_t out[4] = {0, 0, 0, 0};
+        slot_penalty_rows(rows, toks, 1, tails, n_tails, 4, out);
+        check(out[0] == -1 && out[1] == -1 && out[2] == -1 && out[3] == 42, "empty tail");
+    }
+    std::printf("slot_history_test: %d checks passed\n", checks);
+    return 0;
+}

sycl/src/program/generate.cpp

@@ -9230,6 +9230,8 @@ int main(int argc, char **argv) try {
             // fed).  Kept when the request ends (`cached`), so the next turn of that conversation continues from it
             // instead of reading its history again (see `slot_source` in the request path).
             std::vector<int32_t> ids;
+            int64_t reply0 = 0;            ///< the index in `ids` where its penalty window starts (the last prompt token; 0 with STRATA_PENALTY_PROMPT=1)
+            int pen_n = 0;                 ///< its request's penalty window (penalty_last_n) when it has a repeat / frequency / presence penalty, else 0
             bool cached = false;           ///< idle, and its sessions still hold `ids`
             bool cvec = true;              ///< the control vector setting `ids` were read with
             bool img = false;              ///< the conversation has pictures (never reused from the slot)
@@ -9494,6 +9496,12 @@ int main(int argc, char **argv) try {
                 bt_wait0 = ver.ms_wait; bt_pool0 = ver.ms_pool;
                 bt_miss0 = drive.d.multi_misses; bt_hits0 = drive.d.cache_hits; bt_pcie0 = drive.d.pcie_experts;
             }
+            for (int t = 0; t < S; ++t) {   // each penalised slot's reply tail for its penalty window (the verifier builds the rows from it and this window's tokens)
+                const BSlot& ps = bs[(size_t) rows[t]];
+                if (ps.pen_n <= 0 || (t > 0 && rows[t - 1] == rows[t])) continue;
+                const int64_t from = std::max<int64_t>(ps.reply0, (int64_t) ps.ids.size() - ps.pen_n);
+                ver.set_slot_history(rows[t], ps.ids.data() + from, (int64_t) ps.ids.size() - from);
+            }
             const Clock::time_point w0 = Clock::now();
             if (!ver.run_slot_rows(rows, S, tok, pos, win_pool_fn, win_pool_user, outb, err) || drive.d.failed) {
                 std::printf("ERR %s\n", drive.d.failed && drive.d.fail ? drive.d.fail : err.c_str());
@@ -9675,6 +9683,14 @@ int main(int argc, char **argv) try {
                 }
                 PGroup& G = pg[(size_t) pick];
                 strata::core::progress().busy.store(true);
+                if (k == 0) {   // the penalised slots' reply tails for the group's rows (the last stage draws them)
+                    for (int t = 0; t < GS && t < (int) strata::kernels::kVerifyMaxT; ++t) {
+                        const BSlot& ps = bs[(size_t) (pick * GS + t)];
+                        if (ps.pen_n <= 0) continue;
+                        const int64_t from = std::max<int64_t>(ps.reply0, (int64_t) ps.ids.size() - ps.pen_n);
+                        ver.set_slot_history(pick * GS + t, ps.ids.data() + from, (int64_t) ps.ids.size() - from);
+                    }
+                }
                 if (!stage_verifier(k).batch_launch(pick * GS, GS, G.tok, G.pos, err)) {
                     std::printf("ERR %s\n", err.c_str());
                     return false;
@@ -9995,6 +10011,7 @@ int main(int argc, char **argv) try {
             unsigned long long req_seed = 0;
             float req_min_p = 0.0f, req_penalty_repeat = 1.0f, req_penalty_freq = 0.0f, req_penalty_present = 0.0f;
             int req_penalty_last_n = 0;
+            long long req_pen_from = -1;   // pen_from=N: this request continues one that began at token N (the last token of the original prompt): the penalty window starts there
             int req_cvec = 1;   // cvec=0|1: a loaded control vector for this request (on when absent)
             // SYCL port: logprobs=K (0..20) - after each "T id" an "LP logprob id:logprob ..." line with the token's
             // log-probability and the K most likely tokens', from the verify window's head logits (before sampling,
@@ -10033,6 +10050,7 @@ int main(int argc, char **argv) try {
                     else if (key == "top_k") req_top_k = std::atoi(tok.c_str() + eq + 1);
                     else if (key == "min_p") req_min_p = fv;
                     else if (key == "penalty_last_n") req_penalty_last_n = std::atoi(tok.c_str() + eq + 1);
+                    else if (key == "pen_from") req_pen_from = std::atoll(tok.c_str() + eq + 1);
                     else if (key == "penalty_repeat") req_penalty_repeat = fv;
                     else if (key == "penalty_freq") req_penalty_freq = fv;
                     else if (key == "penalty_present") req_penalty_present = fv;
@@ -11118,7 +11136,8 @@ int main(int argc, char **argv) try {
             for (int64_t i = 0; i < n - 1; ++i) consumed.push_back((int32_t) ids[(size_t) i]);
             // #4590: where the penalty window may start in `consumed` (see pen_prompt above): the end of the prompt without its
             // last token, which joins the window as the fed-back head `x`.
-            const size_t pen_gen0 = strata::kernels::penalty_window_from(consumed.size(), pen_prompt);
+            size_t pen_gen0 = strata::kernels::penalty_window_from(consumed.size(), pen_prompt);
+            if (req_pen_from >= 0 && !pen_prompt) pen_gen0 = std::min<size_t>(pen_gen0, (size_t) req_pen_from);   // a continued request: its reply began earlier
             const char* finish = "length";
             const Clock::time_point d0 = Clock::now();
             // STRATA_DECODE_TIMING=1: where a request's decode time goes (one line per request)
@@ -12383,6 +12402,12 @@ int main(int argc, char **argv) try {
                     sl.max_new = admit_max_new;
                     sl.t0 = Clock::now();
                     sl.ids = live;
+                    {   // the penalty window of the slot's decode: the reply, from the last prompt token (the first window fed it), as the solo loop counts it
+                        const bool slot_pen = req_sp.penalty_last_n > 0 && (req_sp.penalty_repeat != 1.0f || req_sp.penalty_freq != 0.0f || req_sp.penalty_present != 0.0f);
+                        sl.pen_n = slot_pen ? std::min(req_sp.penalty_last_n, kPenaltyWindowCap) : 0;
+                        sl.reply0 = (int64_t) strata::kernels::penalty_window_from(live.size() > 0 ? live.size() - 1 : 0, pen_prompt);
+                        if (req_pen_from >= 0 && !pen_prompt) sl.reply0 = std::min<int64_t>(sl.reply0, req_pen_from);   // a continued request (prompt + what it generated): the reply began earlier
+                    }
                     if (batch_mtp) {
                         if (!slot_mtp[(size_t) admit_slot]->draft_first(1, ver.final_R_all(), sl.x,
                                                                         sl.p - 1, sl.draft.data(), err)) {