Skip to content

[Engine] Metal: moe_gu decodes gate then up rows for 3-4 token groups (MOE_SEQ); Q2_0/IQ2_XXS/IQ3_XXS only #4377

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
joseedson18jc wants to merge 1 commit into
base: main
Choose a base branch
from
Open

[Engine] Metal: moe_gu decodes gate then up rows for 3-4 token groups (MOE_SEQ); Q2_0/IQ2_XXS/IQ3_XXS only #4377

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
59 changes: 56 additions & 3 deletions metal/kernels/moe.metal
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -199,9 +199,54 @@ static inline void gu_body(device const uchar* g0, device const uchar* u0, devic
}
}

// NE >= 3: the gate rows' units, then the up rows' units, so only one of them is decoded at a time (M1 Max: the
// joint body spills at 3-4 entries, 2.3x the 2-entry time for 15% more weight bytes)
template <int F, int LPR, int NR, int NE>
static inline void gu_body_seq(device const uchar* g0, device const uchar* u0, device const float* x,
device const float* xs16, device float* h, device const int* etok, int c, int row0, int li,
constant MoeArgs& a, threadgroup const uint* tab) {
constexpr int UNIT = fmt<F>::UNIT;
float ag[NR][NE], au[NR][NE];
for (int r = 0; r < NR; ++r)
for (int t = 0; t < NE; ++t) ag[r][t] = au[r][t] = 0.0f;
const int U = int(a.n_embd) / UNIT;
for (int u = li; u < U; u += LPR) {
{
float q[NR][UNIT];
float4 cq[NR];
for (int r = 0; r < NR; ++r) Dec<F>::run(g0 + r * a.gu_row, u, q[r], cq[r], tab);
for (int t = 0; t < NE; ++t) {
float xs[UNIT], xsum[2];
X<F>::load(x + etok[t] * int(a.n_embd), xs16 + etok[t] * int(a.n_embd / 16), u, xs, xsum);
for (int r = 0; r < NR; ++r) ag[r][t] += unit_combine<UNIT>(q[r], cq[r], xs, xsum);
}
}
{
float q[NR][UNIT];
float4 cq[NR];
for (int r = 0; r < NR; ++r) Dec<F>::run(u0 + r * a.gu_row, u, q[r], cq[r], tab);
for (int t = 0; t < NE; ++t) {
float xs[UNIT], xsum[2];
X<F>::load(x + etok[t] * int(a.n_embd), xs16 + etok[t] * int(a.n_embd / 16), u, xs, xsum);
for (int r = 0; r < NR; ++r) au[r][t] += unit_combine<UNIT>(q[r], cq[r], xs, xsum);
}
}
}
for (int r = 0; r < NR; ++r)
for (int t = 0; t < NE; ++t) {
const float gv = lane_sum<LPR>(ag[r][t]), uv = lane_sum<LPR>(au[r][t]);
if (li == 0) h[(c + t) * int(a.n_ff) + row0 + r] = gv / (1.0f + exp(-gv)) * uv;
}
}

// bench probe (function constant 22): 1 = no table load (wrong values for the i-quants)
constant int MOE_PROBE [[function_constant(22)]];
constant int moe_probe = is_function_constant_defined(MOE_PROBE) ? MOE_PROBE : 0;
// MOE_SEQ (function constant 25): the 3- and 4-entry groups take gu_body_seq. The host sets it only for windows of 3 or
// more tokens, so the pipelines of 1- and 2-token windows compile without that path (with it in the same pipeline the
// M1 Max ran 1 and 2 tokens 8-10% slower)
constant bool MOE_SEQ [[function_constant(25)]];
constant bool moe_seq = is_function_constant_defined(MOE_SEQ) && MOE_SEQ;

template <int F, int LPR, int NR, int SG>
kernel void moe_gu(device const uchar* W [[buffer(0)]], device const float* x [[buffer(1)]],
Expand Down Expand Up @@ -230,12 +275,20 @@ kernel void moe_gu(device const uchar* W [[buffer(0)]], device const float* x [[
switch (ne) {
case 1: gu_body<F, LPR, NR, 1>(g0, u0, x, xs16, h, etok, c, row0, li, a, tab); break;
case 2: gu_body<F, LPR, NR, 2>(g0, u0, x, xs16, h, etok, c, row0, li, a, tab); break;
case 3: gu_body<F, LPR, NR, 3>(g0, u0, x, xs16, h, etok, c, row0, li, a, tab); break;
case 4: gu_body<F, LPR, NR, 4>(g0, u0, x, xs16, h, etok, c, row0, li, a, tab); break;
case 3:
if (moe_seq) gu_body_seq<F, LPR, NR, 3>(g0, u0, x, xs16, h, etok, c, row0, li, a, tab);
else gu_body<F, LPR, NR, 3>(g0, u0, x, xs16, h, etok, c, row0, li, a, tab);
break;
case 4:
if (moe_seq) gu_body_seq<F, LPR, NR, 4>(g0, u0, x, xs16, h, etok, c, row0, li, a, tab);
else gu_body<F, LPR, NR, 4>(g0, u0, x, xs16, h, etok, c, row0, li, a, tab);
break;
default:
for (int k = 0; k < ne; k += 4) {
const int m = min(4, ne - k);
if (m == 4) gu_body<F, LPR, NR, 4>(g0, u0, x, xs16, h, etok + k, c + k, row0, li, a, tab);
if (m == 4 && moe_seq) gu_body_seq<F, LPR, NR, 4>(g0, u0, x, xs16, h, etok + k, c + k, row0, li, a, tab);
else if (m == 4) gu_body<F, LPR, NR, 4>(g0, u0, x, xs16, h, etok + k, c + k, row0, li, a, tab);
else if (m == 3 && moe_seq) gu_body_seq<F, LPR, NR, 3>(g0, u0, x, xs16, h, etok + k, c + k, row0, li, a, tab);
else if (m == 3) gu_body<F, LPR, NR, 3>(g0, u0, x, xs16, h, etok + k, c + k, row0, li, a, tab);
else if (m == 2) gu_body<F, LPR, NR, 2>(g0, u0, x, xs16, h, etok + k, c + k, row0, li, a, tab);
else gu_body<F, LPR, NR, 1>(g0, u0, x, xs16, h, etok + k, c + k, row0, li, a, tab);
Expand Down
8 changes: 7 additions & 1 deletion metal/src/engine/window.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -974,7 +974,13 @@ void Window::layer(std::vector<mtl::Dispatch>& ds, int l, int ntok, int64_t pos0
mtl::Constants ac = stream_ ? mtl::Constants().b(20, true) : mtl::Constants();
ac.b(23, ct);
const std::string ak = std::string(stream_ ? "#arena" : "") + (ct ? "#ct" : "");
g.p = P(lib_, gun, ac, gun + ak);
// function constant 25 (MOE_SEQ): windows of 3+ tokens decode an expert's gate rows, then its up rows; the 1- and
// 2-token pipelines are built without that path (in the same pipeline it slowed them 8-10% on the M1 Max). Only
// Q2_0, IQ2_XXS and IQ3_XXS gate/up with 1 row a lane: on the M1 Max the K-quants, Q8_0, IQ4_XS, IQ2_XS and the 2-row Q2_0 tile ran 10-43% slower with it
const bool seq = ntok >= 3 && gr == 1 && (tg_.type == ref::Q2_0 || tg_.type == ref::IQ2_XXS || tg_.type == ref::IQ3_XXS);
mtl::Constants gc = ac;
if (seq) gc.b(25, true);
g.p = P(lib_, gun, gc, gun + ak + (seq ? "#seq" : ""));
const bool hg = had_gu(l), hd = had_dn(l);
if (hg) had_rows(ds, lib_, "d:", XF_.b, 0, XR_, 0, ntok, N, hsx_.handle ? &hsx_ : nullptr, false); // before the host's turn
g.args = {mtl::Arg::B(EW, ebase), hg ? mtl::Arg::B(XR_) : mtl::Arg::B(XF_.b, XF_.plane_bytes() * xkind(tg_.type)), mtl::Arg::B(h_),
Expand Down
4 changes: 2 additions & 2 deletions metal/tests/expert_parity.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -414,7 +414,7 @@ int main(int argc, char** argv) {
K.plan = cx.pipeline(lib, "moe_plan");
K.prep = cx.pipeline(lib, "xprep");
K.comb = cx.pipeline(lib, "moe_combine");
K.gu = cx.pipeline(lib, kname("moe_gu", gu, a), mtl::Constants().i(22, std::getenv("MOE_PROBE") ? std::atoi(std::getenv("MOE_PROBE")) : 0).b(23, std::getenv("QUANT_CTAB") != nullptr));
K.gu = cx.pipeline(lib, kname("moe_gu", gu, a), mtl::Constants().i(22, std::getenv("MOE_PROBE") ? std::atoi(std::getenv("MOE_PROBE")) : 0).b(23, std::getenv("QUANT_CTAB") != nullptr).b(25, std::getenv("MOE_SEQ") != nullptr));
K.dn = cx.pipeline(lib, kname("moe_down", dn == ref::Q2_0 && !g_dn_prep ? 142 : dn, b));
K.cgu = a;
K.cdn = b;
Expand Down Expand Up @@ -491,7 +491,7 @@ int main(int argc, char** argv) {
}
std::vector<int> ids;
std::vector<float> w;
random_routing(rng, ntok, NE, ids, w, 0.3);
random_routing(rng, ntok, NE, ids, w, std::getenv("PARITY_SHARE") ? std::atof(std::getenv("PARITY_SHARE")) : 0.3);
mtl::Buffer bid = cx.alloc(4 * ids.size()), bw = cx.alloc(4 * w.size()), plan = cx.alloc(4 * P_INTS);
std::memcpy(bid.ptr, ids.data(), 4 * ids.size());
std::memcpy(bw.ptr, w.data(), 4 * w.size());
Expand Down
Loading