前人工作区变更归档（元审计遗留处置·冰朔授权'交给你来处理'·不否定过去原则）：hololake-phase-b-staging删除组/WORK区/submodule指针/0826壳源/neural-chain/skills-mirror/codex恢复帧/.zcode本地配置等140项原样冻结入库——不删除不修改不评判，作为历史一环保留；WX-SECRETS回执经查仅含公开AppID无AppSecret，安全

ZCODE-DEV-20260827/zcode-zc001-flagos-task29-20260827/v10/gelu_and_mul.py

@@ -0,0 +1,29 @@
+# NO_CHIP_SPECIFIC
+"""FlagOS S2 · Task 29 gelu_and_mul — 通用版（国际通用芯片：原生 Triton，tl.math.erf）"""
+import torch
+import triton
+import triton.language as tl
+
+
+@triton.jit
+def _gelu_and_mul_kernel(x_ptr, out_ptr, d, BLOCK: tl.constexpr):
+    row = tl.program_id(0).to(tl.int64)
+    x_row = x_ptr + row * (2 * d)
+    out_row = out_ptr + row * d
+    offs = tl.arange(0, BLOCK)
+    for start in range(0, d, BLOCK):
+        idx = start + offs
+        mask = idx < d
+        x1 = tl.load(x_row + idx, mask=mask, other=0.0).to(tl.float32)
+        x2 = tl.load(x_row + d + idx, mask=mask, other=0.0).to(tl.float32)
+        g = 0.5 * x1 * (1.0 + tl.math.erf(x1 * 0.7071067811865476))
+        tl.store(out_row + idx, (g * x2).to(out_ptr.dtype.element_ty), mask=mask)
+
+
+def gelu_and_mul(x: torch.Tensor, out: torch.Tensor = None) -> torch.Tensor:
+    num_tokens, two_d = x.shape
+    d = two_d // 2
+    if out is None:
+        out = torch.empty((num_tokens, d), dtype=x.dtype, device=x.device)
+    _gelu_and_mul_kernel[(num_tokens,)](x, out, d, BLOCK=1024, num_warps=4)
+    return out

ZCODE-DEV-20260827/zcode-zc001-flagos-task29-20260827/v10/gelu_and_mul_chip.py

@@ -0,0 +1,34 @@
+"""FlagOS S2 · Task 29 gelu_and_mul — {CHIP_CN}版（FlagTree/Triton-TLE 环境）
+
+写法严格照官方《Triton和Triton-TLE入门指南》示例：
+  import triton / triton.language 照常（FlagTree 提供兼容层）；
+  数学原语用 tle.erf（跨芯片统一，TLE-Lite 层）。
+"""
+import torch
+import triton
+import triton.language as tl
+import flagtree.tle as tle
+
+
+@triton.jit
+def _gelu_and_mul_kernel(x_ptr, out_ptr, d, BLOCK: tl.constexpr):
+    row = tl.program_id(0).to(tl.int64)
+    x_row = x_ptr + row * (2 * d)
+    out_row = out_ptr + row * d
+    offs = tl.arange(0, BLOCK)
+    for start in range(0, d, BLOCK):
+        idx = start + offs
+        mask = idx < d
+        x1 = tl.load(x_row + idx, mask=mask, other=0.0).to(tl.float32)
+        x2 = tl.load(x_row + d + idx, mask=mask, other=0.0).to(tl.float32)
+        g = 0.5 * x1 * (1.0 + tle.erf(x1 * 0.7071067811865476))
+        tl.store(out_row + idx, (g * x2).to(out_ptr.dtype.element_ty), mask=mask)
+
+
+def gelu_and_mul(x: torch.Tensor, out: torch.Tensor = None) -> torch.Tensor:
+    num_tokens, two_d = x.shape
+    d = two_d // 2
+    if out is None:
+        out = torch.empty((num_tokens, d), dtype=x.dtype, device=x.device)
+    _gelu_and_mul_kernel[(num_tokens,)](x, out, d, BLOCK=1024, num_warps=4)
+    return out