Skip to content

Commit 7d5d8fc

Browse files
committed
Add cute cache
stack-info: PR: #57, branch: drisspg/stack/15
1 parent 3bd431f commit 7d5d8fc

4 files changed

Lines changed: 299 additions & 4 deletions

File tree

test_cute_cache.py

Lines changed: 112 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,112 @@
1+
#!/usr/bin/env python3
2+
"""Test script to demonstrate CUTE kernel caching"""
3+
4+
import torch
5+
from operator import add, mul
6+
import cutlass
7+
import cutlass.cute as cute
8+
from cutlass.cute.runtime import from_dlpack
9+
10+
from transformer_nuggets.cute import cute_cache, get_cache_stats, clear_cute_cache
11+
12+
13+
@cute.kernel
14+
def simple_elementwise_kernel(
15+
op: cutlass.Constexpr,
16+
gA: cute.Tensor,
17+
gB: cute.Tensor,
18+
gC: cute.Tensor,
19+
):
20+
tidx, _, _ = cute.arch.thread_idx()
21+
bidx, _, _ = cute.arch.block_idx()
22+
bdim, _, _ = cute.arch.block_dim()
23+
24+
thread_idx = bidx * bdim + tidx
25+
26+
m, n = gA.shape
27+
ni = thread_idx % n
28+
mi = thread_idx // n
29+
30+
a_val = gA[mi, ni]
31+
b_val = gB[mi, ni]
32+
33+
gC[mi, ni] = op(a_val, b_val)
34+
35+
36+
@cute_cache()
37+
def cached_elementwise(
38+
op: cutlass.Constexpr,
39+
mA: cute.Tensor,
40+
mB: cute.Tensor,
41+
mC: cute.Tensor,
42+
):
43+
num_threads_per_block = 256
44+
m, n = mA.shape
45+
46+
kernel = simple_elementwise_kernel(op, mA, mB, mC)
47+
kernel.launch(
48+
grid=((m * n) // num_threads_per_block, 1, 1), block=(num_threads_per_block, 1, 1)
49+
)
50+
51+
52+
def main():
53+
print("CUTE Kernel Caching Demo")
54+
print("=" * 50)
55+
56+
# Create test tensors
57+
M, N = 1024, 1024
58+
a = torch.randn(M, N, device="cuda", dtype=torch.float16)
59+
b = torch.randn(M, N, device="cuda", dtype=torch.float16)
60+
c = torch.zeros(M, N, device="cuda", dtype=torch.float16)
61+
62+
# Convert to CUTE tensors
63+
a_ = from_dlpack(a, assumed_align=16)
64+
b_ = from_dlpack(b, assumed_align=16)
65+
c_ = from_dlpack(c, assumed_align=16)
66+
67+
print("\n1. First call with 'add' - should compile:")
68+
cached_elementwise(add, a_, b_, c_)
69+
torch.testing.assert_close(c, a + b)
70+
print("✓ Addition result correct")
71+
72+
print("\n2. Second call with 'add' - should hit cache:")
73+
c.zero_()
74+
cached_elementwise(add, a_, b_, c_)
75+
76+
print("\n3. First call with 'mul' - should compile (different op):")
77+
c.zero_()
78+
cached_elementwise(mul, a_, b_, c_)
79+
torch.testing.assert_close(c, a * b)
80+
print("✓ Multiplication result correct")
81+
82+
print("\n4. Different tensor sizes - should compile:")
83+
a2 = torch.randn(512, 512, device="cuda", dtype=torch.float16)
84+
b2 = torch.randn(512, 512, device="cuda", dtype=torch.float16)
85+
c2 = torch.zeros(512, 512, device="cuda", dtype=torch.float16)
86+
87+
a2_ = from_dlpack(a2, assumed_align=16)
88+
b2_ = from_dlpack(b2, assumed_align=16)
89+
c2_ = from_dlpack(c2, assumed_align=16)
90+
91+
cached_elementwise(add, a2_, b2_, c2_)
92+
93+
print("\n5. Same size as #4 - should hit cache:")
94+
cached_elementwise(add, a2_, b2_, c2_)
95+
96+
print("\n" + "=" * 50)
97+
print("Final Cache Statistics:")
98+
stats = get_cache_stats()
99+
print(f" Total calls: {stats['total']}")
100+
print(f" Cache hits: {stats['hits']}")
101+
print(f" Cache misses: {stats['misses']}")
102+
print(f" Hit rate: {stats['hit_rate']:.2%}")
103+
print(f" Unique kernels cached: {stats['cache_size']}")
104+
105+
print("\nClearing cache...")
106+
clear_cute_cache()
107+
stats = get_cache_stats()
108+
print(f" Cache size after clear: {stats['cache_size']}")
109+
110+
111+
if __name__ == "__main__":
112+
main()
Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,5 @@
1+
from transformer_nuggets.cute.cache import (
2+
cute_cache,
3+
clear_cute_cache,
4+
get_cache_stats,
5+
)

transformer_nuggets/cute/add.py

Lines changed: 14 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -7,6 +7,7 @@
77
from cutlass.cute.runtime import from_dlpack
88

99
from transformer_nuggets.utils.benchmark import benchmark_cuda_function_in_microseconds
10+
from transformer_nuggets.cute.cache import cute_cache, get_cache_stats
1011

1112

1213
@cute.kernel
@@ -72,7 +73,7 @@ def naive_elementwise_add(mA: cute.Tensor, mB: cute.Tensor, mC: cute.Tensor):
7273
)
7374

7475

75-
@cute.jit
76+
@cute_cache()
7677
def elementwise_op(
7778
op: cutlass.Constexpr,
7879
mA: cute.Tensor,
@@ -123,7 +124,9 @@ def benchmark(callable, *, num_warmups=5, num_iterations=200):
123124
c.zero_()
124125

125126
naive_elementwise_add_compiled = cute.compile(naive_elementwise_add, a_, b_, c_)
126-
elementwise_add_compiled = cute.compile(elementwise_op, add, a_, b_, c_)
127+
128+
# Note: elementwise_op is now cached automatically!
129+
# First call will compile, subsequent calls use cache
127130

128131
print("\n" + "=" * 50)
129132
print("BENCHMARK RESULTS")
@@ -135,5 +138,12 @@ def benchmark(callable, *, num_warmups=5, num_iterations=200):
135138
print("2. PyTorch add (baseline):")
136139
benchmark(partial(torch.add, a, b, out=c))
137140

138-
print("3. Optimized elementwise add kernel:")
139-
benchmark(partial(elementwise_add_compiled, a_, b_, c_))
141+
print("3. Optimized elementwise add kernel (with caching):")
142+
benchmark(partial(elementwise_op, add, a_, b_, c_))
143+
144+
print("\nCache Statistics:")
145+
stats = get_cache_stats()
146+
print(f" Cache hits: {stats['hits']}")
147+
print(f" Cache misses: {stats['misses']}")
148+
print(f" Hit rate: {stats['hit_rate']:.2%}")
149+
print(f" Cache size: {stats['cache_size']}")

transformer_nuggets/cute/cache.py

Lines changed: 168 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,168 @@
1+
import functools
2+
import hashlib
3+
import logging
4+
import threading
5+
from typing import Any
6+
from collections.abc import Callable
7+
8+
import cutlass.cute as cute
9+
10+
11+
logger = logging.getLogger(__name__)
12+
13+
14+
class CuteKernelCache:
15+
"""Global cache for compiled CUTE kernels"""
16+
17+
def __init__(self):
18+
self._cache: dict[str, Any] = {}
19+
self._stats = {"hits": 0, "misses": 0}
20+
self._lock = threading.Lock()
21+
22+
def get(self, key: str) -> Any | None:
23+
with self._lock:
24+
if key in self._cache:
25+
self._stats["hits"] += 1
26+
return self._cache[key]
27+
self._stats["misses"] += 1
28+
return None
29+
30+
def set(self, key: str, value: Any) -> None:
31+
with self._lock:
32+
self._cache[key] = value
33+
34+
def clear(self) -> None:
35+
with self._lock:
36+
self._cache.clear()
37+
self._stats = {"hits": 0, "misses": 0}
38+
39+
def get_stats(self) -> dict[str, int | float]:
40+
with self._lock:
41+
total = self._stats["hits"] + self._stats["misses"]
42+
hit_rate = self._stats["hits"] / total if total > 0 else 0
43+
return {
44+
**self._stats,
45+
"total": total,
46+
"hit_rate": hit_rate,
47+
"cache_size": len(self._cache),
48+
}
49+
50+
51+
# Global cache instance
52+
_kernel_cache = CuteKernelCache()
53+
54+
55+
def _generate_cache_key(*args, **kwargs) -> str:
56+
"""Generate a cache key from function arguments"""
57+
key_parts = []
58+
59+
for arg in args:
60+
if isinstance(arg, cute.Tensor):
61+
# For cute tensors, use shape and dtype
62+
key_parts.append(f"tensor_shape={arg.shape}_dtype={arg._dtype}")
63+
elif hasattr(arg, "__name__"):
64+
# For functions/operators like add, mul
65+
key_parts.append(f"op={arg.__name__}")
66+
else:
67+
# For other types, use string representation
68+
key_parts.append(str(arg))
69+
70+
# Add kwargs
71+
for k, v in sorted(kwargs.items()):
72+
key_parts.append(f"{k}={v}")
73+
74+
# Create hash of the key
75+
key_str = "_".join(key_parts)
76+
return hashlib.md5(key_str.encode()).hexdigest()
77+
78+
79+
# Keep cute_jit for backward compatibility
80+
def cute_jit(func: Callable) -> Callable:
81+
"""
82+
Deprecated: Use @cute_cache() instead.
83+
This decorator is kept for backward compatibility.
84+
"""
85+
logger.warning("@cute_jit is deprecated. Use @cute_cache() instead.")
86+
return cute_cache()(func)
87+
88+
89+
def cute_cache(key_params: list | None = None) -> Callable:
90+
"""
91+
Decorator that applies @cute.jit and caches compiled CUTE kernels.
92+
93+
Args:
94+
key_params: List of parameter names to use for cache key generation.
95+
If None, uses all parameters.
96+
97+
Example:
98+
@cute_cache()
99+
def elementwise_op(op, mA, mB, mC):
100+
# Kernel implementation
101+
...
102+
"""
103+
104+
def decorator(func: Callable) -> Callable:
105+
# Apply cute.jit once at decoration time to avoid closure issues
106+
jitted_func = cute.jit(func)
107+
108+
@functools.wraps(func)
109+
def wrapper(*args, **kwargs):
110+
# Generate cache key
111+
if key_params:
112+
# Extract specific params for cache key
113+
import inspect
114+
115+
sig = inspect.signature(func)
116+
bound_args = sig.bind(*args, **kwargs)
117+
bound_args.apply_defaults()
118+
119+
key_args = []
120+
for param in key_params:
121+
if param in bound_args.arguments:
122+
key_args.append(bound_args.arguments[param])
123+
cache_key = _generate_cache_key(*key_args)
124+
else:
125+
# Use all arguments
126+
cache_key = _generate_cache_key(*args, **kwargs)
127+
128+
# Add function name to cache key
129+
cache_key = f"{func.__name__}_{cache_key}"
130+
131+
# Check cache
132+
compiled_kernel = _kernel_cache.get(cache_key)
133+
134+
if compiled_kernel is not None:
135+
logger.debug(f"Cache hit for {func.__name__} (key: {cache_key[:8]}...)")
136+
return compiled_kernel(*args, **kwargs)
137+
138+
logger.debug(
139+
f"Cache miss for {func.__name__} (key: {cache_key[:8]}...) - Compiling..."
140+
)
141+
142+
# Compile the kernel using the pre-jitted function
143+
compiled_kernel = cute.compile(jitted_func, *args, **kwargs)
144+
145+
# Cache the compiled kernel
146+
_kernel_cache.set(cache_key, compiled_kernel)
147+
148+
# Execute the compiled kernel
149+
return compiled_kernel(*args, **kwargs)
150+
151+
# Add utility methods to the wrapper
152+
wrapper.clear_cache = lambda: _kernel_cache.clear()
153+
wrapper.get_cache_stats = lambda: _kernel_cache.get_stats()
154+
155+
return wrapper
156+
157+
return decorator
158+
159+
160+
# Utility functions
161+
def clear_cute_cache():
162+
"""Clear all cached kernels"""
163+
_kernel_cache.clear()
164+
165+
166+
def get_cache_stats():
167+
"""Get cache statistics"""
168+
return _kernel_cache.get_stats()

0 commit comments

Comments
 (0)