| 文件 | 最后提交记录 | 最后更新时间 |
|---|---|---|
Reland "byte permutes in intra-warp layout conversion" (#7933) Reland https://github.com/triton-lang/triton/pull/7809, https://github.com/triton-lang/triton/pull/7825, https://github.com/triton-lang/triton/pull/7861 Add a workaround for ptxas bug and add a regression test | 11 个月前 | |
[Backend] Support broadcasted layouts in warp shuffle gather codegen (#5395) The original implementation side-skirts the issue of dealing with broadcasting layouts. In trying to implement layout selection for gather in the middle end, I found that it's pretty common to have tensors that are too small to avoid broadcasting. E.g. for 4 warps and 32 threads, the tensor needs to have at least 128 elements. It turns out that "enabling" broadcasting support is pretty trivial. Since we know that in a broadcasted layout, broadcasted threads can just index into any other "group" of threads mapped to the same gather column, and that the codegen does not vary based on the broadcasted warps, we can use the pseudo-inverse of the source layout, regardless of what it is computed to. | 1 年前 | |
[Warp Specialization] Allow worker partitions to steal registers from the default partition (#6798) Previously, register reallocation occurs only if the worker partitions have a positive register deficiency, which redistributes the extra registers to the default partition. This PR enables worker partitions to steal from the default partition, which enables, for example, two softmax partitions to go up to 192 registers by stealing from the default partition. | 1 年前 | |
[Dialect] Actually enable TMEM layout check and fix all the tests (#7723) I also removed unnecessary noinline=false from the tests. | 1 年前 | |
[TESTING] Reenable two lit tests (#7641) We split the gather file into two to be able to use different modules for different warp sizes and be able to pipe them to opt. | 1 年前 | |
[Dialect] Actually enable TMEM layout check and fix all the tests (#7723) I also removed unnecessary noinline=false from the tests. | 1 年前 | |
[Dialect] Actually enable TMEM layout check and fix all the tests (#7723) I also removed unnecessary noinline=false from the tests. | 1 年前 | |
[TESTING] Reenable two lit tests (#7641) We split the gather file into two to be able to use different modules for different warp sizes and be able to pipe them to opt. | 1 年前 | |
[TESTING] Reenable two lit tests (#7641) We split the gather file into two to be able to use different modules for different warp sizes and be able to pipe them to opt. | 1 年前 | |
[BACKEND] Add bar.sync before deallocating tmem (#7994) Without a barrier some warp may deallocate tmem while it is still in use causing some other block to override it. | 11 个月前 | |
[BACKEND] hint to LLVM that we can bound threadIdx.x (#7249) | 1 年前 | |
[LAYOUTS] Implement toLinearLayout for TensorMemoryEncodingAttr (#7748) We do so by modelling M/N as describing elements and not the hardware 32bit registers. This allows us to avoid the issue of having two elements pointing to the same register when unpacked=False. We also tighten the MemDescType verifier and the TensorMemoryEncodingAttr verifier to be consistent with the definition we are using. Doing this makes us having to update a ton of lit tests that were silently wrong... | 1 年前 | |
[BACKEND] Fix codegen for ScanOp when there are redundant threads (#5641) This was a mildly tricky bug to track down. Groups of threads with redundant data weren't being masked out, causing them to shuffle data in from threads they weren't supposed to and accumulate them. E.g. if there are 32 threads where the first 16 have unique data and the second half are replicas, lane 16 will shuffle in data from lane 15, 14, 12, etc. and add them in. If the result of the scan is used in such a way that the redundant data is simply discarded, such as stored to global memory, then the invalid values simply aren't observed, but the case that exposed this was a broadcast of the result, causing the invalid results to be observed. | 1 年前 | |
[TensorDesc] Cleanup ttng IR representation (#7036) This moves remaining triton::ExperimentalFoo ops to ttng::Foo, changes ttng::AsyncTMA ops to now take tensordesc arguments, and also removes TensorDescToTMAPtrOp. | 1 年前 | |
[Dialect] Actually enable TMEM layout check and fix all the tests (#7723) I also removed unnecessary noinline=false from the tests. | 1 年前 | |
[Backend] Fix various issues with smem base offsets (#7949) Stacked PRs: * __->__#7949 --- --- --- ### [Backend] Fix various issues with smem base offsets Many ops were using smemObj.getBase() without applying the offsets to get the part of smem to read to write. This is incorrect since the offsets may be nonzero. This PR fixes as many of these issues I could identify and adds tests for this. | 11 个月前 | |
[Gluon] Add support for tmem copy for accumulator tmem lowering (#7956) Extend tmem copy to handle accumulator tmem layout with the proper semantic. This differs from the current scale layout semantic which will have to be fixed later by improving the smem layout representation of the source. | 11 个月前 | |
[LAYOUTS] Implement generalized swizzling for convert_layout (#7565) We generalize the swizzling algorithm to consider the instructions ldmatrix/stmatrix and their transpose versions. To do this, we now require having a dedicated allocator for nvidia, as the required shmem for a convert_layout will now depend on the instructions we can emit. After cleaning up the stmatrix path from the common convert_layout lowering, it became clear that we always take the swizzling path. I changed the allocator to reflect this, and I had to change a ton of tests that used it and now don't require padding. We also implement an improved lowering for the indexing of ldmatrix/stmatrix following the optimisations from ld.shared/st.shared. | 1 年前 | |
[Dialect] Actually enable TMEM layout check and fix all the tests (#7723) I also removed unnecessary noinline=false from the tests. | 1 年前 | |
[Backend] Bump to llvm/llvm-project@bc773632355b (#7881) * Switched Constant{Int|Float}Op type and value order following llvm/llvm-project@a45fda6aeba362926da6cc1b107be92dafb0d490 * Provided triple for TargetLibraryInfoImpl following llvm/llvm-project@c91cbafad2119cace85499e8d231b8e5737f3b41 * Fixed atomic sync scope for NVIDIA following llvm/llvm-project@0f1b16dd5f83fd931ecb111bb925ac9e1d56f589 * Updated MLIR lib names following llvm/llvm-project@e68a20e0b7623738d6af736d3aa02625cba6126a * Updated nvvm.stmatrix op following llvm/llvm-project@2b27377b0bf72e4524774dedf4b03521b07606d5 * Updated ROCDL::Mbcnt{Lo|Hi}Op following llvm/llvm-project@bbe3d64b39d80c2d6132fbad6008b2a6e86fd4d5 Closes https://github.com/triton-lang/triton/pull/7413 Closes https://github.com/triton-lang/triton/pull/7575 Closes https://github.com/triton-lang/triton/pull/7765 --------- Co-authored-by: Yi Qian <yi.qian@amd.com> Co-authored-by: Thomas Raoux <thomas.raoux@openai.com> | 11 个月前 | |
[Dialect] Actually enable TMEM layout check and fix all the tests (#7723) I also removed unnecessary noinline=false from the tests. | 1 年前 | |
[Dialect] Actually enable TMEM layout check and fix all the tests (#7723) I also removed unnecessary noinline=false from the tests. | 1 年前 | |
[LAYOUTS] Implement generalized swizzling for convert_layout (#7565) We generalize the swizzling algorithm to consider the instructions ldmatrix/stmatrix and their transpose versions. To do this, we now require having a dedicated allocator for nvidia, as the required shmem for a convert_layout will now depend on the instructions we can emit. After cleaning up the stmatrix path from the common convert_layout lowering, it became clear that we always take the swizzling path. I changed the allocator to reflect this, and I had to change a ton of tests that used it and now don't require padding. We also implement an improved lowering for the indexing of ldmatrix/stmatrix following the optimisations from ld.shared/st.shared. | 1 年前 | |
Reland "byte permutes in intra-warp layout conversion" (#7933) Reland https://github.com/triton-lang/triton/pull/7809, https://github.com/triton-lang/triton/pull/7825, https://github.com/triton-lang/triton/pull/7861 Add a workaround for ptxas bug and add a regression test | 11 个月前 | |
[ConSan] Support for WGMMA. Checks on non-async shmem and tmem accesses (#7712) This PR covers: * Adding support for wgmma checks (ops marked as read until wgmma_wait) * Adding checks on any shmem and tmem access - this finalizes the basic support for pipelined kernels * Slight refactor of IR - op and attribute names | 1 年前 | |
[Dialect] Actually enable TMEM layout check and fix all the tests (#7723) I also removed unnecessary noinline=false from the tests. | 1 年前 | |
[Backend][NFC] Switch some inline PTX to NVVM ops/intrinsics (#7725) | 1 年前 |
| 文件 | 最后提交记录 | 最后更新时间 |
|---|---|---|
| 11 个月前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 11 个月前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 11 个月前 | ||
| 11 个月前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 11 个月前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 11 个月前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 |