cutileGPT

Pure Tile Programming Philosophy: Think in WHAT, not HOW

A complete GPT implementation proving declarative GPU programming works. Using NVIDIA's CUDA Tile framework, cutileGPT achieves 8.3x speedup on GELU and matches PyTorch performance (within 4%) - all with ~10MB footprint vs PyTorch's ~2GB.

🎨 Tile Programming Philosophy

The Paradigm Shift

# ❌ Traditional CUDA (Imperative HOW)
@cuda.jit
def kernel(x, y, N):
    tid = cuda.threadIdx.x + cuda.blockIdx.x * cuda.blockDim.x
    __shared__ smem[256]
    smem[threadIdx.x] = x[tid]
    __syncthreads()
    # ... manual reduction loops ...

# ✅ Tile Programming (Declarative WHAT)
@ct.kernel
def kernel(X, Y, N):
    x_tile = ct.load(X, ...)      # "Load this data"
    mean = ct.sum(x_tile) / N     # "Compute mean"
    ct.store(Y, ...)              # "Store result"
    # Compiler handles threads, sync, and optimization!

Core Principle: Specify WHAT you want (operations), let the compiler handle HOW (threads, sync, memory).

🚀 Key Results

Performance

Benefits: The Dramatic Simplification

87% less code: Traditional CUDA kernels require ~150 lines with manual thread management, explicit synchronization, and GPU-specific optimizations. Tile Programming reduces this to ~20 lines of clean, declarative code where the compiler handles everything.

Simpler architecture: Complex interconnected components (thread management, block config, sync logic, shared memory) collapse into a single declarative interface. The compiler automatically optimizes for your specific GPU.

📊 Performance Visualizations

Real benchmark results from our GPU (NVIDIA GB10):

GELU Kernel Speedup

8x faster than CuPy on a large tensor (32×512×768 = 12M elements). Tile Programming's declarative approach enables aggressive compiler optimizations.

cutileGPT Performance

Latency and throughput across different model sizes. Larger models benefit more from Tile Programming's efficient kernel fusion.

PyTorch Comparison: Comprehensive Analysis

We benchmarked across 36 configurations (3 model sizes × 4 batch sizes × 3 sequence lengths) to understand performance characteristics across multiple dimensions.

Key Findings:

• Small workloads (batch=1, seq=64): PyTorch faster due to lower kernel launch overhead
• Medium workloads (batch=4-8): Performance gap narrows as computation dominates
• Large workloads (batch=16, seq=256): Near parity with PyTorch (0.977x on medium model)
• Best case: Nano model at batch=8, seq=256 achieves 1.011x (faster than PyTorch!)

Heatmaps show latency and performance ratio across all configurations. Warmer colors (green) indicate better cutileGPT performance, especially visible in large batch scenarios.

Throughput trends: cutileGPT throughput scales well with sequence length, closing the gap with PyTorch as workload size increases. This validates the Tile Programming approach for production workloads.

Trade-off Analysis:

• When to use PyTorch: Small batch inference (batch ≤ 4), latency-critical applications
• When to use cutileGPT: Large batch processing, edge deployment (~10MB vs ~2GB), hardware portability

Footprint Comparison:

• PyTorch minGPT: ~2GB (torch + dependencies)
• cutileGPT: ~10MB (cupy + cuda-tile)
• 200x smaller for edge deployment and serverless

Tile Programming Philosophy

The fundamental shift: specify WHAT (operations), let compiler handle HOW (threads, sync, memory).

⚡ Quick Start

Option 1: Install from PyPI

pip install cutile-gpt[hf]

from cutile_gpt import CutileGPT, GPTConfig

# Load GPT-2 from HuggingFace
model = CutileGPT(GPTConfig.gpt2())
model.load_from_huggingface('gpt2')

# Generate text
import cupy as cp
tokens = cp.array([[15496, 11, 616, 1438, 318]], dtype=cp.int32)  # "Hello, my name is"
generated = model.generate(tokens, max_new_tokens=20)

Option 2: Clone and Run Demo

# Clone and install
git clone --recursive https://github.com/falcons-eyes/cutileGPT.git
cd cutileGPT
uv sync

# Run complete demo
uv run python demo_tile_gpt.py

Output:

✅ Part 1: Individual Tile kernels (LayerNorm, GELU, Linear, Attention)
✅ Part 2: Transformer block test
✅ Part 3: Complete GPT model (forward + generation)
✅ Part 4: Philosophy comparison (Traditional vs Tile)
✅ Part 5: Performance benchmark (8.3x speedup!)

SUCCESS: All Tests Passed!

Use in Your Code

import cupy as cp
from cutile_gpt import CutileGPT, GPTConfig

# Create model with preset config
config = GPTConfig.gpt_nano()
model = CutileGPT(config)

# Or load from HuggingFace
model = CutileGPT(GPTConfig.gpt2())
model.load_from_huggingface('gpt2')

# Forward pass
tokens = cp.array([[100, 200, 300]], dtype=cp.int32)
logits, _ = model.forward(tokens)  # logits: (1, 3, vocab_size)

# Generate text
generated = model.generate(tokens, max_new_tokens=50)

🔧 Installation

Prerequisites

• Python 3.13+
• CUDA Toolkit 13.1+ - required by tileiras, the Tile IR compiler
• NVIDIA Driver r580+
• NVIDIA Blackwell GPU - sm_100 (B200/GB200) or sm_120 (GB10, RTX 50 series)

tileiras currently compiles for Blackwell only, so Hopper (sm_90) and earlier are not supported yet. Upstream lists this as a temporary restriction - see cuTile Python system requirements.

Install from PyPI (Recommended)

# Core package only (minimal dependencies)
pip install cutile-gpt

# With HuggingFace support (transformers, datasets, tiktoken)
pip install cutile-gpt[hf]

# With PyTorch for benchmarking
pip install cutile-gpt[torch]

# With visualization tools (plotly, matplotlib, pandas)
pip install cutile-gpt[viz]

# Everything included
pip install cutile-gpt[all]

Install from Source (Development)

# Clone with submodules
git clone --recursive https://github.com/falcons-eyes/cutileGPT.git
cd cutileGPT

# Or if already cloned
git submodule update --init --recursive

# Install with uv (recommended)
uv sync

# Or with pip
pip install -e ".[all]"

Dependency Structure

💻 Usage

Note: Core features (kernels, Tile API) work with pip install cutile-gpt. HuggingFace loading requires pip install cutile-gpt[hf].

Individual Kernels (Core)

import cupy as cp
from cutile_gpt import cutile_layer_norm, cutile_gelu, cutile_linear_bias

# LayerNorm - Declarative, no manual sync
x = cp.random.randn(4, 128, 768, dtype=cp.float32)
weight = cp.ones(768, dtype=cp.float32)
bias = cp.zeros(768, dtype=cp.float32)
y = cutile_layer_norm(x, weight, bias)

# GELU - 8.3x faster than CuPy!
y = cutile_gelu(x)

# Linear - Tile-based matmul with Tensor Cores
y = cutile_linear_bias(x, weight, bias)

Tile API (Fluent Builder)

from cutile_gpt import tile, configure_tiles, TileConfig

# Fluent API for declarative operations
result = (
    tile(x, "input")
    .linear(weight, bias, out_features=768)
    .gelu()
    .execute()
)

# Configure tile sizes for optimization
configure_tiles(TileConfig(tile_m=128, tile_n=128, use_tma=True))

Data Auto-Profiling

from cutile_gpt import DataAnalyzer

# Auto-detect optimal tile configuration based on data
analyzer = DataAnalyzer()
profile = analyzer.analyze(input_tensor)
print(f"Recommended config: {profile.recommended_config}")

Complete GPT Model

from cutile_gpt import CutileGPT, GPTConfig
import cupy as cp

# Custom config (Core - no extra dependencies)
config = GPTConfig(n_layer=6, n_head=4, n_embd=256)
model = CutileGPT(config)

# Or use presets and load from HuggingFace (requires: pip install cutile-gpt[hf])
model = CutileGPT(GPTConfig.gpt2())
model.load_from_huggingface('gpt2')

# Forward pass
tokens = cp.array([[100, 200, 300]], dtype=cp.int32)
logits, _ = model.forward(tokens)

# Generate
generated = model.generate(
    tokens,
    max_new_tokens=50,
    temperature=0.8,
    top_k=40
)

Benchmark Against PyTorch

# Compare with PyTorch minGPT (requires: pip install cutile-gpt[torch])
uv run python scripts/compare_mingpt.py --benchmark --model tile-medium --batch-size 8 --seq-len 128

# Run HuggingFace inference demo (requires: pip install cutile-gpt[hf])
uv run python scripts/demo_hf_inference.py

📖 API Reference

Core Exports (always available)

from cutile_gpt import (
    # Low-level Kernels
    cutile_gelu,              # GELU activation (8.3x faster)
    cutile_layer_norm,        # Layer normalization
    cutile_linear,            # Matrix multiplication
    cutile_linear_bias,       # Linear with bias
    cutile_embedding,         # Token + position embedding
    cutile_causal_attention,  # Flash Attention
    cutile_fused_mlp,         # Fused Linear→GELU→Linear

    # Tile API (Fluent Builder)
    tile,                     # Create TileOp from tensor
    configure_tiles,          # Set global tile config
    TileConfig,               # Tile size configuration
    TileOp,                   # Fluent operation builder

    # Data Profiling
    DataAnalyzer,             # Auto-detect optimal config
    DataProfile,              # Profile result

    # Model (Core)
    CutileGPT,                # GPT model class
    GPTConfig,                # Model configuration
)

Optional Exports

# Requires: pip install cutile-gpt[hf]
from cutile_gpt import HFWeightLoader  # Load HuggingFace weights
model.load_from_huggingface('gpt2')    # CutileGPT method

# Requires: pip install cutile-gpt[torch]
from cutile_gpt import benchmark_torch  # PyTorch benchmarking

GPTConfig Presets

GPTConfig.gpt_nano()      # 3 layers, 48 dims (testing)
GPTConfig.gpt2()          # 12 layers, 768 dims (117M params)
GPTConfig.gpt2_medium()   # 24 layers, 1024 dims (345M params)
GPTConfig.gpt2_large()    # 36 layers, 1280 dims (774M params)
GPTConfig.gpt2_xl()       # 48 layers, 1600 dims (1.5B params)

🎯 Why cutileGPT?

For Developers

• 87% less code - Focus on WHAT, not HOW
• No manual synchronization - Compiler infers dependencies
• Fewer bugs - No thread indexing errors
• Readable - Clear algorithmic intent

For Deployment

• 200x smaller - ~10MB vs PyTorch's ~2GB
• Edge-ready - Embedded devices
• Serverless-friendly - Lambda-compatible
• Fast builds - Docker-friendly

For Performance

• 8.3x GELU speedup - Compiler-optimized math
• PyTorch competitive - Within 4% on full model
• Auto-tuning - Optimal for each GPU
• Flash Attention - O(N) memory, not O(N²)

For Future

• Hardware portable - Same code, different GPUs
• Compiler updates - Free performance improvements
• No vendor lock-in - Standard tile operations
• Educational - Learn modern GPU programming

📁 Project Structure

cutileGPT/
├── cutile_gpt/                      # 🎯 Core Implementation
│   ├── __init__.py                  # Package exports
│   ├── api/                         # 🔧 High-level Tile API
│   │   ├── tile_op.py               # Fluent Builder API (tile().linear().gelu())
│   │   ├── config.py                # TileConfig, TensorSpec, Layout, DType
│   │   └── profiler.py              # DataAnalyzer for auto-optimization
│   │
│   ├── models/                      # 🧠 GPT Model Implementations
│   │   ├── gpt.py                   # CutileGPT (HuggingFace + minGPT support)
│   │   └── config.py                # GPTConfig with presets
│   │
│   ├── kernels/                     # ⚡ Low-level CUDA Kernels
│   │   ├── gelu.py                  # GELU activation (8.3x speedup)
│   │   ├── layernorm.py             # Layer normalization
│   │   ├── linear.py                # Matrix multiplication
│   │   ├── attention.py             # Flash Attention (O(N) memory)
│   │   ├── embedding.py             # Token + position embeddings
│   │   └── fused_mlp.py             # Fused Linear→GELU→Linear
│   │
│   ├── utils/                       # 🛠️ Utilities
│   │   ├── hf_loader.py             # HuggingFace weight loader
│   │   └── benchmark.py             # Performance benchmarking
│   │
│   └── examples/                    # 📚 Educational Examples
│       ├── linear_tile.py           # Matrix multiplication tutorial
│       ├── attention_tile.py        # Attention tutorial
│       ├── layernorm_tile.py        # LayerNorm tutorial
│       └── gelu_tile.py             # GELU tutorial
│
├── scripts/                         # 🎮 Demo & Benchmark Scripts
│   ├── compare_mingpt.py            # PyTorch minGPT comparison
│   └── demo_hf_inference.py         # HuggingFace inference demo
│
├── demo_tile_gpt.py                 # 🎮 Complete Demo
├── docs/                            # 📖 Documentation
├── profiling_results/               # 📊 Performance data
├── mlir_research/                   # 🧪 Optional MLIR research
└── external/                        # Git submodules (cutile-python, minGPT)

Start here:

• 🎮 demo_tile_gpt.py - Run the complete demo
• 🔧 cutile_gpt/api/ - High-level Tile API reference
• 🧠 cutile_gpt/models/ - GPT model implementation
• 📖 docs/TILE_PHILOSOPHY_DEMO.md - Philosophy deep dive
• 📁 docs/PROJECT_STRUCTURE.md - Complete directory guide

🔬 What is Tile Programming?

Tile Programming is a declarative approach to GPU programming:

1. Specify WHAT operations you want (load, reduce, multiply)
2. Let compiler decide HOW to execute (threads, sync, memory)
3. Achieve better performance through compiler optimization

Example: LayerNorm

# Traditional CUDA: ~150 lines
# - Manual thread indexing (threadIdx.x, blockIdx.x)
# - Explicit shared memory (__shared__ float smem[256])
# - Manual reduction loops (for s = 128; s > 0; s >>= 1)
# - Multiple __syncthreads() calls

# Tile Programming: ~20 lines
@ct.kernel
def layernorm_kernel(X, W, B, Y, eps, N):
    bid = ct.bid(0)  # Block ID only, NO thread IDs!

    x = ct.load(X, index=(bid, 0), shape=(1, N))
    mean = ct.sum(x) / N
    var = ct.sum(x * x) / N - mean * mean
    x_norm = (x - mean) / ct.sqrt(var + eps)
    y = x_norm * W + B
    ct.store(Y, index=(bid, 0), tile=y)

Benefits: 87% code reduction, no manual sync, fewer bugs, better performance.

🏗️ Architecture Layers

cutileGPT is organized into clean hierarchical layers:

┌─────────────────────────────────────────────────────────────┐
│                     User Application                         │
├─────────────────────────────────────────────────────────────┤
│  models/        │ CutileGPT, GPTConfig                       │
│                 │ High-level model with HuggingFace support  │
├─────────────────────────────────────────────────────────────┤
│  api/           │ tile().linear().gelu().execute()           │
│                 │ Fluent Builder + DataAnalyzer              │
├─────────────────────────────────────────────────────────────┤
│  kernels/       │ cutile_gelu, cutile_linear, cutile_attn    │
│                 │ Low-level CUDA Tile kernels                │
├─────────────────────────────────────────────────────────────┤
│  cuda.tile      │ NVIDIA's Tile Programming Framework        │
└─────────────────────────────────────────────────────────────┘

Choose your level:

• High-level: Use CutileGPT for complete models with HuggingFace weights
• Mid-level: Use tile() API for custom declarative operations
• Low-level: Use cutile_* kernels for maximum control

🎓 What We've Proven

cutileGPT demonstrates that Tile Programming Philosophy is practical:

✅ Declarative GPU Programming Works

• Complete GPT with ZERO explicit thread management
• Every operation specifies WHAT, compiler handles HOW
• No manual synchronization anywhere

✅ Performance is Competitive

• 8.3x speedup on GELU kernel vs CuPy
• Competitive with PyTorch on full model
• Compiler optimization is effective

✅ Code is Maintainable

• 87% code reduction vs traditional CUDA
• Readable and clear algorithmic intent
• Easy to modify and extend

✅ The Future of GPU Programming

• Declarative > Imperative - Higher abstraction
• Compiler > Manual - Better optimization
• Portable > Specific - Hardware-independent

🛣️ Roadmap

Completed ✅

• Pure Tile Programming Philosophy GPT
• 8.3x GELU speedup over CuPy
• PyTorch competitive performance
• Flash Attention (O(N) memory)
• Complete demo with all tests passing
• Tile API - Fluent Builder interface (tile().linear().gelu().execute())
• Data Profiler - Auto-detection of optimal tile configurations
• HuggingFace Integration - Load pre-trained GPT-2 weights
• Hierarchical Architecture - Clean separation (api, models, kernels, utils)

Future Work 🔮

• FP16/BF16 support for 2-3x speedup
• KV cache for efficient generation
• Multi-GPU support via NCCL
• INT8 quantization kernels
• Auto-tuning for tile sizes

📚 Learn More

• 🎮 demo_tile_gpt.py - Run the demo!
• 🔧 cutile_gpt/api/ - Tile API reference (Fluent Builder, Config, Profiler)
• 🧠 cutile_gpt/models/ - GPT model & config documentation
• ⚡ cutile_gpt/kernels/ - Low-level kernel implementations
• 📚 cutile_gpt/examples/ - Educational tile programming tutorials
• 📖 docs/TILE_PHILOSOPHY_DEMO.md - Complete philosophy documentation
• 🏗️ docs/ARCHITECTURE_VISION.md - Project vision & roadmap

🤝 Contributing

Bug reports, benchmark numbers from other Blackwell GPUs, and documentation fixes are all welcome - see CONTRIBUTING.md for setup and the checks CI runs.

Published performance numbers come from a single NVIDIA GB10, so results from different hardware are especially useful; there is a benchmark result issue template for exactly that. Open-ended questions belong in Discussions.

You do not need a GPU to contribute - CI is GPU-free, and docs, packaging, and the visualization scripts all run without one.

📄 License

Apache-2.0 - See LICENSE for details.

🙏 Acknowledgments

• NVIDIA CUDA Tile - Declarative GPU programming framework
• Andrej Karpathy's minGPT - Reference architecture
• CuPy - NumPy-compatible GPU arrays
• Flash Attention - Online softmax algorithm (Dao et al., 2022)