Skip to content

Commit e171ef7

Browse files
Update docs
1 parent 9fc6abf commit e171ef7

File tree

195 files changed

+1750
-508
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

195 files changed

+1750
-508
lines changed

_sources/autoapi/tilelang/autotuner/param/index.rst.txt

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -73,7 +73,7 @@ Module Contents
7373

7474
.. attribute:: execution_backend
7575

76-
Execution backend to use for kernel execution (default: "cython").
76+
Execution backend to use for kernel execution (default: "auto").
7777

7878
.. attribute:: target
7979

@@ -103,8 +103,8 @@ Module Contents
103103

104104

105105
.. py:attribute:: execution_backend
106-
:type: Literal['dlpack', 'ctypes', 'cython']
107-
:value: 'cython'
106+
:type: Literal['auto', 'tvm_ffi', 'ctypes', 'cython', 'nvrtc', 'torch']
107+
:value: 'auto'
108108

109109

110110

_sources/autoapi/tilelang/autotuner/tuner/index.rst.txt

Lines changed: 4 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -132,7 +132,7 @@ Module Contents
132132

133133

134134

135-
.. py:method:: set_compile_args(out_idx = None, target = 'auto', execution_backend = 'cython', target_host = None, verbose = False, pass_configs = None)
135+
.. py:method:: set_compile_args(out_idx = None, target = 'auto', execution_backend = 'auto', target_host = None, verbose = False, pass_configs = None)
136136
137137
Set compilation arguments for the auto-tuner.
138138

@@ -349,8 +349,9 @@ Module Contents
349349
:type target: Union[str, Target], optional
350350
:param target_host: Target host for cross-compilation. Defaults to None.
351351
:type target_host: Union[str, Target], optional
352-
:param execution_backend: Backend for kernel execution and argument passing. Defaults to "cython".
353-
:type execution_backend: Literal["dlpack", "ctypes", "cython"], optional
352+
:param execution_backend: Backend for kernel execution and argument passing. Use "auto" to pick a sensible
353+
default per target (cuda->tvm_ffi, metal->torch, others->cython).
354+
:type execution_backend: Literal["auto", "tvm_ffi", "ctypes", "cython", "nvrtc", "torch"], optional
354355
:param verbose: Enables verbose logging during compilation. Defaults to False.
355356
:type verbose: bool, optional
356357
:param pass_configs: Configurations for TVM's pass context. Defaults to None.

_sources/autoapi/tilelang/cache/index.rst.txt

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -30,7 +30,7 @@ Functions
3030
Package Contents
3131
----------------
3232

33-
.. py:function:: cached(func = None, out_idx = None, *args, target = 'auto', target_host = None, execution_backend = 'cython', verbose = False, pass_configs = None, compile_flags = None)
33+
.. py:function:: cached(func = None, out_idx = None, *args, target = 'auto', target_host = None, execution_backend = 'auto', verbose = False, pass_configs = None, compile_flags = None)
3434
3535
Caches and reuses compiled kernels (using KernelCache class).
3636

_sources/autoapi/tilelang/cache/kernel_cache/index.rst.txt

Lines changed: 14 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -14,8 +14,9 @@ Attributes
1414

1515
.. autoapisummary::
1616

17-
tilelang.cache.kernel_cache.KERNEL_PATH
18-
tilelang.cache.kernel_cache.WRAPPED_KERNEL_PATH
17+
tilelang.cache.kernel_cache.DEVICE_KERNEL_PATH
18+
tilelang.cache.kernel_cache.HOST_KERNEL_PATH
19+
tilelang.cache.kernel_cache.EXECUTABLE_PATH
1920
tilelang.cache.kernel_cache.KERNEL_LIB_PATH
2021
tilelang.cache.kernel_cache.KERNEL_CUBIN_PATH
2122
tilelang.cache.kernel_cache.KERNEL_PY_PATH
@@ -33,12 +34,16 @@ Classes
3334
Module Contents
3435
---------------
3536

36-
.. py:data:: KERNEL_PATH
37-
:value: 'kernel.cu'
37+
.. py:data:: DEVICE_KERNEL_PATH
38+
:value: 'device_kernel.cu'
3839

3940

40-
.. py:data:: WRAPPED_KERNEL_PATH
41-
:value: 'wrapped_kernel.cu'
41+
.. py:data:: HOST_KERNEL_PATH
42+
:value: 'host_kernel.cu'
43+
44+
45+
.. py:data:: EXECUTABLE_PATH
46+
:value: 'executable.so'
4247

4348

4449
.. py:data:: KERNEL_LIB_PATH
@@ -68,12 +73,12 @@ Module Contents
6873

6974

7075
.. py:attribute:: execution_backend
71-
:type: Literal['dlpack', 'ctypes', 'cython', 'nvrtc']
72-
:value: 'cython'
76+
:type: Literal['tvm_ffi', 'ctypes', 'cython', 'nvrtc', 'torch']
77+
:value: 'tvm_ffi'
7378

7479

7580

76-
.. py:method:: cached(func = None, out_idx = None, *args, target = 'auto', target_host = None, execution_backend = 'cython', verbose = False, pass_configs = None, compile_flags = None)
81+
.. py:method:: cached(func = None, out_idx = None, *args, target = 'auto', target_host = None, execution_backend = 'auto', verbose = False, pass_configs = None, compile_flags = None)
7782
7883
Caches and reuses compiled kernels to avoid redundant compilation.
7984

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

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

1717
tilelang.contrib.dlpack.convert_func
18-
tilelang.contrib.dlpack.to_pytorch_func
1918

2019

2120
Module Contents
@@ -34,14 +33,3 @@ Module Contents
3433
:type to_dlpack_func: Function
3534

3635

37-
.. py:function:: to_pytorch_func(tvm_func)
38-
39-
Convert a tvm function into one that accepts PyTorch tensors
40-
41-
:param tvm_func: Built tvm function operating on arrays
42-
:type tvm_func: Function
43-
44-
:returns: **wrapped_func** -- Wrapped tvm function that operates on PyTorch tensors
45-
:rtype: Function
46-
47-

_sources/autoapi/tilelang/jit/adapter/base/index.rst.txt

Lines changed: 26 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -44,9 +44,34 @@ Module Contents
4444
.. py:attribute:: result_idx
4545
4646
47+
.. py:method:: get_current_stream_functor()
48+
:staticmethod:
49+
50+
51+
Return a callable that reads Torch's current CUDA stream pointer.
52+
53+
The returned lambda yields the raw CUDA stream handle of the current
54+
PyTorch stream on the active device. It's a thunk (evaluated at call
55+
time) so that any upstream stream guards are respected. If CUDA is
56+
unavailable, it returns a lambda that yields 0.
57+
58+
59+
60+
.. py:method:: get_current_device_functor()
61+
:staticmethod:
62+
63+
64+
Return a callable that yields Torch's current device.
65+
66+
Similar to the stream functor, we capture a callable that, when called,
67+
fetches the current device according to PyTorch. On CPU or when CUDA is
68+
unavailable, returns ``torch.device('cpu')``.
69+
70+
71+
4772
.. py:method:: __call__(*args, **kwds)
4873
4974
50-
.. py:method:: get_kernel_source()
75+
.. py:method:: get_kernel_source(kernel_only = True)
5176
5277

_sources/autoapi/tilelang/jit/adapter/ctypes/adapter/index.rst.txt

Lines changed: 10 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -20,7 +20,7 @@ Classes
2020
Module Contents
2121
---------------
2222

23-
.. py:class:: CtypesKernelAdapter(params, result_idx, target, func_or_mod, host_mod = None, device_mod = None, kernel_global_source = None, verbose = False, pass_configs = None, compile_flags = None)
23+
.. py:class:: CtypesKernelAdapter(params, result_idx, target, func_or_mod, host_mod = None, device_mod = None, host_kernel_source = None, device_kernel_source = None, verbose = False, pass_configs = None, compile_flags = None)
2424
2525
Bases: :py:obj:`tilelang.jit.adapter.base.BaseKernelAdapter`
2626

@@ -44,20 +44,20 @@ Module Contents
4444

4545

4646

47-
.. py:attribute:: kernel_global_source
47+
.. py:attribute:: host_kernel_source
4848
:type: str | None
4949
:value: None
5050

5151

5252

53-
.. py:attribute:: lib
54-
:type: ctypes.CDLL | None
53+
.. py:attribute:: device_kernel_source
54+
:type: str | None
5555
:value: None
5656

5757

5858

59-
.. py:attribute:: wrapped_source
60-
:type: str | None
59+
.. py:attribute:: lib
60+
:type: ctypes.CDLL | None
6161
:value: None
6262

6363

@@ -103,7 +103,10 @@ Module Contents
103103
.. py:attribute:: lib_generator
104104
105105
106-
.. py:method:: from_database(params, result_idx, target, func_or_mod, kernel_global_source, kernel_lib_path, verbose = False, pass_configs = None, compile_flags = None)
106+
.. py:attribute:: wrapped_source
107+
108+
109+
.. py:method:: from_database(params, result_idx, target, func_or_mod, host_kernel_source, device_kernel_source, kernel_lib_path, verbose = False, pass_configs = None, compile_flags = None)
107110
:classmethod:
108111

109112

_sources/autoapi/tilelang/jit/adapter/cython/adapter/index.rst.txt

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -44,7 +44,7 @@ Module Contents
4444
A symbolic expression can be a simple tvm.Var, or an tvm.PrimExpr containing tvm.Var.
4545

4646

47-
.. py:class:: CythonKernelAdapter(params, result_idx, target, func_or_mod, host_mod = None, device_mod = None, kernel_global_source = None, verbose = False, pass_configs = None, compile_flags = None)
47+
.. py:class:: CythonKernelAdapter(params, result_idx, target, func_or_mod, host_mod = None, device_mod = None, device_kernel_source = None, verbose = False, pass_configs = None, compile_flags = None)
4848
4949
Bases: :py:obj:`tilelang.jit.adapter.base.BaseKernelAdapter`
5050

@@ -69,20 +69,20 @@ Module Contents
6969

7070

7171

72-
.. py:attribute:: kernel_global_source
72+
.. py:attribute:: host_kernel_source
7373
:type: str | None
7474
:value: None
7575

7676

7777

78-
.. py:attribute:: lib
79-
:type: ctypes.CDLL | None
78+
.. py:attribute:: device_kernel_source
79+
:type: str | None
8080
:value: None
8181

8282

8383

84-
.. py:attribute:: wrapped_source
85-
:type: str | None
84+
.. py:attribute:: lib
85+
:type: ctypes.CDLL | None
8686
:value: None
8787

8888

@@ -155,7 +155,7 @@ Module Contents
155155
.. py:attribute:: cython_wrapper
156156
157157
158-
.. py:method:: from_database(params, result_idx, target, func_or_mod, kernel_global_source, kernel_lib_path, verbose = False, pass_configs = None, compile_flags = None)
158+
.. py:method:: from_database(params, result_idx, target, func_or_mod, host_kernel_source, device_kernel_source, kernel_lib_path, verbose = False, pass_configs = None, compile_flags = None)
159159
:classmethod:
160160

161161

_sources/autoapi/tilelang/jit/adapter/dlpack/index.rst.txt

Lines changed: 0 additions & 31 deletions
This file was deleted.

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -13,10 +13,10 @@ Submodules
1313
/autoapi/tilelang/jit/adapter/base/index
1414
/autoapi/tilelang/jit/adapter/ctypes/index
1515
/autoapi/tilelang/jit/adapter/cython/index
16-
/autoapi/tilelang/jit/adapter/dlpack/index
1716
/autoapi/tilelang/jit/adapter/libgen/index
1817
/autoapi/tilelang/jit/adapter/nvrtc/index
1918
/autoapi/tilelang/jit/adapter/torch/index
19+
/autoapi/tilelang/jit/adapter/tvm_ffi/index
2020
/autoapi/tilelang/jit/adapter/utils/index
2121
/autoapi/tilelang/jit/adapter/wrapper/index
2222

0 commit comments

Comments
 (0)