Confirm successful installation by checking the skill directory location:
.cursor/skills/flash-moe-inference
Restart Cursor to activate flash-moe-inference. Access via /flash-moe-inference in your agent's command palette.
β
Security Notice
We perform automated surface-level scans (Gen AI Scanner, Socket, Snyk) during installation. These checks detect common vulnerabilities but do not guarantee complete security. Always review skill source code and verify the publisher's reputation before production use.
Skills execute code in your environment. Always review source, verify the publisher, and test in isolation before production.
Flash-MoE is a pure C/Objective-C/Metal inference engine that runs Qwen3.5-397B-A17B (397B parameter Mixture-of-Experts) on a MacBook Pro with 48GB RAM at 4.4+ tokens/second. It streams 209GB of expert weights from NVMe SSD on demand β no Python, no ML frameworks, just C, Objective-C, and hand-tuned Metal shaders.
Requirements
Hardware: Apple Silicon Mac (M3 Max or similar), 48GB+ unified memory, 1TB+ SSD with ~210GB free
Model: Qwen3.5-397B-A17B safetensors weights (download separately from HuggingFace)
Installation & Build
# Clone the repogit clone https://github.com/danveloper/flash-moe
cd flash-moe/metal_infer
# Build everythingmake# Verify build artifactsls infer chat main
The Makefile compiles infer.m, chat.m, main.m with Metal shader compilation for shaders.metal.
Weight Preparation
Step 1: Extract non-expert weights
# From the metal_infer/ directory# Point to your downloaded Qwen3.5-397B safetensors directorypython3 extract_weights.py /path/to/Qwen3.5-397B-A17B-Instruct/
# Produces:# model_weights.bin (~5.5GB, mmap'd at runtime)# model_weights.json (tensor manifest)# vocab.bin (vocabulary)# tokenizer.bin (BPE tokenizer data)
Step 2: Pack expert weights (4-bit, production)
# From repo rootpython3 repack_experts.py /path/to/Qwen3.5-397B-A17B-Instruct/ metal_infer/packed_experts/
# Produces packed_experts/ directory (~209GB)# Each expert is a separate file: layer_XX_expert_YYYY.bin
Step 3: Optional 2-bit requantization (faster but breaks JSON/tool calling)
The shaders.metal file contains hand-written kernels. Key kernels:
// 4-bit dequantized matrix-vector multiply (FMA-optimized)
// Key insight: fma(nibble, scale*x, bias*x) instead of (nibble*scale + bias)*x
// Pre-compute scale*x and bias*x to fuse dequant+multiply in one FMA instruction
kernel void matvec_4bit_fma(
device const uint8_t* weights [[buffer(0)]],
device const float* scales [[buffer(1)]],
device const float* biases [[buffer(2)]],
device const float* x [[buffer(3)]],
device float* out [[buffer(4)]],
uint tid [[thread_position_in_threadgroup]],
uint gid [[threadgroup_position_in_grid]])
{
// ... tiled SIMD-reduced FMA kernel
// 12% faster than naive (nibble * scale + bias) * x
}
// Fused SwiGLU activation
kernel void swiglu(device float* gate [[buffer(0)]],
device const float* up [[buffer(1)]],
uint gid [[thread_position_in_grid]])
{
float g = gate[gid];
gate[gid] = (g / (1.0f + exp(-g))) * up[gid];
}
// RMS normalization (two-pass)
kernel void rms_norm_pass1(...) // sum of squares reduction
kernel void rms_norm_pass2(...) // apply normalization
// GPU RoPE (fused with Q deinterleave and K normalization)
kernel void rope_qk(...)
// MoE combine + residual + sigmoid gate (fused)
kernel void moe_combine_residual(...)
SSD Expert Streaming Pattern
The core innovation β loading only K=4 active experts per layer from SSD:
// Parallel expert loading using GCD dispatch groups// From infer.m (conceptual pattern)dispatch_group_t group =dispatch_group_create();dispatch_queue_t ioQueue =dispatch_get_global_queue(QOS_CLASS_USER_INITIATED,0);for(int k =0; k < K_EXPERTS; k++){int expert_id = top_k_indices[k];dispatch_group_async(group, ioQueue,^{// Each expert: ~6.75MB at 4-bitchar path[256];snprintf(path,sizeof(path),"packed_experts/layer_%02d_expert_%04d.bin", layer, expert_id);int fd =open(path, O_RDONLY);// pread() β non-blocking, OS page cache handles LRUpread(fd, expert_buffer[k], expert_size,0);close(fd);});}dispatch_group_wait(group, DISPATCH_TIME_FOREVER);// GPU compute follows β serial pipeline is hardware-optimal on Apple Silicon
Why pread() not mmap(): mmap incurs per-page fault overhead on cold data (~5x slower). Direct pread() with OS page cache achieves ~71% hit rate naturally.
GatedDeltaNet Linear Attention (BLAS)
The recurrence update uses Accelerate BLAS β 64% faster than scalar:
// GatedDeltaNet state update per head (conceptual pattern)// state: 128Γ128 float matrix, 64 heads// From infer.m#import<Accelerate/Accelerate.h>for(int h =0; h <64; h++){float* S = state + h *128*128;// 128Γ128 state matrixfloat* q = Q + h *128;float* k = K + h *128;float* v = V + h *128;// Ξ²Β·(kβv) outer product update// cblas_sger: S += beta * (k β v)cblas_sger(CblasRowMajor,128,128, beta[h], k,1, v,1, S,128);// Decay: S = alpha * Scblas_sscal(128*128, alpha[h], S,1);// Output: o = S @ qcblas_sgemv(CblasRowMajor, CblasNoTrans,128,128,1.0f, S,128, q,1,0.0f, output + h *128,1);}
Performance Configuration
4-bit (production default)
Quality: Excellent β full tool calling, correct JSON
Speed: 4.36 tok/s
Disk: 209GB
2-bit (speed testing only)
Quality: Good β but breaks JSON/tool calling (\name\ instead of "name")
Speed: 5.74 tok/s (7.05 peak single token with warm cache)
Disk: 120GB
Uses F_NOCACHE flag to avoid page cache thrashing
What NOT to Try (Learned from 58 Experiments)
Approach
Why it fails
mmap() expert f
β
Make data-driven prioritization decisions faster
Stakeholder Communication
Draft PRDs, status updates, and stakeholder presentations
βΊAccess to product documentation and roadmap tools (Jira, Notion, etc.)
βΊUnderstanding of product management frameworks (RICE, Jobs-to-be-Done, etc.)
βΊStakeholder contact information and communication channels
Time Estimate
30-60 minutes to see productivity improvements
Steps
1Install product management skill
2Start with user story generation for known feature
3Progress to competitive analysis: research 2-3 competitors
4Use for roadmap prioritization: apply RICE/ICE scoring
5Draft stakeholder communications and refine based on feedback
6Build template library for recurring PM tasks
7Share effective prompts with product team
Common Pitfalls
β Not validating competitive researchβverify facts before sharing
β Accepting user stories without involving engineering team
β Over-relying on frameworks without qualitative judgment
β Not customizing outputs to company culture and communication style
β Skipping stakeholder validation of generated requirements
Best Practices
β Do
+Validate research and competitive analysis with real data
+Collaborate with engineering when generating technical requirements
+Customize frameworks and templates to your company context
+Use skill for first drafts, refine with stakeholder input
+Document successful prompt patterns for PM tasks
+Combine AI efficiency with human judgment and intuition
β Don't
βDon't publish competitive analysis without fact-checking
βDon't finalize user stories without engineering review
βDon't make prioritization decisions solely on AI scoring
βDon't skip customer validation of generated requirements
βDon't ignore company-specific context and culture
π‘ Pro Tips
β Provide context: company goals, constraints, customer feedback
β Ask for alternatives: 'Show 3 ways to prioritize this roadmap'
β Request stakeholder-specific formatting: 'Executive summary vs. engineering spec'
β Use skill for 70% generation + 30% customization to company needs
When to Use This
β Use when
Use for user story writing, competitive research, roadmap prioritization, stakeholder communication, and PRD drafting. Best for reducing repetitive documentation and research work.
β Avoid when
Avoid for strategic product vision (requires deep customer empathy), pricing decisions (needs market and financial expertise), or when face-to-face customer discovery is more valuable than speed.
Learning Path
1Basic: user stories, feature specs, status updates
