Niko1221/Strata
1
0 Fork 0
mirror of https://github.com/Niko1221/Strata.git synced 2026-11-07 14:21:06 -10:00
Code Issues Packages Projects Releases 1 Wiki Activity

SYCL: the #4185 decode controls default on for Battlemage (STRATA_<name>=0 turns each off), STRATA_SYCL_IQ_DECODE_FAST default ON in CMake and passed by the SYCL build scripts; STRATA_SELECT_GEMM_MIN sets the block threshold of the #3248 GEMM block scores

Browse Source 
Controls of glorysdj's pull request #4185: unset means on for a Battlemage device and off for every other device (strata::bm_default_on); a set value still decides. Measured results are in phaseA-tests/pr-priorities/RESULTS-INTEL.md.
This commit is contained in:
Niko1221 committed 2026-11-01 00:50:05 +01:00
1 parent 5a41f3b391
commit 8fffe81d4b
10 files changed
+34 -45
Download Patch File Download Diff File Expand all files Collapse all files

No files matched your search

packaging/build_sycl_linux.sh
+1 -1
View File
@@ -42,7 +42,7 @@ if [ -n "${LLAMA_CPP:-}" ]; then ggml="-DFETCHCONTENT_SOURCE_DIR_STRATA_LLAMACPP
elif [ -d "$repo/third_party/llama.cpp/ggml" ]; then ggml="-DSTRATA_GGML_DIR=$repo/third_party/llama.cpp"; fi
# shellcheck disable=SC2086
cmake -S sycl -B "$build" -G Ninja -DCMAKE_BUILD_TYPE=Release -DCMAKE_C_COMPILER=icx -DCMAKE_CXX_COMPILER=icpx \
-DSTRATA_SYCL_PARITY=OFF ${aot:+"-DSTRATA_SYCL_AOT=$aot"} $ggml
-DSTRATA_SYCL_PARITY=OFF -DSTRATA_SYCL_IQ_DECODE_FAST=${IQ_DECODE_FAST:-ON} ${aot:+"-DSTRATA_SYCL_AOT=$aot"} $ggml
cmake --build "$build" --target strata -j "$jobs"
ver=$(sed -n 's/^[[:space:]]*project([^)]*VERSION[[:space:]]*\([0-9][0-9.]*\).*/\1/p' CMakeLists.txt | head -1)
sycl/CMakeLists.txt
+4 -1
View File
@@ -29,7 +29,10 @@ endif()
add_compile_definitions(STRATA_VERSION="${PROJECT_VERSION}")
set(STRATA_SYCL_AOT "" CACHE STRING "AOT device for the GPU code, e.g. bmg-g31 (empty: SPIR-V, JIT at run time)")
option(STRATA_SYCL_IQ_DECODE_FAST "Opt in to the packed IQ/Q2 decode helpers" OFF)
# PR #4185 (glorysdj): the packed Q2 / IQ4 lookup / IQ3_S sign arithmetic of the mixed-IQ3 decode kernels. Integer arithmetic with the
# same results as the code it replaces (bit-exact, checked end to end on 2x Arc Pro B65), faster on Battlemage; ON by default in every
# build, including the JIT release package. -DSTRATA_SYCL_IQ_DECODE_FAST=OFF builds the earlier helpers.
option(STRATA_SYCL_IQ_DECODE_FAST "Packed IQ/Q2 decode helpers (bit-exact integer arithmetic, faster on Battlemage)" ON)
if(STRATA_SYCL_IQ_DECODE_FAST)
add_compile_definitions(STRATA_Q2_WORD_FAST=1 STRATA_IQ4NL_LUT_SWAR=1 STRATA_IQ4XS_LUT_SWAR=1
STRATA_IQ4_LUT_SWAR=1 STRATA_IQ3S_SIGN_FAST=1 STRATA_IQ3S_SIGN_APPLY_FAST=1)
sycl/include/strata/sycl_queue.hpp
+7
View File
@@ -156,6 +156,13 @@ inline IntelGpuGen intel_gpu_gen(const sycl::device& d) {
}
return IntelGpuGen::Unknown;
}
// The Battlemage decode controls of PR #4185 (glorysdj): unset = on for a Battlemage device, off for every other device (Arc
// A-series included); a set value turns the control on ("1") or off ("0") whatever the device, as before.
inline bool bm_default_on(const char* name, const sycl::device& d) {
const char* v = std::getenv(name);
if (v != nullptr && v[0] != '') return v[0] != '0';
return intel_gpu_gen(d) == IntelGpuGen::Battlemage;
}
} // namespace strata
#include <cstdint>
sycl/src/kernels/cuda/fused_gr.dp.cpp
+3 -12
View File
@@ -1978,14 +1978,8 @@ bool fused_gr_read_multi(const FusedGrArgs* a, int n_tok, float* xn_scratch, voi
return false; // SYCL port: never writes the q8_1 images (STRATA_QFUSE is not supported here)
}
m.part = down_partials(st);
static const bool reuse_norm = [] {
const char* v = std::getenv("STRATA_GR_CACHED_NORM");
return v && v[0] == '1';
}();
static const bool reuse_lo = [] {
const char* v = std::getenv("STRATA_GR_LO_REUSE");
return v && v[0] == '1';
}();
static const bool reuse_norm = strata::bm_default_on("STRATA_GR_CACHED_NORM", st->get_device());
static const bool reuse_lo = strata::bm_default_on("STRATA_GR_LO_REUSE", st->get_device());
if ((reuse_norm || reuse_lo) && gr_norm_split() && gr_down_sliced() &&
!gr_more_groups(st) && !gr_large_grf(st) && !gr_up_ilp() && gr_up_cols() == UPM_COLS &&
strata::intel_gpu_gen(st->get_device()) == strata::IntelGpuGen::Battlemage &&
@@ -2062,10 +2056,7 @@ bool fused_gr_read_multi(const FusedGrArgs* a, int n_tok, float* xn_scratch, voi
const sycl::nd_range<3> nr(sycl::range(1, 1, nred) * sycl::range(1, 1, 128),
sycl::range(1, 1, 128));
if (m.cached_lo) {
static const bool reduce_ilp = [] {
const char* v = std::getenv("STRATA_GR_REDUCE_ILP");
return v && v[0] == '1';
}();
static const bool reduce_ilp = strata::bm_default_on("STRATA_GR_REDUCE_ILP", st->get_device());
if (reduce_ilp) {
strata::kprof::kp(st, __FILE__, __func__, __LINE__)->template parallel_for<dpct_kernel_name<class gr_down_reduce_lo_ilp>>(nr,
[=](sycl::nd_item<3>) [[sycl::reqd_sub_group_size(32)]] {
sycl/src/kernels/cuda/iq_kernels.dp.cpp
+8 -7
View File
@@ -3441,8 +3441,7 @@ bool g_old_kernels = env_on("STRATA_OLD_IQ_MMVQ");
// every format; STRATA_EXPERT_SPLIT=1 takes upstream's multi kernels, launched with their own grid (GU_ROWS rows).
bool g_split_multi = env_on("STRATA_EXPERT_SPLIT");
bool g_no_sub16_gu = env_on("STRATA_NO_SUB16_GU");
bool g_iq3s_gu_slm = env_on("STRATA_IQ3S_GU_SLM");
bool g_iq4nl_down_slm = env_on("STRATA_IQ4NL_DOWN_SLM");
// PR #4185 controls: on by default for a Battlemage device (strata::bm_default_on), =0 turns each off
// STRATA_IQ_STAGE_GRID=0 disables staging 64-bit i-quant codebook tables into shared memory
bool g_stage_grid = [] {
const char* v = std::getenv("STRATA_IQ_STAGE_GRID");
@@ -4829,7 +4828,7 @@ void launch_gu_row_reuse(unsigned groups, dpct::queue_ptr s,
const unsigned gx = (unsigned) ((2 * L.n_ff + 63) / 64);
auto props = sycl::ext::oneapi::experimental::properties{
sycl::ext::oneapi::experimental::use_root_sync};
static const bool e3 = TG == 21 || env_on("STRATA_EXPERT_ROW_E3");
static const bool e3 = TG == 21 || strata::bm_default_on("STRATA_EXPERT_ROW_E3", s->get_device());
if (e3) {
static const bool logged = [] {
std::fprintf(stderr, "strata: %s exact three-entry row reuse enabled\n",
@@ -4873,7 +4872,7 @@ void launch_gu_port(unsigned groups, dpct::queue_ptr s, const unsigned long long
const NativeExpertLayout &L, float *gate, float *up) {
#if !STRATA_GRID_SLM
if constexpr ((TG == 22 || TG == 18) && LN == 8) {
static const bool reuse = env_on(TG == 22 ? "STRATA_IQ2S_GU_ROW_REUSE" : "STRATA_IQ3XXS_GU_ROW_REUSE");
static const bool reuse = strata::bm_default_on(TG == 22 ? "STRATA_IQ2S_GU_ROW_REUSE" : "STRATA_IQ3XXS_GU_ROW_REUSE", s->get_device());
if (reuse && !g_old_kernels && !g_split_multi &&
strata::intel_gpu_gen(s->get_device()) == strata::IntelGpuGen::Battlemage) {
static const bool note = [] {
@@ -4986,7 +4985,7 @@ void launch_gu(dpct::dim3 grid, dpct::queue_ptr s,
float *up) {
if constexpr (TG == 21) {
#if !STRATA_GRID_SLM
static const bool reuse = env_on("STRATA_IQ3S_GU_ROW_REUSE");
static const bool reuse = strata::bm_default_on("STRATA_IQ3S_GU_ROW_REUSE", s->get_device());
if (reuse && !g_old_kernels && !g_split_multi && gu_lanes() == 8 &&
L.n_embd == 2560 && L.n_ff == 640 &&
strata::intel_gpu_gen(s->get_device()) == strata::IntelGpuGen::Battlemage) {
@@ -4999,7 +4998,8 @@ void launch_gu(dpct::dim3 grid, dpct::queue_ptr s,
return;
}
#endif
if (!g_old_kernels && g_iq3s_gu_slm && L.n_embd == 2560 && L.n_ff == 640) {
static const bool iq3s_gu_slm = strata::bm_default_on("STRATA_IQ3S_GU_SLM", s->get_device());
if (!g_old_kernels && iq3s_gu_slm && L.n_embd == 2560 && L.n_ff == 640) {
auto exp_props =
sycl::ext::oneapi::experimental::properties{sycl::ext::oneapi::experimental::use_root_sync};
const size_t gx = (size_t) ((2 * L.n_ff + 31) / 32);
@@ -5058,7 +5058,8 @@ void launch_down(dpct::dim3 grid, dpct::queue_ptr s,
const block_q8_1 *hq, const NativeExpertLayout &L,
float *out) {
if constexpr (TD == 20) {
if (!g_old_kernels && g_iq4nl_down_slm && L.n_embd == 2560 && L.n_ff == 640) {
static const bool iq4nl_down_slm = strata::bm_default_on("STRATA_IQ4NL_DOWN_SLM", s->get_device());
if (!g_old_kernels && iq4nl_down_slm && L.n_embd == 2560 && L.n_ff == 640) {
launch_down_iq4nl_slm(grid, s, grp_ptr, grp_start, n_groups, ent_dst, hq, L, out);
return;
}
sycl/src/kernels/cuda/native_mmvq.dp.cpp
+2 -8
View File
@@ -1754,10 +1754,7 @@ void launch_q6k_wide(const void* weights, const void* x_q8_1, float* y, int n_in
if (rpw == 1 && q6k_a2()) {
const unsigned blocks = unsigned((std::size_t(n_out) + WARPS - 1) / WARPS);
if constexpr (NCOLS > 1) {
static const bool shape_grf = [] {
const char* v = std::getenv("STRATA_Q6_SHAPE_GRF");
return v && v[0] == '1';
}();
static const bool shape_grf = strata::bm_default_on("STRATA_Q6_SHAPE_GRF", s->get_device());
const bool eligible = (n_in == 6144 && n_out == 2560) ||
(n_in == 2560 && (n_out == 512 || n_out == 640)) ||
(NCOLS >= 3 && n_in == 2560 && n_out == 6144) ||
@@ -2013,10 +2010,7 @@ void launch_wide(const void* weights, const void* x_q8_1, float* y, int n_in, in
const auto* w = static_cast<const typename F::Block*>(weights);
const auto* x = static_cast<const Q81Block*>(x_q8_1);
int rpw = wide_k_rpw();
static const bool shape_reuse = [] {
const char* v = std::getenv("STRATA_MMVQ_SHAPE_RPW");
return v && v[0] == '1';
}();
static const bool shape_reuse = strata::bm_default_on("STRATA_MMVQ_SHAPE_RPW", s->get_device());
if (shape_reuse && std::getenv("STRATA_MMVQ_WIDE_K_RPW") == nullptr &&
strata::intel_gpu_gen(s->get_device()) == strata::IntelGpuGen::Battlemage) {
bool eligible = false;
sycl/src/kernels/cuda/qsa_decode_attn.dp.cpp
+3 -6
View File
@@ -677,11 +677,8 @@ __dpct_inline__ void attn_merge_active_kernel(
attn[(size_t) h * HD + d] = L > 0.0f ? acc / L : 0.0f;
}
bool merge_active_on() {
static const bool on = [] {
const char* v = std::getenv("STRATA_ATTN_MERGE_ACTIVE");
return v && v[0] == '1';
}();
bool merge_active_on(const sycl::device& d) {
static const bool on = strata::bm_default_on("STRATA_ATTN_MERGE_ACTIVE", d);
return on;
}
@@ -1076,7 +1073,7 @@ void qsa_decode_attn_batch(const float* q, const QsaAttnPools& pools, const int3
auto exp_props = sycl::ext::oneapi::experimental::properties{
sycl::ext::oneapi::experimental::use_root_sync};
if (merge_active_on() && n_q <= 8) {
if (merge_active_on(st->get_device()) && n_q <= 8) {
strata::kprof::kp(st, __FILE__, __func__, __LINE__)->template parallel_for<dpct_kernel_name<class attn_merge_active>>(
sycl::nd_range<3>(
sycl::range(1, (unsigned)n_q, (unsigned)s.n_head) *
sycl/src/kernels/cuda/verify_kernels.dp.cpp
+2 -8
View File
@@ -1892,10 +1892,7 @@ void gdn_step_norm_multi(float* state, const float* h, int conv_channels, const
std::fprintf(stderr, "gdn_step_norm_multi: invalid arguments\n");
std::exit(1);
}
static const bool qk_reg = [] {
const char* v = std::getenv("STRATA_GDN_QK_REG");
return v && v[0] == '1';
}();
static const bool qk_reg = strata::bm_default_on("STRATA_GDN_QK_REG", strata::q_of(stream)->get_device());
if (qk_reg && !xq && !n_keep && t_out_begin == 0 &&
h_k == 16 && h_v == 48 && conv_channels == 10240 &&
strata::intel_gpu_gen(strata::q_of(stream)->get_device()) == strata::IntelGpuGen::Battlemage) {
@@ -2467,10 +2464,7 @@ void resident_plan(const int32_t* ids, int n_entries, int k, const int32_t* res_
int32_t* use = nullptr; // the layer's slice of the device routing counts, if on
if (g_mirror_table != nullptr && in_mirror_table(res_layer)) mir = g_mirror_table + (res_layer - g_mirror_res);
if (g_usage != nullptr && in_mirror_table(res_layer)) use = g_usage + (res_layer - g_mirror_res);
static const bool compact = [] {
const char* e = std::getenv("STRATA_RESIDENT_PLAN_COMPACT");
return e != nullptr && std::atoi(e) != 0;
}();
static const bool compact = strata::bm_default_on("STRATA_RESIDENT_PLAN_COMPACT", strata::q_of(stream)->get_device());
const int threads = compact ? std::max(32, ((n_entries + 31) / 32) * 32) : kResidentPlanMax;
{
auto exp_props = sycl::ext::oneapi::experimental::properties{
sycl/src/prefill/prefill.cpp
+2 -1
View File
@@ -3323,7 +3323,8 @@ bool Prefill::run_impl(const int64_t *tokens, int64_t n, int64_t pos0,
// 3x faster at 40K), so by default only for a batch reaching 2,048 blocks.
// STRATA_SELECT_GEMM=0: the warp kernel always, =1: the GEMM always.
static const int sel_gemm_mode = [] { const char* v = std::getenv("STRATA_SELECT_GEMM"); return v ? std::atoi(v) : -1; }();
const bool sel_gemm = sel_gemm_mode > 0 || (sel_gemm_mode < 0 && active >= 2048);
static const int64_t sel_gemm_min = [] { const char* v = std::getenv("STRATA_SELECT_GEMM_MIN"); return v ? (int64_t) std::atoll(v) : (int64_t) 2048; }(); // blocks (4 cells each)
const bool sel_gemm = sel_gemm_mode > 0 || (sel_gemm_mode < 0 && active >= sel_gemm_min);
const int64_t sel_tile = std::min<int64_t>(8192, (int64_t) ((uint64_t) m.attn_batch *
strata::kernels::qsa_decode_attn_scratch_floats(m.cap, s) / (uint64_t) (nb * 4)));
bool sel_done = false;
sycl/tools/build.sh
+2 -1
View File
@@ -1,6 +1,7 @@
#!/usr/bin/env bash
# Configure + build the SYCL port inside strata-sycl-dev. sycl/build.sh [target...]
# AOT=bmg-g31 ahead-of-time device code (empty: SPIR-V, JIT at the first start)
# IQ_DECODE_FAST=OFF the earlier IQ/Q2 decode helpers (default ON: bit-exact packed arithmetic, PR #4185)
# I8_ESIMD=OFF no fused int8 ESIMD expert kernels (stubs; the prompt's experts take the XMX path). AUTO (default) or ON otherwise.
# ALLOW_BAD_OCLOC=1 build although this ocloc is known to crash on the Battlemage kernels (see below)
set -uo pipefail
@@ -37,7 +38,7 @@ fi
repo=${REPO:-$(git rev-parse --show-toplevel)}
b=${BUILD_DIR:-$repo/build-sycl}
[ -f $b/build.ninja ] || cmake -S $repo/sycl -B $b -G Ninja -DCMAKE_C_COMPILER=icx -DCMAKE_CXX_COMPILER=icpx \
-DSTRATA_SYCL_AOT="${AOT:-}" -DSTRATA_SYCL_SPIN_MAX="${SPIN_MAX:-}" -DSTRATA_SYCL_I8_ESIMD="${I8_ESIMD:-AUTO}" 2>&1 | tail -15
-DSTRATA_SYCL_AOT="${AOT:-}" -DSTRATA_SYCL_SPIN_MAX="${SPIN_MAX:-}" -DSTRATA_SYCL_I8_ESIMD="${I8_ESIMD:-AUTO}" -DSTRATA_SYCL_IQ_DECODE_FAST="${IQ_DECODE_FAST:-ON}" 2>&1 | tail -15
# an existing build folder keeps its cache: a switch given now reaches it (only this one, so AOT and the rest stay as configured)
[ -z "${I8_ESIMD:-}" ] || cmake -S $repo/sycl -B $b -DSTRATA_SYCL_I8_ESIMD="$I8_ESIMD" >/dev/null 2>&1
cmake --build $b -j ${JOBS:-12} ${1:+--target "$@"} -- -k 0 2>&1 | tee $b/build.log | grep -E '^FAILED|error:|^ninja: build stopped|Linking|^\[[0-9]+/[0-9]+\] Linking' | tail -40