Skip to content

Commit d2d5ae9

Browse files
Update docs
1 parent 28b8f37 commit d2d5ae9

File tree

7 files changed

+416
-26
lines changed

7 files changed

+416
-26
lines changed

_sources/autoapi/tilelang/contrib/nvcc/index.rst.txt

Lines changed: 49 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -15,6 +15,9 @@ Functions
1515
.. autoapisummary::
1616

1717
tilelang.contrib.nvcc.compile_cuda
18+
tilelang.contrib.nvcc.default_compile_options
19+
tilelang.contrib.nvcc.get_ptx_from_source
20+
tilelang.contrib.nvcc.get_sass_from_source
1821
tilelang.contrib.nvcc.find_cuda_path
1922
tilelang.contrib.nvcc.get_cuda_version
2023
tilelang.contrib.nvcc.tilelang_callback_cuda_compile
@@ -56,6 +59,52 @@ Module Contents
5659
:rtype: bytearray
5760

5861

62+
.. py:function:: default_compile_options(compile_flags = None)
63+
64+
Build a set of default NVCC compile options for TileLang generated sources.
65+
66+
Includes C++ standard and common include paths (TileLang templates, CUTLASS,
67+
CUDA include). Merges user-provided compile flags if given.
68+
69+
:param compile_flags: Additional flags to include. Items are split on whitespace.
70+
:type compile_flags: Optional[List[str]]
71+
72+
:returns: A list of flags suitable for NVCC's command line.
73+
:rtype: List[str]
74+
75+
76+
.. py:function:: get_ptx_from_source(code, compile_flags = None, verbose = False)
77+
78+
Compile CUDA C++ source to PTX using NVCC and return as text.
79+
80+
:param code: CUDA C++ kernel source code.
81+
:type code: str
82+
:param compile_flags: Additional flags merged with defaults.
83+
:type compile_flags: Optional[List[str]]
84+
:param verbose: Print NVCC output when True.
85+
:type verbose: bool
86+
87+
:returns: PTX text.
88+
:rtype: str
89+
90+
91+
.. py:function:: get_sass_from_source(code, compile_flags = None, verbose = False)
92+
93+
Compile CUDA C++ source to CUBIN and disassemble to SASS.
94+
95+
Uses nvdisasm if available; otherwise falls back to cuobjdump.
96+
97+
:param code: CUDA C++ kernel source code.
98+
:type code: str
99+
:param compile_flags: Additional flags merged with defaults.
100+
:type compile_flags: Optional[List[str]]
101+
:param verbose: Print tool outputs when True.
102+
:type verbose: bool
103+
104+
:returns: SASS text.
105+
:rtype: str
106+
107+
59108
.. py:function:: find_cuda_path()
60109
61110
Utility function to find cuda path

_sources/autoapi/tilelang/jit/kernel/index.rst.txt

Lines changed: 101 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -191,6 +191,61 @@ Module Contents
191191
.. py:method:: run_once(func = None)
192192
193193
194+
.. py:method:: show_source(which = 'kernel')
195+
196+
Print generated source code to stdout.
197+
198+
:param which: Select which source to print. Defaults to "kernel".
199+
:type which: Literal["kernel", "host", "both"], optional
200+
201+
.. rubric:: Examples
202+
203+
>>> jit_kernel.show_source() # print kernel source
204+
>>> jit_kernel.show_source("host") # print host source
205+
>>> jit_kernel.show_source("both") # print both sources
206+
207+
208+
209+
.. py:method:: export_sources(kernel_path = None, host_path = None)
210+
211+
Export generated source code to files.
212+
213+
:param kernel_path: Destination file path to write the kernel source. If None, skips writing kernel code.
214+
:type kernel_path: Optional[str]
215+
:param host_path: Destination file path to write the host source. If None, skips writing host code.
216+
:type host_path: Optional[str]
217+
218+
.. rubric:: Examples
219+
220+
>>> jit_kernel.export_sources(kernel_path="/tmp/kernel.cu")
221+
>>> jit_kernel.export_sources(host_path="/tmp/host.cc")
222+
>>> jit_kernel.export_sources(
223+
... kernel_path="/tmp/kernel.cu",
224+
... host_path="/tmp/host.cc",
225+
... )
226+
227+
228+
229+
.. py:method:: print_source_code(which = 'kernel', file = None)
230+
231+
Deprecated: use show_source() or export_sources() instead.
232+
233+
:param which: Kept for backward compatibility with printing behavior.
234+
:type which: Literal["kernel", "host", "both"], optional
235+
:param file: If provided, behaves like export_sources(kernel_path=file).
236+
:type file: Optional[str]
237+
238+
.. rubric:: Examples
239+
240+
>>> # New API (preferred)
241+
>>> jit_kernel.show_source("both")
242+
>>> jit_kernel.export_sources(kernel_path="/tmp/kernel.cu")
243+
244+
>>> # Old API (still works but deprecated)
245+
>>> jit_kernel.print_source_code(file="/tmp/kernel.cu")
246+
247+
248+
194249
.. py:method:: update_tuner_result(latency, config, ref_latency)
195250
196251
Updates the tuning results for this kernel.
@@ -247,3 +302,49 @@ Module Contents
247302

248303

249304

305+
.. py:method:: show_ptx()
306+
307+
Print compiled PTX for the kernel (CUDA only).
308+
309+
.. rubric:: Examples
310+
311+
>>> jit_kernel.show_ptx()
312+
313+
314+
315+
.. py:method:: export_ptx(path)
316+
317+
Export compiled PTX to a file (CUDA only).
318+
319+
:param path: Destination file path to write PTX.
320+
:type path: str
321+
322+
.. rubric:: Examples
323+
324+
>>> jit_kernel.export_ptx("/tmp/kernel.ptx")
325+
326+
327+
328+
.. py:method:: show_sass()
329+
330+
Print disassembled SASS for the kernel (CUDA only).
331+
332+
.. rubric:: Examples
333+
334+
>>> jit_kernel.show_sass()
335+
336+
337+
338+
.. py:method:: export_sass(path)
339+
340+
Export disassembled SASS to a file (CUDA only).
341+
342+
:param path: Destination file path to write SASS.
343+
:type path: str
344+
345+
.. rubric:: Examples
346+
347+
>>> jit_kernel.export_sass("/tmp/kernel.sass")
348+
349+
350+

0 commit comments

Comments
 (0)