@@ -595,52 +595,49 @@ <h2>Functions<a class="headerlink" href="#functions" title="Link to this heading
595595< tr class ="row-odd "> < td > < p > < a class ="reference internal " href ="#tilelang.transform.LowerDeviceStorageAccessInfo " title ="tilelang.transform.LowerDeviceStorageAccessInfo "> < code class ="xref py py-obj docutils literal notranslate "> < span class ="pre "> LowerDeviceStorageAccessInfo</ span > </ code > </ a > ()</ p > </ td >
596596< td > < p > Lower attached storage access information on device.</ p > </ td >
597597</ tr >
598- < tr class ="row-even "> < td > < p > < a class ="reference internal " href ="#tilelang.transform.LoopVectorizeDynamic " title ="tilelang.transform.LoopVectorizeDynamic "> < code class ="xref py py-obj docutils literal notranslate "> < span class ="pre "> LoopVectorizeDynamic</ span > </ code > </ a > ()</ p > </ td >
599- < td > < p > Try to vectorize loop with dynamic shape.</ p > </ td >
600- </ tr >
601- < tr class ="row-odd "> < td > < p > < a class ="reference internal " href ="#tilelang.transform.ConfigIndexBitwidth " title ="tilelang.transform.ConfigIndexBitwidth "> < code class ="xref py py-obj docutils literal notranslate "> < span class ="pre "> ConfigIndexBitwidth</ span > </ code > </ a > ()</ p > </ td >
598+ < tr class ="row-even "> < td > < p > < a class ="reference internal " href ="#tilelang.transform.ConfigIndexBitwidth " title ="tilelang.transform.ConfigIndexBitwidth "> < code class ="xref py py-obj docutils literal notranslate "> < span class ="pre "> ConfigIndexBitwidth</ span > </ code > </ a > ()</ p > </ td >
602599< td > < p > Config index bitwidth.</ p > </ td >
603600</ tr >
604- < tr class ="row-even "> < td > < p > < a class ="reference internal " href ="#tilelang.transform.FlattenBuffer " title ="tilelang.transform.FlattenBuffer "> < code class ="xref py py-obj docutils literal notranslate "> < span class ="pre "> FlattenBuffer</ span > </ code > </ a > ()</ p > </ td >
601+ < tr class ="row-odd "> < td > < p > < a class ="reference internal " href ="#tilelang.transform.FlattenBuffer " title ="tilelang.transform.FlattenBuffer "> < code class ="xref py py-obj docutils literal notranslate "> < span class ="pre "> FlattenBuffer</ span > </ code > </ a > ()</ p > </ td >
605602< td > < p > FlattenBuffer</ p > </ td >
606603</ tr >
607- < tr class ="row-odd "> < td > < p > < a class ="reference internal " href ="#tilelang.transform.EliminateStorageSyncForMBarrier " title ="tilelang.transform.EliminateStorageSyncForMBarrier "> < code class ="xref py py-obj docutils literal notranslate "> < span class ="pre "> EliminateStorageSyncForMBarrier</ span > </ code > </ a > ()</ p > </ td >
604+ < tr class ="row-even "> < td > < p > < a class ="reference internal " href ="#tilelang.transform.EliminateStorageSyncForMBarrier " title ="tilelang.transform.EliminateStorageSyncForMBarrier "> < code class ="xref py py-obj docutils literal notranslate "> < span class ="pre "> EliminateStorageSyncForMBarrier</ span > </ code > </ a > ()</ p > </ td >
608605< td > < p > EliminateStorageSyncForMBarrier</ p > </ td >
609606</ tr >
610- < tr class ="row-even "> < td > < p > < a class ="reference internal " href ="#tilelang.transform.MergeSharedMemoryAllocations " title ="tilelang.transform.MergeSharedMemoryAllocations "> < code class ="xref py py-obj docutils literal notranslate "> < span class ="pre "> MergeSharedMemoryAllocations</ span > </ code > </ a > ([...])</ p > </ td >
607+ < tr class ="row-odd "> < td > < p > < a class ="reference internal " href ="#tilelang.transform.MergeSharedMemoryAllocations " title ="tilelang.transform.MergeSharedMemoryAllocations "> < code class ="xref py py-obj docutils literal notranslate "> < span class ="pre "> MergeSharedMemoryAllocations</ span > </ code > </ a > ([...])</ p > </ td >
611608< td > < p > MergeSharedMemoryAllocations</ p > </ td >
612609</ tr >
613- < tr class ="row-odd "> < td > < p > < a class ="reference internal " href ="#tilelang.transform.LowerL2Persistent " title ="tilelang.transform.LowerL2Persistent "> < code class ="xref py py-obj docutils literal notranslate "> < span class ="pre "> LowerL2Persistent</ span > </ code > </ a > ()</ p > </ td >
610+ < tr class ="row-even "> < td > < p > < a class ="reference internal " href ="#tilelang.transform.LowerL2Persistent " title ="tilelang.transform.LowerL2Persistent "> < code class ="xref py py-obj docutils literal notranslate "> < span class ="pre "> LowerL2Persistent</ span > </ code > </ a > ()</ p > </ td >
614611< td > < p > LowerL2Persistent</ p > </ td >
615612</ tr >
616- < tr class ="row-even "> < td > < p > < a class ="reference internal " href ="#tilelang.transform.PersistThreadblock " title ="tilelang.transform.PersistThreadblock "> < code class ="xref py py-obj docutils literal notranslate "> < span class ="pre "> PersistThreadblock</ span > </ code > </ a > ()</ p > </ td >
613+ < tr class ="row-odd "> < td > < p > < a class ="reference internal " href ="#tilelang.transform.PersistThreadblock " title ="tilelang.transform.PersistThreadblock "> < code class ="xref py py-obj docutils literal notranslate "> < span class ="pre "> PersistThreadblock</ span > </ code > </ a > ()</ p > </ td >
617614< td > < p > PersistThreadblock</ p > </ td >
618615</ tr >
619- < tr class ="row-odd "> < td > < p > < a class ="reference internal " href ="#tilelang.transform.AlignDynamicSharedMemoryAllocations " title ="tilelang.transform.AlignDynamicSharedMemoryAllocations "> < code class ="xref py py-obj docutils literal notranslate "> < span class ="pre "> AlignDynamicSharedMemoryAllocations</ span > </ code > </ a > ([align_bytes])</ p > </ td >
616+ < tr class ="row-even "> < td > < p > < a class ="reference internal " href ="#tilelang.transform.AlignDynamicSharedMemoryAllocations " title ="tilelang.transform.AlignDynamicSharedMemoryAllocations "> < code class ="xref py py-obj docutils literal notranslate "> < span class ="pre "> AlignDynamicSharedMemoryAllocations</ span > </ code > </ a > ([align_bytes])</ p > </ td >
620617< td > < p > AlignDynamicSharedMemoryAllocations</ p > </ td >
621618</ tr >
622- < tr class ="row-even "> < td > < p > < a class ="reference internal " href ="#tilelang.transform.LowerSharedBarrier " title ="tilelang.transform.LowerSharedBarrier "> < code class ="xref py py-obj docutils literal notranslate "> < span class ="pre "> LowerSharedBarrier</ span > </ code > </ a > ()</ p > </ td >
619+ < tr class ="row-odd "> < td > < p > < a class ="reference internal " href ="#tilelang.transform.LowerSharedBarrier " title ="tilelang.transform.LowerSharedBarrier "> < code class ="xref py py-obj docutils literal notranslate "> < span class ="pre "> LowerSharedBarrier</ span > </ code > </ a > ()</ p > </ td >
623620< td > < p > LowerSharedBarrier</ p > </ td >
624621</ tr >
625- < tr class ="row-odd "> < td > < p > < a class ="reference internal " href ="#tilelang.transform.StorageRewrite " title ="tilelang.transform.StorageRewrite "> < code class ="xref py py-obj docutils literal notranslate "> < span class ="pre "> StorageRewrite</ span > </ code > </ a > ()</ p > </ td >
622+ < tr class ="row-even "> < td > < p > < a class ="reference internal " href ="#tilelang.transform.StorageRewrite " title ="tilelang.transform.StorageRewrite "> < code class ="xref py py-obj docutils literal notranslate "> < span class ="pre "> StorageRewrite</ span > </ code > </ a > ()</ p > </ td >
626623< td > < p > StorageRewrite</ p > </ td >
627624</ tr >
628- < tr class ="row-even "> < td > < p > < a class ="reference internal " href ="#tilelang.transform.LowerOpaqueBlock " title ="tilelang.transform.LowerOpaqueBlock "> < code class ="xref py py-obj docutils literal notranslate "> < span class ="pre "> LowerOpaqueBlock</ span > </ code > </ a > ()</ p > </ td >
625+ < tr class ="row-odd "> < td > < p > < a class ="reference internal " href ="#tilelang.transform.LowerOpaqueBlock " title ="tilelang.transform.LowerOpaqueBlock "> < code class ="xref py py-obj docutils literal notranslate "> < span class ="pre "> LowerOpaqueBlock</ span > </ code > </ a > ()</ p > </ td >
629626< td > < p > LowerOpaqueBlock</ p > </ td >
630627</ tr >
631- < tr class ="row-odd "> < td > < p > < a class ="reference internal " href ="#tilelang.transform.LowerThreadAllreduce " title ="tilelang.transform.LowerThreadAllreduce "> < code class ="xref py py-obj docutils literal notranslate "> < span class ="pre "> LowerThreadAllreduce</ span > </ code > </ a > ()</ p > </ td >
628+ < tr class ="row-even "> < td > < p > < a class ="reference internal " href ="#tilelang.transform.LowerThreadAllreduce " title ="tilelang.transform.LowerThreadAllreduce "> < code class ="xref py py-obj docutils literal notranslate "> < span class ="pre "> LowerThreadAllreduce</ span > </ code > </ a > ()</ p > </ td >
632629< td > < p > LowerThreadAllreduce</ p > </ td >
633630</ tr >
634- < tr class ="row-even "> < td > < p > < a class ="reference internal " href ="#tilelang.transform.LowerIntrin " title ="tilelang.transform.LowerIntrin "> < code class ="xref py py-obj docutils literal notranslate "> < span class ="pre "> LowerIntrin</ span > </ code > </ a > ()</ p > </ td >
631+ < tr class ="row-odd "> < td > < p > < a class ="reference internal " href ="#tilelang.transform.LowerIntrin " title ="tilelang.transform.LowerIntrin "> < code class ="xref py py-obj docutils literal notranslate "> < span class ="pre "> LowerIntrin</ span > </ code > </ a > ()</ p > </ td >
635632< td > < p > LowerIntrin</ p > </ td >
636633</ tr >
637- < tr class ="row-odd "> < td > < p > < a class ="reference internal " href ="#tilelang.transform.LowerDeviceKernelLaunch " title ="tilelang.transform.LowerDeviceKernelLaunch "> < code class ="xref py py-obj docutils literal notranslate "> < span class ="pre "> LowerDeviceKernelLaunch</ span > </ code > </ a > ()</ p > </ td >
634+ < tr class ="row-even "> < td > < p > < a class ="reference internal " href ="#tilelang.transform.LowerDeviceKernelLaunch " title ="tilelang.transform.LowerDeviceKernelLaunch "> < code class ="xref py py-obj docutils literal notranslate "> < span class ="pre "> LowerDeviceKernelLaunch</ span > </ code > </ a > ()</ p > </ td >
638635< td > < p > Create and return a transform pass that lowers device kernel launch constructs to target-specific IR.</ p > </ td >
639636</ tr >
640- < tr class ="row-even "> < td > < p > < a class ="reference internal " href ="#tilelang.transform.LowerSharedTmem " title ="tilelang.transform.LowerSharedTmem "> < code class ="xref py py-obj docutils literal notranslate "> < span class ="pre "> LowerSharedTmem</ span > </ code > </ a > ()</ p > </ td >
637+ < tr class ="row-odd "> < td > < p > < a class ="reference internal " href ="#tilelang.transform.LowerSharedTmem " title ="tilelang.transform.LowerSharedTmem "> < code class ="xref py py-obj docutils literal notranslate "> < span class ="pre "> LowerSharedTmem</ span > </ code > </ a > ()</ p > </ td >
641638< td > < p > LowerSharedTmem</ p > </ td >
642639</ tr >
643- < tr class ="row-odd "> < td > < p > < a class ="reference internal " href ="#tilelang.transform.LayoutReducer " title ="tilelang.transform.LayoutReducer "> < code class ="xref py py-obj docutils literal notranslate "> < span class ="pre "> LayoutReducer</ span > </ code > </ a > ()</ p > </ td >
640+ < tr class ="row-even "> < td > < p > < a class ="reference internal " href ="#tilelang.transform.LayoutReducer " title ="tilelang.transform.LayoutReducer "> < code class ="xref py py-obj docutils literal notranslate "> < span class ="pre "> LayoutReducer</ span > </ code > </ a > ()</ p > </ td >
644641< td > < p > Return a TVM transform pass that performs layout reduction/normalization.</ p > </ td >
645642</ tr >
646643</ tbody >
@@ -1062,21 +1059,6 @@ <h3>Returns:<a class="headerlink" href="#returns" title="Link to this heading">
10621059</ div >
10631060</ dd > </ dl >
10641061
1065- < dl class ="py function ">
1066- < dt class ="sig sig-object py " id ="tilelang.transform.LoopVectorizeDynamic ">
1067- < span class ="sig-prename descclassname "> < span class ="pre "> tilelang.transform.</ span > </ span > < span class ="sig-name descname "> < span class ="pre "> LoopVectorizeDynamic</ span > </ span > < span class ="sig-paren "> (</ span > < span class ="sig-paren "> )</ span > < a class ="headerlink " href ="#tilelang.transform.LoopVectorizeDynamic " title ="Link to this definition "> ¶</ a > </ dt >
1068- < dd > < p > Try to vectorize loop with dynamic shape.</ p >
1069- < dl class ="field-list simple ">
1070- < dt class ="field-odd "> Returns< span class ="colon "> :</ span > </ dt >
1071- < dd class ="field-odd "> < p > < ul class ="simple ">
1072- < li > < p > < strong > fpass</ strong > (< em > tvm.transform.Pass</ em > ) – The result pass</ p > </ li >
1073- < li > < p > < em > —-</ em > </ p > </ li >
1074- </ ul >
1075- </ p >
1076- </ dd >
1077- </ dl >
1078- </ dd > </ dl >
1079-
10801062< dl class ="py function ">
10811063< dt class ="sig sig-object py " id ="tilelang.transform.ConfigIndexBitwidth ">
10821064< span class ="sig-prename descclassname "> < span class ="pre "> tilelang.transform.</ span > </ span > < span class ="sig-name descname "> < span class ="pre "> ConfigIndexBitwidth</ span > </ span > < span class ="sig-paren "> (</ span > < span class ="sig-paren "> )</ span > < a class ="headerlink " href ="#tilelang.transform.ConfigIndexBitwidth " title ="Link to this definition "> ¶</ a > </ dt >
@@ -1318,7 +1300,6 @@ <h3>Returns:<a class="headerlink" href="#returns" title="Link to this heading">
13181300< li > < a class ="reference internal " href ="#tilelang.transform.VectorizeLoop "> < code class ="docutils literal notranslate "> < span class ="pre "> VectorizeLoop()</ span > </ code > </ a > </ li >
13191301< li > < a class ="reference internal " href ="#tilelang.transform.InjectPTXAsyncCopy "> < code class ="docutils literal notranslate "> < span class ="pre "> InjectPTXAsyncCopy()</ span > </ code > </ a > </ li >
13201302< li > < a class ="reference internal " href ="#tilelang.transform.LowerDeviceStorageAccessInfo "> < code class ="docutils literal notranslate "> < span class ="pre "> LowerDeviceStorageAccessInfo()</ span > </ code > </ a > </ li >
1321- < li > < a class ="reference internal " href ="#tilelang.transform.LoopVectorizeDynamic "> < code class ="docutils literal notranslate "> < span class ="pre "> LoopVectorizeDynamic()</ span > </ code > </ a > </ li >
13221303< li > < a class ="reference internal " href ="#tilelang.transform.ConfigIndexBitwidth "> < code class ="docutils literal notranslate "> < span class ="pre "> ConfigIndexBitwidth()</ span > </ code > </ a > </ li >
13231304< li > < a class ="reference internal " href ="#tilelang.transform.FlattenBuffer "> < code class ="docutils literal notranslate "> < span class ="pre "> FlattenBuffer()</ span > </ code > </ a > </ li >
13241305< li > < a class ="reference internal " href ="#tilelang.transform.EliminateStorageSyncForMBarrier "> < code class ="docutils literal notranslate "> < span class ="pre "> EliminateStorageSyncForMBarrier()</ span > </ code > </ a > </ li >
0 commit comments