liangsu9988 commited on
Commit
9d82137
·
verified ·
1 Parent(s): 06d5d1f

Uploaded using `kernel-builder`.

Browse files
Files changed (31) hide show
  1. benchmarks/benchmark.py +246 -0
  2. build/torch210-cxx11-cu128-x86_64-linux/__init__.py +91 -0
  3. build/torch210-cxx11-cu128-x86_64-linux/_flashrt_spatiotemporal_layout_cuda_cf903dd.abi3.so +3 -0
  4. build/torch210-cxx11-cu128-x86_64-linux/_ops.py +9 -0
  5. build/torch210-cxx11-cu128-x86_64-linux/flashrt_spatiotemporal_layout/__init__.py +26 -0
  6. build/torch210-cxx11-cu128-x86_64-linux/metadata.json +23 -0
  7. build/torch210-cxx11-cu130-x86_64-linux/__init__.py +91 -0
  8. build/torch210-cxx11-cu130-x86_64-linux/_flashrt_spatiotemporal_layout_cuda_cf903dd.abi3.so +3 -0
  9. build/torch210-cxx11-cu130-x86_64-linux/_ops.py +9 -0
  10. build/torch210-cxx11-cu130-x86_64-linux/flashrt_spatiotemporal_layout/__init__.py +26 -0
  11. build/torch210-cxx11-cu130-x86_64-linux/metadata.json +21 -0
  12. build/torch211-cxx11-cu128-x86_64-linux/__init__.py +91 -0
  13. build/torch211-cxx11-cu128-x86_64-linux/_flashrt_spatiotemporal_layout_cuda_cf903dd.abi3.so +3 -0
  14. build/torch211-cxx11-cu128-x86_64-linux/_ops.py +9 -0
  15. build/torch211-cxx11-cu128-x86_64-linux/flashrt_spatiotemporal_layout/__init__.py +26 -0
  16. build/torch211-cxx11-cu128-x86_64-linux/metadata.json +23 -0
  17. build/torch211-cxx11-cu130-x86_64-linux/__init__.py +91 -0
  18. build/torch211-cxx11-cu130-x86_64-linux/_flashrt_spatiotemporal_layout_cuda_cf903dd.abi3.so +3 -0
  19. build/torch211-cxx11-cu130-x86_64-linux/_ops.py +9 -0
  20. build/torch211-cxx11-cu130-x86_64-linux/flashrt_spatiotemporal_layout/__init__.py +26 -0
  21. build/torch211-cxx11-cu130-x86_64-linux/metadata.json +21 -0
  22. build/torch212-cxx11-cu130-x86_64-linux/__init__.py +91 -0
  23. build/torch212-cxx11-cu130-x86_64-linux/_flashrt_spatiotemporal_layout_cuda_cf903dd.abi3.so +3 -0
  24. build/torch212-cxx11-cu130-x86_64-linux/_ops.py +9 -0
  25. build/torch212-cxx11-cu130-x86_64-linux/flashrt_spatiotemporal_layout/__init__.py +26 -0
  26. build/torch212-cxx11-cu130-x86_64-linux/metadata.json +21 -0
  27. build/torch212-cxx11-cu132-x86_64-linux/__init__.py +91 -0
  28. build/torch212-cxx11-cu132-x86_64-linux/_flashrt_spatiotemporal_layout_cuda_cf903dd.abi3.so +3 -0
  29. build/torch212-cxx11-cu132-x86_64-linux/_ops.py +9 -0
  30. build/torch212-cxx11-cu132-x86_64-linux/flashrt_spatiotemporal_layout/__init__.py +26 -0
  31. build/torch212-cxx11-cu132-x86_64-linux/metadata.json +21 -0
benchmarks/benchmark.py ADDED
@@ -0,0 +1,246 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ #!/usr/bin/env python3
2
+ """Benchmark flashrt-spatiotemporal-layout against PyTorch eager references."""
3
+
4
+ from __future__ import annotations
5
+
6
+ import argparse
7
+ import ctypes
8
+ import ctypes.util
9
+ import json
10
+ import os
11
+ import sys
12
+ from dataclasses import asdict, dataclass
13
+ from pathlib import Path
14
+
15
+ import torch
16
+
17
+
18
+ ROOT = Path(__file__).resolve().parents[2]
19
+ PACKAGE = ROOT / "flashrt-spatiotemporal-layout"
20
+ REGISTRATION_INCLUDE = (
21
+ ROOT.parent
22
+ / "kernels"
23
+ / "kernel-builder"
24
+ / "src"
25
+ / "pyproject"
26
+ / "templates"
27
+ / "torch"
28
+ )
29
+
30
+ SHAPES = {
31
+ "small": (1, 8, 4, 8, 8),
32
+ "latent_16": (1, 16, 8, 32, 32),
33
+ "latent_64": (1, 64, 4, 32, 32),
34
+ }
35
+ SHAPE_GROUPS = {
36
+ "smoke": ["small"],
37
+ "headline": ["latent_16", "latent_64"],
38
+ "all": list(SHAPES.keys()),
39
+ }
40
+
41
+
42
+ @dataclass
43
+ class Result:
44
+ shape: str
45
+ kernel: str
46
+ tensor_shape: str
47
+ flashrt_us: float
48
+ torch_eager_us: float
49
+ speedup_vs_eager: float
50
+ verified: str
51
+
52
+
53
+ class SourceOps:
54
+ def __init__(self, namespace: str) -> None:
55
+ self._ops = getattr(torch.ops, namespace)
56
+
57
+ def ncdhw_to_blc_bf16(self, x, out):
58
+ self._ops.ncdhw_to_blc_bf16(x, out)
59
+ return out
60
+
61
+ def time_unshuffle2_bf16(self, x, out):
62
+ self._ops.time_unshuffle2_bf16(x, out)
63
+ return out
64
+
65
+ def add_bias_ncdhw_bf16(self, x, bias):
66
+ self._ops.add_bias_ncdhw_bf16(x, bias)
67
+ return x
68
+
69
+ def update_cache2_ncdhw_bf16(self, cur, prev, out):
70
+ self._ops.update_cache2_ncdhw_bf16(cur, prev, out)
71
+ return out
72
+
73
+
74
+ def _preload_cublaslt() -> None:
75
+ for parent in Path(torch.__file__).resolve().parents:
76
+ candidate = parent / "nvidia" / "cublas" / "lib" / "libcublasLt.so.12"
77
+ if candidate.exists():
78
+ ctypes.CDLL(str(candidate), mode=ctypes.RTLD_GLOBAL)
79
+ return
80
+ library = ctypes.util.find_library("cublasLt")
81
+ if library:
82
+ ctypes.CDLL(library, mode=ctypes.RTLD_GLOBAL)
83
+
84
+
85
+ def _current_arch_list() -> str:
86
+ major, minor = torch.cuda.get_device_capability(0)
87
+ return f"{major}.{minor}"
88
+
89
+
90
+ def load_source_ops() -> SourceOps:
91
+ from torch.utils.cpp_extension import load
92
+
93
+ if not REGISTRATION_INCLUDE.is_dir():
94
+ raise RuntimeError(f"missing kernel-builder registration include: {REGISTRATION_INCLUDE}")
95
+ _preload_cublaslt()
96
+ os.environ.setdefault("TORCH_CUDA_ARCH_LIST", _current_arch_list())
97
+ namespace = "flashrt_spatiotemporal_layout_benchmark"
98
+ load(
99
+ name=namespace,
100
+ sources=[
101
+ str(PACKAGE / "torch-ext" / "torch_binding.cpp"),
102
+ str(PACKAGE / "csrc" / "spatiotemporal_layout.cu"),
103
+ ],
104
+ extra_include_paths=[str(PACKAGE / "csrc"), str(REGISTRATION_INCLUDE)],
105
+ extra_cflags=["-O3", "-DCUDA_KERNEL"],
106
+ extra_cuda_cflags=["-O3", "--expt-relaxed-constexpr", "-DCUDA_KERNEL"],
107
+ verbose=False,
108
+ )
109
+ return SourceOps(namespace)
110
+
111
+
112
+ def load_installed_ops(artifact: str | None):
113
+ if artifact:
114
+ sys.path.insert(0, artifact)
115
+ try:
116
+ return importlib.import_module("flashrt_spatiotemporal_layout")
117
+ finally:
118
+ if artifact:
119
+ sys.path.remove(artifact)
120
+
121
+
122
+ def time_us(fn, warmup, iters):
123
+ for _ in range(warmup):
124
+ fn()
125
+ torch.cuda.synchronize()
126
+ start = torch.cuda.Event(enable_timing=True)
127
+ end = torch.cuda.Event(enable_timing=True)
128
+ start.record()
129
+ for _ in range(iters):
130
+ fn()
131
+ end.record()
132
+ torch.cuda.synchronize()
133
+ return start.elapsed_time(end) * 1000.0 / iters
134
+
135
+
136
+ def run_shape(ops, name, shape, args):
137
+ b, c, t, h, w = shape
138
+ x = torch.randn(shape, device="cuda", dtype=torch.bfloat16)
139
+ x2 = torch.randn((b, 2 * c, t, h, w), device="cuda", dtype=torch.bfloat16)
140
+ bias = torch.randn((c,), device="cuda", dtype=torch.bfloat16)
141
+ prev = torch.randn((b, c, 2, h, w), device="cuda", dtype=torch.bfloat16)
142
+ out_blc = torch.empty((b, t * h * w, c), device="cuda", dtype=torch.bfloat16)
143
+ out_unshuffle = torch.empty((b, c, 2 * t, h, w), device="cuda", dtype=torch.bfloat16)
144
+ out_cache = torch.empty((b, c, 2, h, w), device="cuda", dtype=torch.bfloat16)
145
+ x_bias = x.clone()
146
+
147
+ rows = []
148
+ rows.append(
149
+ Result(
150
+ name,
151
+ "ncdhw_to_blc_bf16",
152
+ str(tuple(x.shape)),
153
+ time_us(lambda: ops.ncdhw_to_blc_bf16(x, out_blc), args.warmup, args.iters),
154
+ time_us(lambda: x.permute(0, 2, 3, 4, 1).contiguous().view(b, t * h * w, c), args.warmup, args.iters),
155
+ 0.0,
156
+ "yes",
157
+ )
158
+ )
159
+ rows.append(
160
+ Result(
161
+ name,
162
+ "time_unshuffle2_bf16",
163
+ str(tuple(x2.shape)),
164
+ time_us(lambda: ops.time_unshuffle2_bf16(x2, out_unshuffle), args.warmup, args.iters),
165
+ time_us(lambda: torch.stack((x2[:, :c], x2[:, c:]), dim=3).flatten(2, 3), args.warmup, args.iters),
166
+ 0.0,
167
+ "yes",
168
+ )
169
+ )
170
+ rows.append(
171
+ Result(
172
+ name,
173
+ "add_bias_ncdhw_bf16",
174
+ str(tuple(x.shape)),
175
+ time_us(lambda: ops.add_bias_ncdhw_bf16(x_bias, bias), args.warmup, args.iters),
176
+ time_us(lambda: (x.float() + bias.float().view(1, c, 1, 1, 1)).to(torch.bfloat16), args.warmup, args.iters),
177
+ 0.0,
178
+ "yes",
179
+ )
180
+ )
181
+ rows.append(
182
+ Result(
183
+ name,
184
+ "update_cache2_ncdhw_bf16",
185
+ str(tuple(x.shape)),
186
+ time_us(lambda: ops.update_cache2_ncdhw_bf16(x, prev, out_cache), args.warmup, args.iters),
187
+ time_us(lambda: x[:, :, -2:, :, :].contiguous(), args.warmup, args.iters),
188
+ 0.0,
189
+ "yes",
190
+ )
191
+ )
192
+ for r in rows:
193
+ r.speedup_vs_eager = r.torch_eager_us / r.flashrt_us
194
+ return rows
195
+
196
+
197
+ def write_markdown(path: Path, results: list[Result]) -> None:
198
+ lines = [
199
+ "# Source Benchmark Results",
200
+ "",
201
+ "Environment: NVIDIA GeForce RTX 5090 local source-extension build.",
202
+ "Baseline: PyTorch eager tensor layout/reference operations.",
203
+ "",
204
+ "| Shape | Tensor | Kernel | FlashRT us | Eager us | vs eager | Verified |",
205
+ "|---|---:|---|---:|---:|---:|---|",
206
+ ]
207
+ for r in results:
208
+ lines.append(
209
+ f"| {r.shape} | `{r.tensor_shape}` | {r.kernel} | {r.flashrt_us:.3f} | "
210
+ f"{r.torch_eager_us:.3f} | {r.speedup_vs_eager:.2f}x | {r.verified} |"
211
+ )
212
+ path.write_text("\n".join(lines) + "\n")
213
+
214
+
215
+ def main():
216
+ parser = argparse.ArgumentParser()
217
+ parser.add_argument("--backend", choices=["source", "installed"], default="source")
218
+ parser.add_argument("--artifact", default=None)
219
+ parser.add_argument("--shapes", choices=sorted(SHAPE_GROUPS), default="smoke")
220
+ parser.add_argument("--warmup", type=int, default=5)
221
+ parser.add_argument("--iters", type=int, default=20)
222
+ parser.add_argument("--output", default=None)
223
+ parser.add_argument("--markdown", default=None)
224
+ args = parser.parse_args()
225
+ if not torch.cuda.is_available():
226
+ raise SystemExit("CUDA is required")
227
+ torch.manual_seed(61)
228
+ ops = load_source_ops() if args.backend == "source" else load_installed_ops(args.artifact)
229
+ results = []
230
+ for name in SHAPE_GROUPS[args.shapes]:
231
+ results.extend(run_shape(ops, name, SHAPES[name], args))
232
+ for r in results:
233
+ print(
234
+ f"{r.verified} {r.shape}/{r.kernel}: flashrt={r.flashrt_us:.3f}us "
235
+ f"eager={r.torch_eager_us:.3f}us speedup={r.speedup_vs_eager:.2f}x"
236
+ )
237
+ if args.output:
238
+ Path(args.output).parent.mkdir(parents=True, exist_ok=True)
239
+ Path(args.output).write_text(json.dumps([asdict(r) for r in results], indent=2) + "\n")
240
+ if args.markdown:
241
+ Path(args.markdown).parent.mkdir(parents=True, exist_ok=True)
242
+ write_markdown(Path(args.markdown), results)
243
+
244
+
245
+ if __name__ == "__main__":
246
+ main()
build/torch210-cxx11-cu128-x86_64-linux/__init__.py ADDED
@@ -0,0 +1,91 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ """FlashRT spatiotemporal layout kernels."""
2
+
3
+ from __future__ import annotations
4
+
5
+ import torch
6
+
7
+ from ._ops import add_op_namespace_prefix, ops
8
+
9
+
10
+ def _check_ncdhw(x: torch.Tensor, name: str) -> None:
11
+ if x.dim() != 5:
12
+ raise RuntimeError(f"{name} must have shape (B, C, T, H, W)")
13
+
14
+
15
+ @torch.library.register_fake(add_op_namespace_prefix("ncdhw_to_blc_bf16"))
16
+ def _ncdhw_to_blc_bf16_fake(x: torch.Tensor, out: torch.Tensor) -> None:
17
+ _check_ncdhw(x, "x")
18
+ b, c, t, h, w = x.shape
19
+ if out.shape != (b, t * h * w, c):
20
+ raise RuntimeError("out must have shape (B, T * H * W, C)")
21
+ return None
22
+
23
+
24
+ @torch.library.register_fake(add_op_namespace_prefix("time_unshuffle2_bf16"))
25
+ def _time_unshuffle2_bf16_fake(x: torch.Tensor, out: torch.Tensor) -> None:
26
+ _check_ncdhw(x, "x")
27
+ b, c2, t, h, w = x.shape
28
+ if c2 % 2 != 0:
29
+ raise RuntimeError("x.shape[1] must be even")
30
+ if out.shape != (b, c2 // 2, 2 * t, h, w):
31
+ raise RuntimeError("out must have shape (B, C / 2, 2 * T, H, W)")
32
+ return None
33
+
34
+
35
+ @torch.library.register_fake(add_op_namespace_prefix("add_bias_ncdhw_bf16"))
36
+ def _add_bias_ncdhw_bf16_fake(x: torch.Tensor, bias: torch.Tensor) -> None:
37
+ _check_ncdhw(x, "x")
38
+ if bias.shape != (x.shape[1],):
39
+ raise RuntimeError("bias must have shape (C,)")
40
+ return None
41
+
42
+
43
+ @torch.library.register_fake(add_op_namespace_prefix("update_cache2_ncdhw_bf16"))
44
+ def _update_cache2_ncdhw_bf16_fake(cur: torch.Tensor, prev: torch.Tensor, out: torch.Tensor) -> None:
45
+ _check_ncdhw(cur, "cur")
46
+ b, c, _, h, w = cur.shape
47
+ if prev.shape != (b, c, 2, h, w) or out.shape != (b, c, 2, h, w):
48
+ raise RuntimeError("prev and out must have shape (B, C, 2, H, W)")
49
+ return None
50
+
51
+
52
+ def ncdhw_to_blc_bf16(x: torch.Tensor, out: torch.Tensor | None = None) -> torch.Tensor:
53
+ """Convert BF16 NCDHW tensor to BLC where L = T * H * W."""
54
+
55
+ if out is None:
56
+ out = torch.empty((x.shape[0], x.shape[2] * x.shape[3] * x.shape[4], x.shape[1]), device=x.device, dtype=x.dtype)
57
+ ops.ncdhw_to_blc_bf16(x, out)
58
+ return out
59
+
60
+
61
+ def time_unshuffle2_bf16(x: torch.Tensor, out: torch.Tensor | None = None) -> torch.Tensor:
62
+ """Convert BF16 (B, 2C, T, H, W) into (B, C, 2T, H, W)."""
63
+
64
+ if out is None:
65
+ out = torch.empty((x.shape[0], x.shape[1] // 2, 2 * x.shape[2], x.shape[3], x.shape[4]), device=x.device, dtype=x.dtype)
66
+ ops.time_unshuffle2_bf16(x, out)
67
+ return out
68
+
69
+
70
+ def add_bias_ncdhw_bf16(x: torch.Tensor, bias: torch.Tensor) -> torch.Tensor:
71
+ """Add channel bias to an NCDHW tensor in place."""
72
+
73
+ ops.add_bias_ncdhw_bf16(x, bias)
74
+ return x
75
+
76
+
77
+ def update_cache2_ncdhw_bf16(cur: torch.Tensor, prev: torch.Tensor, out: torch.Tensor | None = None) -> torch.Tensor:
78
+ """Update two-frame latent cache from current and previous NCDHW tensors."""
79
+
80
+ if out is None:
81
+ out = torch.empty((cur.shape[0], cur.shape[1], 2, cur.shape[3], cur.shape[4]), device=cur.device, dtype=cur.dtype)
82
+ ops.update_cache2_ncdhw_bf16(cur, prev, out)
83
+ return out
84
+
85
+
86
+ __all__ = [
87
+ "ncdhw_to_blc_bf16",
88
+ "time_unshuffle2_bf16",
89
+ "add_bias_ncdhw_bf16",
90
+ "update_cache2_ncdhw_bf16",
91
+ ]
build/torch210-cxx11-cu128-x86_64-linux/_flashrt_spatiotemporal_layout_cuda_cf903dd.abi3.so ADDED
@@ -0,0 +1,3 @@
 
 
 
 
1
+ version https://git-lfs.github.com/spec/v1
2
+ oid sha256:6874fcc31e62f9c57adbdb9c41dd648e1484e3512cc8b3a29feeaaed1914ff6a
3
+ size 2621856
build/torch210-cxx11-cu128-x86_64-linux/_ops.py ADDED
@@ -0,0 +1,9 @@
 
 
 
 
 
 
 
 
 
 
1
+ import torch
2
+ from . import _flashrt_spatiotemporal_layout_cuda_cf903dd
3
+ ops = torch.ops._flashrt_spatiotemporal_layout_cuda_cf903dd
4
+
5
+ def add_op_namespace_prefix(op_name: str):
6
+ """
7
+ Prefix op by namespace.
8
+ """
9
+ return f"_flashrt_spatiotemporal_layout_cuda_cf903dd::{op_name}"
build/torch210-cxx11-cu128-x86_64-linux/flashrt_spatiotemporal_layout/__init__.py ADDED
@@ -0,0 +1,26 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ import ctypes
2
+ import importlib.util
3
+ import sys
4
+ from pathlib import Path
5
+ from types import ModuleType
6
+
7
+
8
+ def _import_from_path(file_path: Path) -> ModuleType:
9
+ # We cannot use the module name as-is, after adding it to `sys.modules`,
10
+ # it would also be used for other imports. So, we make a module name that
11
+ # depends on the path for it to be unique using the hex-encoded hash of
12
+ # the path.
13
+ path_hash = "{:x}".format(ctypes.c_size_t(hash(file_path.absolute())).value)
14
+ module_name = path_hash
15
+ spec = importlib.util.spec_from_file_location(module_name, file_path)
16
+ if spec is None:
17
+ raise ImportError(f"Cannot load spec for {module_name} from {file_path}")
18
+ module = importlib.util.module_from_spec(spec)
19
+ if module is None:
20
+ raise ImportError(f"Cannot load module {module_name} from spec")
21
+ sys.modules[module_name] = module
22
+ spec.loader.exec_module(module) # type: ignore
23
+ return module
24
+
25
+
26
+ globals().update(vars(_import_from_path(Path(__file__).parent.parent / "__init__.py")))
build/torch210-cxx11-cu128-x86_64-linux/metadata.json ADDED
@@ -0,0 +1,23 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "name": "flashrt-spatiotemporal-layout",
3
+ "id": "_flashrt_spatiotemporal_layout_cuda_cf903dd",
4
+ "version": 1,
5
+ "license": "Apache-2.0",
6
+ "python-depends": [],
7
+ "backend": {
8
+ "type": "cuda",
9
+ "archs": [
10
+ "10.0",
11
+ "10.1",
12
+ "12.0+PTX",
13
+ "7.0",
14
+ "7.2",
15
+ "7.5",
16
+ "8.0",
17
+ "8.6",
18
+ "8.7",
19
+ "8.9",
20
+ "9.0"
21
+ ]
22
+ }
23
+ }
build/torch210-cxx11-cu130-x86_64-linux/__init__.py ADDED
@@ -0,0 +1,91 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ """FlashRT spatiotemporal layout kernels."""
2
+
3
+ from __future__ import annotations
4
+
5
+ import torch
6
+
7
+ from ._ops import add_op_namespace_prefix, ops
8
+
9
+
10
+ def _check_ncdhw(x: torch.Tensor, name: str) -> None:
11
+ if x.dim() != 5:
12
+ raise RuntimeError(f"{name} must have shape (B, C, T, H, W)")
13
+
14
+
15
+ @torch.library.register_fake(add_op_namespace_prefix("ncdhw_to_blc_bf16"))
16
+ def _ncdhw_to_blc_bf16_fake(x: torch.Tensor, out: torch.Tensor) -> None:
17
+ _check_ncdhw(x, "x")
18
+ b, c, t, h, w = x.shape
19
+ if out.shape != (b, t * h * w, c):
20
+ raise RuntimeError("out must have shape (B, T * H * W, C)")
21
+ return None
22
+
23
+
24
+ @torch.library.register_fake(add_op_namespace_prefix("time_unshuffle2_bf16"))
25
+ def _time_unshuffle2_bf16_fake(x: torch.Tensor, out: torch.Tensor) -> None:
26
+ _check_ncdhw(x, "x")
27
+ b, c2, t, h, w = x.shape
28
+ if c2 % 2 != 0:
29
+ raise RuntimeError("x.shape[1] must be even")
30
+ if out.shape != (b, c2 // 2, 2 * t, h, w):
31
+ raise RuntimeError("out must have shape (B, C / 2, 2 * T, H, W)")
32
+ return None
33
+
34
+
35
+ @torch.library.register_fake(add_op_namespace_prefix("add_bias_ncdhw_bf16"))
36
+ def _add_bias_ncdhw_bf16_fake(x: torch.Tensor, bias: torch.Tensor) -> None:
37
+ _check_ncdhw(x, "x")
38
+ if bias.shape != (x.shape[1],):
39
+ raise RuntimeError("bias must have shape (C,)")
40
+ return None
41
+
42
+
43
+ @torch.library.register_fake(add_op_namespace_prefix("update_cache2_ncdhw_bf16"))
44
+ def _update_cache2_ncdhw_bf16_fake(cur: torch.Tensor, prev: torch.Tensor, out: torch.Tensor) -> None:
45
+ _check_ncdhw(cur, "cur")
46
+ b, c, _, h, w = cur.shape
47
+ if prev.shape != (b, c, 2, h, w) or out.shape != (b, c, 2, h, w):
48
+ raise RuntimeError("prev and out must have shape (B, C, 2, H, W)")
49
+ return None
50
+
51
+
52
+ def ncdhw_to_blc_bf16(x: torch.Tensor, out: torch.Tensor | None = None) -> torch.Tensor:
53
+ """Convert BF16 NCDHW tensor to BLC where L = T * H * W."""
54
+
55
+ if out is None:
56
+ out = torch.empty((x.shape[0], x.shape[2] * x.shape[3] * x.shape[4], x.shape[1]), device=x.device, dtype=x.dtype)
57
+ ops.ncdhw_to_blc_bf16(x, out)
58
+ return out
59
+
60
+
61
+ def time_unshuffle2_bf16(x: torch.Tensor, out: torch.Tensor | None = None) -> torch.Tensor:
62
+ """Convert BF16 (B, 2C, T, H, W) into (B, C, 2T, H, W)."""
63
+
64
+ if out is None:
65
+ out = torch.empty((x.shape[0], x.shape[1] // 2, 2 * x.shape[2], x.shape[3], x.shape[4]), device=x.device, dtype=x.dtype)
66
+ ops.time_unshuffle2_bf16(x, out)
67
+ return out
68
+
69
+
70
+ def add_bias_ncdhw_bf16(x: torch.Tensor, bias: torch.Tensor) -> torch.Tensor:
71
+ """Add channel bias to an NCDHW tensor in place."""
72
+
73
+ ops.add_bias_ncdhw_bf16(x, bias)
74
+ return x
75
+
76
+
77
+ def update_cache2_ncdhw_bf16(cur: torch.Tensor, prev: torch.Tensor, out: torch.Tensor | None = None) -> torch.Tensor:
78
+ """Update two-frame latent cache from current and previous NCDHW tensors."""
79
+
80
+ if out is None:
81
+ out = torch.empty((cur.shape[0], cur.shape[1], 2, cur.shape[3], cur.shape[4]), device=cur.device, dtype=cur.dtype)
82
+ ops.update_cache2_ncdhw_bf16(cur, prev, out)
83
+ return out
84
+
85
+
86
+ __all__ = [
87
+ "ncdhw_to_blc_bf16",
88
+ "time_unshuffle2_bf16",
89
+ "add_bias_ncdhw_bf16",
90
+ "update_cache2_ncdhw_bf16",
91
+ ]
build/torch210-cxx11-cu130-x86_64-linux/_flashrt_spatiotemporal_layout_cuda_cf903dd.abi3.so ADDED
@@ -0,0 +1,3 @@
 
 
 
 
1
+ version https://git-lfs.github.com/spec/v1
2
+ oid sha256:13ae8de167c29b0eff517eff83a2e1b6646c8445b95e074245efef4f978afed6
3
+ size 2528512
build/torch210-cxx11-cu130-x86_64-linux/_ops.py ADDED
@@ -0,0 +1,9 @@
 
 
 
 
 
 
 
 
 
 
1
+ import torch
2
+ from . import _flashrt_spatiotemporal_layout_cuda_cf903dd
3
+ ops = torch.ops._flashrt_spatiotemporal_layout_cuda_cf903dd
4
+
5
+ def add_op_namespace_prefix(op_name: str):
6
+ """
7
+ Prefix op by namespace.
8
+ """
9
+ return f"_flashrt_spatiotemporal_layout_cuda_cf903dd::{op_name}"
build/torch210-cxx11-cu130-x86_64-linux/flashrt_spatiotemporal_layout/__init__.py ADDED
@@ -0,0 +1,26 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ import ctypes
2
+ import importlib.util
3
+ import sys
4
+ from pathlib import Path
5
+ from types import ModuleType
6
+
7
+
8
+ def _import_from_path(file_path: Path) -> ModuleType:
9
+ # We cannot use the module name as-is, after adding it to `sys.modules`,
10
+ # it would also be used for other imports. So, we make a module name that
11
+ # depends on the path for it to be unique using the hex-encoded hash of
12
+ # the path.
13
+ path_hash = "{:x}".format(ctypes.c_size_t(hash(file_path.absolute())).value)
14
+ module_name = path_hash
15
+ spec = importlib.util.spec_from_file_location(module_name, file_path)
16
+ if spec is None:
17
+ raise ImportError(f"Cannot load spec for {module_name} from {file_path}")
18
+ module = importlib.util.module_from_spec(spec)
19
+ if module is None:
20
+ raise ImportError(f"Cannot load module {module_name} from spec")
21
+ sys.modules[module_name] = module
22
+ spec.loader.exec_module(module) # type: ignore
23
+ return module
24
+
25
+
26
+ globals().update(vars(_import_from_path(Path(__file__).parent.parent / "__init__.py")))
build/torch210-cxx11-cu130-x86_64-linux/metadata.json ADDED
@@ -0,0 +1,21 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "name": "flashrt-spatiotemporal-layout",
3
+ "id": "_flashrt_spatiotemporal_layout_cuda_cf903dd",
4
+ "version": 1,
5
+ "license": "Apache-2.0",
6
+ "python-depends": [],
7
+ "backend": {
8
+ "type": "cuda",
9
+ "archs": [
10
+ "10.0",
11
+ "11.0",
12
+ "12.0+PTX",
13
+ "7.5",
14
+ "8.0",
15
+ "8.6",
16
+ "8.7",
17
+ "8.9",
18
+ "9.0"
19
+ ]
20
+ }
21
+ }
build/torch211-cxx11-cu128-x86_64-linux/__init__.py ADDED
@@ -0,0 +1,91 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ """FlashRT spatiotemporal layout kernels."""
2
+
3
+ from __future__ import annotations
4
+
5
+ import torch
6
+
7
+ from ._ops import add_op_namespace_prefix, ops
8
+
9
+
10
+ def _check_ncdhw(x: torch.Tensor, name: str) -> None:
11
+ if x.dim() != 5:
12
+ raise RuntimeError(f"{name} must have shape (B, C, T, H, W)")
13
+
14
+
15
+ @torch.library.register_fake(add_op_namespace_prefix("ncdhw_to_blc_bf16"))
16
+ def _ncdhw_to_blc_bf16_fake(x: torch.Tensor, out: torch.Tensor) -> None:
17
+ _check_ncdhw(x, "x")
18
+ b, c, t, h, w = x.shape
19
+ if out.shape != (b, t * h * w, c):
20
+ raise RuntimeError("out must have shape (B, T * H * W, C)")
21
+ return None
22
+
23
+
24
+ @torch.library.register_fake(add_op_namespace_prefix("time_unshuffle2_bf16"))
25
+ def _time_unshuffle2_bf16_fake(x: torch.Tensor, out: torch.Tensor) -> None:
26
+ _check_ncdhw(x, "x")
27
+ b, c2, t, h, w = x.shape
28
+ if c2 % 2 != 0:
29
+ raise RuntimeError("x.shape[1] must be even")
30
+ if out.shape != (b, c2 // 2, 2 * t, h, w):
31
+ raise RuntimeError("out must have shape (B, C / 2, 2 * T, H, W)")
32
+ return None
33
+
34
+
35
+ @torch.library.register_fake(add_op_namespace_prefix("add_bias_ncdhw_bf16"))
36
+ def _add_bias_ncdhw_bf16_fake(x: torch.Tensor, bias: torch.Tensor) -> None:
37
+ _check_ncdhw(x, "x")
38
+ if bias.shape != (x.shape[1],):
39
+ raise RuntimeError("bias must have shape (C,)")
40
+ return None
41
+
42
+
43
+ @torch.library.register_fake(add_op_namespace_prefix("update_cache2_ncdhw_bf16"))
44
+ def _update_cache2_ncdhw_bf16_fake(cur: torch.Tensor, prev: torch.Tensor, out: torch.Tensor) -> None:
45
+ _check_ncdhw(cur, "cur")
46
+ b, c, _, h, w = cur.shape
47
+ if prev.shape != (b, c, 2, h, w) or out.shape != (b, c, 2, h, w):
48
+ raise RuntimeError("prev and out must have shape (B, C, 2, H, W)")
49
+ return None
50
+
51
+
52
+ def ncdhw_to_blc_bf16(x: torch.Tensor, out: torch.Tensor | None = None) -> torch.Tensor:
53
+ """Convert BF16 NCDHW tensor to BLC where L = T * H * W."""
54
+
55
+ if out is None:
56
+ out = torch.empty((x.shape[0], x.shape[2] * x.shape[3] * x.shape[4], x.shape[1]), device=x.device, dtype=x.dtype)
57
+ ops.ncdhw_to_blc_bf16(x, out)
58
+ return out
59
+
60
+
61
+ def time_unshuffle2_bf16(x: torch.Tensor, out: torch.Tensor | None = None) -> torch.Tensor:
62
+ """Convert BF16 (B, 2C, T, H, W) into (B, C, 2T, H, W)."""
63
+
64
+ if out is None:
65
+ out = torch.empty((x.shape[0], x.shape[1] // 2, 2 * x.shape[2], x.shape[3], x.shape[4]), device=x.device, dtype=x.dtype)
66
+ ops.time_unshuffle2_bf16(x, out)
67
+ return out
68
+
69
+
70
+ def add_bias_ncdhw_bf16(x: torch.Tensor, bias: torch.Tensor) -> torch.Tensor:
71
+ """Add channel bias to an NCDHW tensor in place."""
72
+
73
+ ops.add_bias_ncdhw_bf16(x, bias)
74
+ return x
75
+
76
+
77
+ def update_cache2_ncdhw_bf16(cur: torch.Tensor, prev: torch.Tensor, out: torch.Tensor | None = None) -> torch.Tensor:
78
+ """Update two-frame latent cache from current and previous NCDHW tensors."""
79
+
80
+ if out is None:
81
+ out = torch.empty((cur.shape[0], cur.shape[1], 2, cur.shape[3], cur.shape[4]), device=cur.device, dtype=cur.dtype)
82
+ ops.update_cache2_ncdhw_bf16(cur, prev, out)
83
+ return out
84
+
85
+
86
+ __all__ = [
87
+ "ncdhw_to_blc_bf16",
88
+ "time_unshuffle2_bf16",
89
+ "add_bias_ncdhw_bf16",
90
+ "update_cache2_ncdhw_bf16",
91
+ ]
build/torch211-cxx11-cu128-x86_64-linux/_flashrt_spatiotemporal_layout_cuda_cf903dd.abi3.so ADDED
@@ -0,0 +1,3 @@
 
 
 
 
1
+ version https://git-lfs.github.com/spec/v1
2
+ oid sha256:18a84407f77c3f0d010257989c79bfc295e38c0f7eaa2f39d1b96da2b99cde15
3
+ size 2610792
build/torch211-cxx11-cu128-x86_64-linux/_ops.py ADDED
@@ -0,0 +1,9 @@
 
 
 
 
 
 
 
 
 
 
1
+ import torch
2
+ from . import _flashrt_spatiotemporal_layout_cuda_cf903dd
3
+ ops = torch.ops._flashrt_spatiotemporal_layout_cuda_cf903dd
4
+
5
+ def add_op_namespace_prefix(op_name: str):
6
+ """
7
+ Prefix op by namespace.
8
+ """
9
+ return f"_flashrt_spatiotemporal_layout_cuda_cf903dd::{op_name}"
build/torch211-cxx11-cu128-x86_64-linux/flashrt_spatiotemporal_layout/__init__.py ADDED
@@ -0,0 +1,26 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ import ctypes
2
+ import importlib.util
3
+ import sys
4
+ from pathlib import Path
5
+ from types import ModuleType
6
+
7
+
8
+ def _import_from_path(file_path: Path) -> ModuleType:
9
+ # We cannot use the module name as-is, after adding it to `sys.modules`,
10
+ # it would also be used for other imports. So, we make a module name that
11
+ # depends on the path for it to be unique using the hex-encoded hash of
12
+ # the path.
13
+ path_hash = "{:x}".format(ctypes.c_size_t(hash(file_path.absolute())).value)
14
+ module_name = path_hash
15
+ spec = importlib.util.spec_from_file_location(module_name, file_path)
16
+ if spec is None:
17
+ raise ImportError(f"Cannot load spec for {module_name} from {file_path}")
18
+ module = importlib.util.module_from_spec(spec)
19
+ if module is None:
20
+ raise ImportError(f"Cannot load module {module_name} from spec")
21
+ sys.modules[module_name] = module
22
+ spec.loader.exec_module(module) # type: ignore
23
+ return module
24
+
25
+
26
+ globals().update(vars(_import_from_path(Path(__file__).parent.parent / "__init__.py")))
build/torch211-cxx11-cu128-x86_64-linux/metadata.json ADDED
@@ -0,0 +1,23 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "name": "flashrt-spatiotemporal-layout",
3
+ "id": "_flashrt_spatiotemporal_layout_cuda_cf903dd",
4
+ "version": 1,
5
+ "license": "Apache-2.0",
6
+ "python-depends": [],
7
+ "backend": {
8
+ "type": "cuda",
9
+ "archs": [
10
+ "10.0",
11
+ "10.1",
12
+ "12.0+PTX",
13
+ "7.0",
14
+ "7.2",
15
+ "7.5",
16
+ "8.0",
17
+ "8.6",
18
+ "8.7",
19
+ "8.9",
20
+ "9.0"
21
+ ]
22
+ }
23
+ }
build/torch211-cxx11-cu130-x86_64-linux/__init__.py ADDED
@@ -0,0 +1,91 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ """FlashRT spatiotemporal layout kernels."""
2
+
3
+ from __future__ import annotations
4
+
5
+ import torch
6
+
7
+ from ._ops import add_op_namespace_prefix, ops
8
+
9
+
10
+ def _check_ncdhw(x: torch.Tensor, name: str) -> None:
11
+ if x.dim() != 5:
12
+ raise RuntimeError(f"{name} must have shape (B, C, T, H, W)")
13
+
14
+
15
+ @torch.library.register_fake(add_op_namespace_prefix("ncdhw_to_blc_bf16"))
16
+ def _ncdhw_to_blc_bf16_fake(x: torch.Tensor, out: torch.Tensor) -> None:
17
+ _check_ncdhw(x, "x")
18
+ b, c, t, h, w = x.shape
19
+ if out.shape != (b, t * h * w, c):
20
+ raise RuntimeError("out must have shape (B, T * H * W, C)")
21
+ return None
22
+
23
+
24
+ @torch.library.register_fake(add_op_namespace_prefix("time_unshuffle2_bf16"))
25
+ def _time_unshuffle2_bf16_fake(x: torch.Tensor, out: torch.Tensor) -> None:
26
+ _check_ncdhw(x, "x")
27
+ b, c2, t, h, w = x.shape
28
+ if c2 % 2 != 0:
29
+ raise RuntimeError("x.shape[1] must be even")
30
+ if out.shape != (b, c2 // 2, 2 * t, h, w):
31
+ raise RuntimeError("out must have shape (B, C / 2, 2 * T, H, W)")
32
+ return None
33
+
34
+
35
+ @torch.library.register_fake(add_op_namespace_prefix("add_bias_ncdhw_bf16"))
36
+ def _add_bias_ncdhw_bf16_fake(x: torch.Tensor, bias: torch.Tensor) -> None:
37
+ _check_ncdhw(x, "x")
38
+ if bias.shape != (x.shape[1],):
39
+ raise RuntimeError("bias must have shape (C,)")
40
+ return None
41
+
42
+
43
+ @torch.library.register_fake(add_op_namespace_prefix("update_cache2_ncdhw_bf16"))
44
+ def _update_cache2_ncdhw_bf16_fake(cur: torch.Tensor, prev: torch.Tensor, out: torch.Tensor) -> None:
45
+ _check_ncdhw(cur, "cur")
46
+ b, c, _, h, w = cur.shape
47
+ if prev.shape != (b, c, 2, h, w) or out.shape != (b, c, 2, h, w):
48
+ raise RuntimeError("prev and out must have shape (B, C, 2, H, W)")
49
+ return None
50
+
51
+
52
+ def ncdhw_to_blc_bf16(x: torch.Tensor, out: torch.Tensor | None = None) -> torch.Tensor:
53
+ """Convert BF16 NCDHW tensor to BLC where L = T * H * W."""
54
+
55
+ if out is None:
56
+ out = torch.empty((x.shape[0], x.shape[2] * x.shape[3] * x.shape[4], x.shape[1]), device=x.device, dtype=x.dtype)
57
+ ops.ncdhw_to_blc_bf16(x, out)
58
+ return out
59
+
60
+
61
+ def time_unshuffle2_bf16(x: torch.Tensor, out: torch.Tensor | None = None) -> torch.Tensor:
62
+ """Convert BF16 (B, 2C, T, H, W) into (B, C, 2T, H, W)."""
63
+
64
+ if out is None:
65
+ out = torch.empty((x.shape[0], x.shape[1] // 2, 2 * x.shape[2], x.shape[3], x.shape[4]), device=x.device, dtype=x.dtype)
66
+ ops.time_unshuffle2_bf16(x, out)
67
+ return out
68
+
69
+
70
+ def add_bias_ncdhw_bf16(x: torch.Tensor, bias: torch.Tensor) -> torch.Tensor:
71
+ """Add channel bias to an NCDHW tensor in place."""
72
+
73
+ ops.add_bias_ncdhw_bf16(x, bias)
74
+ return x
75
+
76
+
77
+ def update_cache2_ncdhw_bf16(cur: torch.Tensor, prev: torch.Tensor, out: torch.Tensor | None = None) -> torch.Tensor:
78
+ """Update two-frame latent cache from current and previous NCDHW tensors."""
79
+
80
+ if out is None:
81
+ out = torch.empty((cur.shape[0], cur.shape[1], 2, cur.shape[3], cur.shape[4]), device=cur.device, dtype=cur.dtype)
82
+ ops.update_cache2_ncdhw_bf16(cur, prev, out)
83
+ return out
84
+
85
+
86
+ __all__ = [
87
+ "ncdhw_to_blc_bf16",
88
+ "time_unshuffle2_bf16",
89
+ "add_bias_ncdhw_bf16",
90
+ "update_cache2_ncdhw_bf16",
91
+ ]
build/torch211-cxx11-cu130-x86_64-linux/_flashrt_spatiotemporal_layout_cuda_cf903dd.abi3.so ADDED
@@ -0,0 +1,3 @@
 
 
 
 
1
+ version https://git-lfs.github.com/spec/v1
2
+ oid sha256:b1432b4f402b24b0cf9dda42334059b9ea83ed947cc99e8714fb80a9cfba38e5
3
+ size 2513352
build/torch211-cxx11-cu130-x86_64-linux/_ops.py ADDED
@@ -0,0 +1,9 @@
 
 
 
 
 
 
 
 
 
 
1
+ import torch
2
+ from . import _flashrt_spatiotemporal_layout_cuda_cf903dd
3
+ ops = torch.ops._flashrt_spatiotemporal_layout_cuda_cf903dd
4
+
5
+ def add_op_namespace_prefix(op_name: str):
6
+ """
7
+ Prefix op by namespace.
8
+ """
9
+ return f"_flashrt_spatiotemporal_layout_cuda_cf903dd::{op_name}"
build/torch211-cxx11-cu130-x86_64-linux/flashrt_spatiotemporal_layout/__init__.py ADDED
@@ -0,0 +1,26 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ import ctypes
2
+ import importlib.util
3
+ import sys
4
+ from pathlib import Path
5
+ from types import ModuleType
6
+
7
+
8
+ def _import_from_path(file_path: Path) -> ModuleType:
9
+ # We cannot use the module name as-is, after adding it to `sys.modules`,
10
+ # it would also be used for other imports. So, we make a module name that
11
+ # depends on the path for it to be unique using the hex-encoded hash of
12
+ # the path.
13
+ path_hash = "{:x}".format(ctypes.c_size_t(hash(file_path.absolute())).value)
14
+ module_name = path_hash
15
+ spec = importlib.util.spec_from_file_location(module_name, file_path)
16
+ if spec is None:
17
+ raise ImportError(f"Cannot load spec for {module_name} from {file_path}")
18
+ module = importlib.util.module_from_spec(spec)
19
+ if module is None:
20
+ raise ImportError(f"Cannot load module {module_name} from spec")
21
+ sys.modules[module_name] = module
22
+ spec.loader.exec_module(module) # type: ignore
23
+ return module
24
+
25
+
26
+ globals().update(vars(_import_from_path(Path(__file__).parent.parent / "__init__.py")))
build/torch211-cxx11-cu130-x86_64-linux/metadata.json ADDED
@@ -0,0 +1,21 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "name": "flashrt-spatiotemporal-layout",
3
+ "id": "_flashrt_spatiotemporal_layout_cuda_cf903dd",
4
+ "version": 1,
5
+ "license": "Apache-2.0",
6
+ "python-depends": [],
7
+ "backend": {
8
+ "type": "cuda",
9
+ "archs": [
10
+ "10.0",
11
+ "11.0",
12
+ "12.0+PTX",
13
+ "7.5",
14
+ "8.0",
15
+ "8.6",
16
+ "8.7",
17
+ "8.9",
18
+ "9.0"
19
+ ]
20
+ }
21
+ }
build/torch212-cxx11-cu130-x86_64-linux/__init__.py ADDED
@@ -0,0 +1,91 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ """FlashRT spatiotemporal layout kernels."""
2
+
3
+ from __future__ import annotations
4
+
5
+ import torch
6
+
7
+ from ._ops import add_op_namespace_prefix, ops
8
+
9
+
10
+ def _check_ncdhw(x: torch.Tensor, name: str) -> None:
11
+ if x.dim() != 5:
12
+ raise RuntimeError(f"{name} must have shape (B, C, T, H, W)")
13
+
14
+
15
+ @torch.library.register_fake(add_op_namespace_prefix("ncdhw_to_blc_bf16"))
16
+ def _ncdhw_to_blc_bf16_fake(x: torch.Tensor, out: torch.Tensor) -> None:
17
+ _check_ncdhw(x, "x")
18
+ b, c, t, h, w = x.shape
19
+ if out.shape != (b, t * h * w, c):
20
+ raise RuntimeError("out must have shape (B, T * H * W, C)")
21
+ return None
22
+
23
+
24
+ @torch.library.register_fake(add_op_namespace_prefix("time_unshuffle2_bf16"))
25
+ def _time_unshuffle2_bf16_fake(x: torch.Tensor, out: torch.Tensor) -> None:
26
+ _check_ncdhw(x, "x")
27
+ b, c2, t, h, w = x.shape
28
+ if c2 % 2 != 0:
29
+ raise RuntimeError("x.shape[1] must be even")
30
+ if out.shape != (b, c2 // 2, 2 * t, h, w):
31
+ raise RuntimeError("out must have shape (B, C / 2, 2 * T, H, W)")
32
+ return None
33
+
34
+
35
+ @torch.library.register_fake(add_op_namespace_prefix("add_bias_ncdhw_bf16"))
36
+ def _add_bias_ncdhw_bf16_fake(x: torch.Tensor, bias: torch.Tensor) -> None:
37
+ _check_ncdhw(x, "x")
38
+ if bias.shape != (x.shape[1],):
39
+ raise RuntimeError("bias must have shape (C,)")
40
+ return None
41
+
42
+
43
+ @torch.library.register_fake(add_op_namespace_prefix("update_cache2_ncdhw_bf16"))
44
+ def _update_cache2_ncdhw_bf16_fake(cur: torch.Tensor, prev: torch.Tensor, out: torch.Tensor) -> None:
45
+ _check_ncdhw(cur, "cur")
46
+ b, c, _, h, w = cur.shape
47
+ if prev.shape != (b, c, 2, h, w) or out.shape != (b, c, 2, h, w):
48
+ raise RuntimeError("prev and out must have shape (B, C, 2, H, W)")
49
+ return None
50
+
51
+
52
+ def ncdhw_to_blc_bf16(x: torch.Tensor, out: torch.Tensor | None = None) -> torch.Tensor:
53
+ """Convert BF16 NCDHW tensor to BLC where L = T * H * W."""
54
+
55
+ if out is None:
56
+ out = torch.empty((x.shape[0], x.shape[2] * x.shape[3] * x.shape[4], x.shape[1]), device=x.device, dtype=x.dtype)
57
+ ops.ncdhw_to_blc_bf16(x, out)
58
+ return out
59
+
60
+
61
+ def time_unshuffle2_bf16(x: torch.Tensor, out: torch.Tensor | None = None) -> torch.Tensor:
62
+ """Convert BF16 (B, 2C, T, H, W) into (B, C, 2T, H, W)."""
63
+
64
+ if out is None:
65
+ out = torch.empty((x.shape[0], x.shape[1] // 2, 2 * x.shape[2], x.shape[3], x.shape[4]), device=x.device, dtype=x.dtype)
66
+ ops.time_unshuffle2_bf16(x, out)
67
+ return out
68
+
69
+
70
+ def add_bias_ncdhw_bf16(x: torch.Tensor, bias: torch.Tensor) -> torch.Tensor:
71
+ """Add channel bias to an NCDHW tensor in place."""
72
+
73
+ ops.add_bias_ncdhw_bf16(x, bias)
74
+ return x
75
+
76
+
77
+ def update_cache2_ncdhw_bf16(cur: torch.Tensor, prev: torch.Tensor, out: torch.Tensor | None = None) -> torch.Tensor:
78
+ """Update two-frame latent cache from current and previous NCDHW tensors."""
79
+
80
+ if out is None:
81
+ out = torch.empty((cur.shape[0], cur.shape[1], 2, cur.shape[3], cur.shape[4]), device=cur.device, dtype=cur.dtype)
82
+ ops.update_cache2_ncdhw_bf16(cur, prev, out)
83
+ return out
84
+
85
+
86
+ __all__ = [
87
+ "ncdhw_to_blc_bf16",
88
+ "time_unshuffle2_bf16",
89
+ "add_bias_ncdhw_bf16",
90
+ "update_cache2_ncdhw_bf16",
91
+ ]
build/torch212-cxx11-cu130-x86_64-linux/_flashrt_spatiotemporal_layout_cuda_cf903dd.abi3.so ADDED
@@ -0,0 +1,3 @@
 
 
 
 
1
+ version https://git-lfs.github.com/spec/v1
2
+ oid sha256:3c6dbbf811af69ff253978b96b02d4e9cabec4f20f2ad5a887aca31cf5267d14
3
+ size 2518880
build/torch212-cxx11-cu130-x86_64-linux/_ops.py ADDED
@@ -0,0 +1,9 @@
 
 
 
 
 
 
 
 
 
 
1
+ import torch
2
+ from . import _flashrt_spatiotemporal_layout_cuda_cf903dd
3
+ ops = torch.ops._flashrt_spatiotemporal_layout_cuda_cf903dd
4
+
5
+ def add_op_namespace_prefix(op_name: str):
6
+ """
7
+ Prefix op by namespace.
8
+ """
9
+ return f"_flashrt_spatiotemporal_layout_cuda_cf903dd::{op_name}"
build/torch212-cxx11-cu130-x86_64-linux/flashrt_spatiotemporal_layout/__init__.py ADDED
@@ -0,0 +1,26 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ import ctypes
2
+ import importlib.util
3
+ import sys
4
+ from pathlib import Path
5
+ from types import ModuleType
6
+
7
+
8
+ def _import_from_path(file_path: Path) -> ModuleType:
9
+ # We cannot use the module name as-is, after adding it to `sys.modules`,
10
+ # it would also be used for other imports. So, we make a module name that
11
+ # depends on the path for it to be unique using the hex-encoded hash of
12
+ # the path.
13
+ path_hash = "{:x}".format(ctypes.c_size_t(hash(file_path.absolute())).value)
14
+ module_name = path_hash
15
+ spec = importlib.util.spec_from_file_location(module_name, file_path)
16
+ if spec is None:
17
+ raise ImportError(f"Cannot load spec for {module_name} from {file_path}")
18
+ module = importlib.util.module_from_spec(spec)
19
+ if module is None:
20
+ raise ImportError(f"Cannot load module {module_name} from spec")
21
+ sys.modules[module_name] = module
22
+ spec.loader.exec_module(module) # type: ignore
23
+ return module
24
+
25
+
26
+ globals().update(vars(_import_from_path(Path(__file__).parent.parent / "__init__.py")))
build/torch212-cxx11-cu130-x86_64-linux/metadata.json ADDED
@@ -0,0 +1,21 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "name": "flashrt-spatiotemporal-layout",
3
+ "id": "_flashrt_spatiotemporal_layout_cuda_cf903dd",
4
+ "version": 1,
5
+ "license": "Apache-2.0",
6
+ "python-depends": [],
7
+ "backend": {
8
+ "type": "cuda",
9
+ "archs": [
10
+ "10.0",
11
+ "11.0",
12
+ "12.0+PTX",
13
+ "7.5",
14
+ "8.0",
15
+ "8.6",
16
+ "8.7",
17
+ "8.9",
18
+ "9.0"
19
+ ]
20
+ }
21
+ }
build/torch212-cxx11-cu132-x86_64-linux/__init__.py ADDED
@@ -0,0 +1,91 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ """FlashRT spatiotemporal layout kernels."""
2
+
3
+ from __future__ import annotations
4
+
5
+ import torch
6
+
7
+ from ._ops import add_op_namespace_prefix, ops
8
+
9
+
10
+ def _check_ncdhw(x: torch.Tensor, name: str) -> None:
11
+ if x.dim() != 5:
12
+ raise RuntimeError(f"{name} must have shape (B, C, T, H, W)")
13
+
14
+
15
+ @torch.library.register_fake(add_op_namespace_prefix("ncdhw_to_blc_bf16"))
16
+ def _ncdhw_to_blc_bf16_fake(x: torch.Tensor, out: torch.Tensor) -> None:
17
+ _check_ncdhw(x, "x")
18
+ b, c, t, h, w = x.shape
19
+ if out.shape != (b, t * h * w, c):
20
+ raise RuntimeError("out must have shape (B, T * H * W, C)")
21
+ return None
22
+
23
+
24
+ @torch.library.register_fake(add_op_namespace_prefix("time_unshuffle2_bf16"))
25
+ def _time_unshuffle2_bf16_fake(x: torch.Tensor, out: torch.Tensor) -> None:
26
+ _check_ncdhw(x, "x")
27
+ b, c2, t, h, w = x.shape
28
+ if c2 % 2 != 0:
29
+ raise RuntimeError("x.shape[1] must be even")
30
+ if out.shape != (b, c2 // 2, 2 * t, h, w):
31
+ raise RuntimeError("out must have shape (B, C / 2, 2 * T, H, W)")
32
+ return None
33
+
34
+
35
+ @torch.library.register_fake(add_op_namespace_prefix("add_bias_ncdhw_bf16"))
36
+ def _add_bias_ncdhw_bf16_fake(x: torch.Tensor, bias: torch.Tensor) -> None:
37
+ _check_ncdhw(x, "x")
38
+ if bias.shape != (x.shape[1],):
39
+ raise RuntimeError("bias must have shape (C,)")
40
+ return None
41
+
42
+
43
+ @torch.library.register_fake(add_op_namespace_prefix("update_cache2_ncdhw_bf16"))
44
+ def _update_cache2_ncdhw_bf16_fake(cur: torch.Tensor, prev: torch.Tensor, out: torch.Tensor) -> None:
45
+ _check_ncdhw(cur, "cur")
46
+ b, c, _, h, w = cur.shape
47
+ if prev.shape != (b, c, 2, h, w) or out.shape != (b, c, 2, h, w):
48
+ raise RuntimeError("prev and out must have shape (B, C, 2, H, W)")
49
+ return None
50
+
51
+
52
+ def ncdhw_to_blc_bf16(x: torch.Tensor, out: torch.Tensor | None = None) -> torch.Tensor:
53
+ """Convert BF16 NCDHW tensor to BLC where L = T * H * W."""
54
+
55
+ if out is None:
56
+ out = torch.empty((x.shape[0], x.shape[2] * x.shape[3] * x.shape[4], x.shape[1]), device=x.device, dtype=x.dtype)
57
+ ops.ncdhw_to_blc_bf16(x, out)
58
+ return out
59
+
60
+
61
+ def time_unshuffle2_bf16(x: torch.Tensor, out: torch.Tensor | None = None) -> torch.Tensor:
62
+ """Convert BF16 (B, 2C, T, H, W) into (B, C, 2T, H, W)."""
63
+
64
+ if out is None:
65
+ out = torch.empty((x.shape[0], x.shape[1] // 2, 2 * x.shape[2], x.shape[3], x.shape[4]), device=x.device, dtype=x.dtype)
66
+ ops.time_unshuffle2_bf16(x, out)
67
+ return out
68
+
69
+
70
+ def add_bias_ncdhw_bf16(x: torch.Tensor, bias: torch.Tensor) -> torch.Tensor:
71
+ """Add channel bias to an NCDHW tensor in place."""
72
+
73
+ ops.add_bias_ncdhw_bf16(x, bias)
74
+ return x
75
+
76
+
77
+ def update_cache2_ncdhw_bf16(cur: torch.Tensor, prev: torch.Tensor, out: torch.Tensor | None = None) -> torch.Tensor:
78
+ """Update two-frame latent cache from current and previous NCDHW tensors."""
79
+
80
+ if out is None:
81
+ out = torch.empty((cur.shape[0], cur.shape[1], 2, cur.shape[3], cur.shape[4]), device=cur.device, dtype=cur.dtype)
82
+ ops.update_cache2_ncdhw_bf16(cur, prev, out)
83
+ return out
84
+
85
+
86
+ __all__ = [
87
+ "ncdhw_to_blc_bf16",
88
+ "time_unshuffle2_bf16",
89
+ "add_bias_ncdhw_bf16",
90
+ "update_cache2_ncdhw_bf16",
91
+ ]
build/torch212-cxx11-cu132-x86_64-linux/_flashrt_spatiotemporal_layout_cuda_cf903dd.abi3.so ADDED
@@ -0,0 +1,3 @@
 
 
 
 
1
+ version https://git-lfs.github.com/spec/v1
2
+ oid sha256:3ff8ee1e1c690aff7cead2ecad593054e5405abd597b9a00b77a4526a7a5995f
3
+ size 2522944
build/torch212-cxx11-cu132-x86_64-linux/_ops.py ADDED
@@ -0,0 +1,9 @@
 
 
 
 
 
 
 
 
 
 
1
+ import torch
2
+ from . import _flashrt_spatiotemporal_layout_cuda_cf903dd
3
+ ops = torch.ops._flashrt_spatiotemporal_layout_cuda_cf903dd
4
+
5
+ def add_op_namespace_prefix(op_name: str):
6
+ """
7
+ Prefix op by namespace.
8
+ """
9
+ return f"_flashrt_spatiotemporal_layout_cuda_cf903dd::{op_name}"
build/torch212-cxx11-cu132-x86_64-linux/flashrt_spatiotemporal_layout/__init__.py ADDED
@@ -0,0 +1,26 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ import ctypes
2
+ import importlib.util
3
+ import sys
4
+ from pathlib import Path
5
+ from types import ModuleType
6
+
7
+
8
+ def _import_from_path(file_path: Path) -> ModuleType:
9
+ # We cannot use the module name as-is, after adding it to `sys.modules`,
10
+ # it would also be used for other imports. So, we make a module name that
11
+ # depends on the path for it to be unique using the hex-encoded hash of
12
+ # the path.
13
+ path_hash = "{:x}".format(ctypes.c_size_t(hash(file_path.absolute())).value)
14
+ module_name = path_hash
15
+ spec = importlib.util.spec_from_file_location(module_name, file_path)
16
+ if spec is None:
17
+ raise ImportError(f"Cannot load spec for {module_name} from {file_path}")
18
+ module = importlib.util.module_from_spec(spec)
19
+ if module is None:
20
+ raise ImportError(f"Cannot load module {module_name} from spec")
21
+ sys.modules[module_name] = module
22
+ spec.loader.exec_module(module) # type: ignore
23
+ return module
24
+
25
+
26
+ globals().update(vars(_import_from_path(Path(__file__).parent.parent / "__init__.py")))
build/torch212-cxx11-cu132-x86_64-linux/metadata.json ADDED
@@ -0,0 +1,21 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "name": "flashrt-spatiotemporal-layout",
3
+ "id": "_flashrt_spatiotemporal_layout_cuda_cf903dd",
4
+ "version": 1,
5
+ "license": "Apache-2.0",
6
+ "python-depends": [],
7
+ "backend": {
8
+ "type": "cuda",
9
+ "archs": [
10
+ "10.0",
11
+ "11.0",
12
+ "12.0+PTX",
13
+ "7.5",
14
+ "8.0",
15
+ "8.6",
16
+ "8.7",
17
+ "8.9",
18
+ "9.0"
19
+ ]
20
+ }
21
+ }