HF Space: run cuobjdump/nvdisasm on uploaded cubins

.gitattributes

@@ -0,0 +1 @@
+*.cubin filter=lfs diff=lfs merge=lfs -text

.gitignore

@@ -0,0 +1 @@
+__pycache__/

README.md

@@ -0,0 +1,16 @@
+---
+title: cuobjdump / nvdisasm
+emoji: 🔍
+colorFrom: green
+colorTo: gray
+sdk: gradio
+sdk_version: 6.29.1
+app_file: app.py
+pinned: false
+short_description: Disassemble CUDA cubins with cuobjdump and nvdisasm
+---
+
+Upload a CUDA cubin, fatbin or ELF object and run NVIDIA's `cuobjdump` or `nvdisasm` on it.
+Both tools are installed from PyPI (`nvidia-cuda-cuobjdump`, `nvidia-cuda-nvdisasm`); no GPU required.
+
+`examples/` holds two cubins assembled with `ptxas` from the included PTX sources.

app.py

@@ -0,0 +1,56 @@
+import shlex
+import subprocess
+from pathlib import Path
+
+import gradio as gr
+import nvidia
+
+
+def find_tool(name):
+    path = next(p for d in nvidia.__path__ if (p := Path(d) / "cu13/bin" / name).exists())
+    path.chmod(0o755)
+    return str(path)
+
+
+TOOLS = {name: find_tool(name) for name in ("cuobjdump", "nvdisasm")}
+FLAGS = {
+    "cuobjdump": ["-sass", "-ptx", "-elf", "-res-usage", "-symbols", "-lelf", "-lptx"],
+    "nvdisasm": ["-c", "-g", "-gi", "-hex", "-print-life-ranges", "-print-raw", "-print-line-info", "-cfg", "-bbcfg"],
+}
+DEFAULT_FLAGS = {"cuobjdump": ["-sass"], "nvdisasm": ["-c"]}
+EXAMPLES = sorted(str(p) for p in Path(__file__).with_name("examples").glob("*.cubin"))
+
+
+def version(tool):
+    return subprocess.run([TOOLS[tool], "--version"], capture_output=True, text=True).stdout.split("release ")[-1].splitlines()[0]
+
+
+def run(file, tool, flags, extra):
+    if file is None:
+        return "Upload a cubin first."
+    try:
+        r = subprocess.run([TOOLS[tool], *flags, *shlex.split(extra or ""), file],
+                           capture_output=True, text=True, timeout=60)
+    except subprocess.TimeoutExpired:
+        return "Timed out after 60s."
+    return r.stdout + r.stderr + (f"\n[exit code {r.returncode}]" if r.returncode else "")
+
+
+with gr.Blocks(title="cuobjdump / nvdisasm") as demo:
+    gr.Markdown("# cuobjdump / nvdisasm\n"
+                "Disassemble CUDA cubins, fatbins and ELF objects with NVIDIA's tools from PyPI.\n\n"
+                + "\n".join(f"- `{t}`: {version(t)}" for t in TOOLS))
+    with gr.Row():
+        with gr.Column():
+            file = gr.File(label="cubin / fatbin / CUDA ELF", type="filepath")
+            tool = gr.Radio(list(TOOLS), value="cuobjdump", label="Tool")
+            flags = gr.CheckboxGroup(FLAGS["cuobjdump"], value=DEFAULT_FLAGS["cuobjdump"], label="Flags")
+            extra = gr.Textbox(label="Extra arguments", placeholder="-fun vecadd")
+            btn = gr.Button("Run", variant="primary")
+        out = gr.Code(label="Output", lines=30, max_lines=60)
+    gr.Examples([[e, t] for e in EXAMPLES for t in TOOLS], [file, tool])
+    tool.change(lambda t: gr.CheckboxGroup(FLAGS[t], value=DEFAULT_FLAGS[t]), tool, flags)
+    btn.click(run, [file, tool, flags, extra], out)
+
+if __name__ == "__main__":
+    demo.launch()

examples/saxpy.ptx

@@ -0,0 +1,38 @@
+.version 8.0
+.target sm_90
+.address_size 64
+
+.visible .entry saxpy(
+	.param .f32 alpha,
+	.param .u64 x,
+	.param .u64 y,
+	.param .u32 n
+)
+{
+	.reg .pred %p;
+	.reg .b32 %r<6>;
+	.reg .f32 %f<5>;
+	.reg .b64 %rd<8>;
+
+	ld.param.f32 %f1, [alpha];
+	ld.param.u64 %rd1, [x];
+	ld.param.u64 %rd2, [y];
+	ld.param.u32 %r1, [n];
+	mov.u32 %r2, %ctaid.x;
+	mov.u32 %r3, %ntid.x;
+	mov.u32 %r4, %tid.x;
+	mad.lo.s32 %r5, %r2, %r3, %r4;
+	setp.ge.s32 %p, %r5, %r1;
+	@%p bra DONE;
+	cvta.to.global.u64 %rd3, %rd1;
+	cvta.to.global.u64 %rd4, %rd2;
+	mul.wide.s32 %rd5, %r5, 4;
+	add.s64 %rd6, %rd3, %rd5;
+	add.s64 %rd7, %rd4, %rd5;
+	ld.global.f32 %f2, [%rd6];
+	ld.global.f32 %f3, [%rd7];
+	fma.rn.f32 %f4, %f1, %f2, %f3;
+	st.global.f32 [%rd7], %f4;
+DONE:
+	ret;
+}

examples/saxpy_sm90.cubin

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:9020ab700f01f18e53ab948267b26a34b86ea3ac5952879e3c840bb703898ebe
+size 3848

examples/vecadd.ptx

@@ -0,0 +1,40 @@
+.version 8.0
+.target sm_80
+.address_size 64
+
+.visible .entry vecadd(
+	.param .u64 a,
+	.param .u64 b,
+	.param .u64 c,
+	.param .u32 n
+)
+{
+	.reg .pred %p;
+	.reg .b32 %r<6>;
+	.reg .f32 %f<4>;
+	.reg .b64 %rd<10>;
+
+	ld.param.u64 %rd1, [a];
+	ld.param.u64 %rd2, [b];
+	ld.param.u64 %rd3, [c];
+	ld.param.u32 %r1, [n];
+	mov.u32 %r2, %ctaid.x;
+	mov.u32 %r3, %ntid.x;
+	mov.u32 %r4, %tid.x;
+	mad.lo.s32 %r5, %r2, %r3, %r4;
+	setp.ge.s32 %p, %r5, %r1;
+	@%p bra DONE;
+	cvta.to.global.u64 %rd4, %rd1;
+	cvta.to.global.u64 %rd5, %rd2;
+	cvta.to.global.u64 %rd6, %rd3;
+	mul.wide.s32 %rd7, %r5, 4;
+	add.s64 %rd8, %rd4, %rd7;
+	ld.global.f32 %f1, [%rd8];
+	add.s64 %rd8, %rd5, %rd7;
+	ld.global.f32 %f2, [%rd8];
+	add.f32 %f3, %f1, %f2;
+	add.s64 %rd9, %rd6, %rd7;
+	st.global.f32 [%rd9], %f3;
+DONE:
+	ret;
+}

examples/vecadd_sm80.cubin

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:f1d8c85fb66f436d80ed32bfbd16809226361377520f64ea675dac791622f0f7
+size 3368

requirements.txt

@@ -0,0 +1,2 @@
+nvidia-cuda-cuobjdump
+nvidia-cuda-nvdisasm