Code refactoring #18

Galaxy1458 · 2025-06-12T07:43:52Z

No description provided.

We unify the support rules for `fp8` dtypes within triton on a per-backend basis. While doing that, we heavily simplify the check code for `tl.dot`, given that we now have much stronger preconditions on the possible dtypes.

This adds two new triton IR operators: 1. `ExperimentalTensormapCreateOp` which creates a descriptor and stores it in global memory 2. `ExperimentalTensormapFenceproxyAcquireOp` which produces the required fence to use the updated descriptor I then use these to expose 3 new functions in `tl.extra.cuda`. 1. `experimental_device_tensormap_create1d` 2. `experimental_device_tensormap_create2d` 3. `experimental_tensormap_fenceproxy_acquire` which match up with the existing host-side tensormap creation API.

Select epilogue results based on iteration predication and allow speculative execution. For instance, when pipelining with num_stages==3 ``` load (0) load(1) local_store(0) %res = for (0..K-1) { dot(i) load(i+2) local_store(i+1) } %d1 = dot(K-2) local_store(K-1) %s1 = select %valid_iteration1, %d1, %res#0 %d0 = dot(K-1) %s0 = select %valid_iteration0, %d0, %s1 ``` This mirrors upstream change llvm/llvm-project#106436

…unch.py` (#4641) Signed-off-by: Anatoly Myachev <anatoly.myachev@intel.com>

Previously, we would search for `libcupti.so` in the `LD_LIBRARY_PATH`. However, we already have `libcupti.so` downloaded from the conda package, and it is the version that exactly matches the CUPTI header file. Therefore, we can simply copy and paste `libcupti.so` to `third_party/nvidia/lib/cupti` and search there first. This eliminates the need to set `LD_LIBRARY_PATH` in many cases.

In a recent change to the [LLVM AMD backend](llvm/llvm-project#83131), we moved the `AMDGPUAttributor` pass into the optimization pipeline (as opposed to the codegen pipeline). Since this is a pass specific for `AMD` targets, we want to pass the `TargetMachine` when building the pipeline, i.e., during the call to `optimize_module`. Failure to do so will result in an increase of number of registers used. Also, we spoke with our LLVM backend team, and they advised to always pass the `TargetMachine` when building the LLVM optimization pipeline. This PR is addressing this issue, in the following way: - I added optional parameters to the `optimize_module` funciton (similar to those passed to `translate_to_asm`) - if those params are passed in, then we will create the `TargetMachine` and pass it to the `PassBuilder` - Otherwise the `TargetMachine` will still be `nullptr` (as it was before) Please note that, as it stands now, this change will only effect the AMD backend.

…res (#4654) If the loop has multiple TMA stores we can re-use the allocation as long as we wait for the store to be finished before starting to use the alloc. We were already doing that so re-using shared memory allocation is a net win. There may be more trade-off to do in the future between shared memory usage and latency hiding.

Remove the restriction that the split dim needs to be the fastest moving one. As long as all the registers are within a thread we can implement splitOp as a no-op. This allows more layout propagation.

…nction (#4656) Signed-off-by: Anatoly Myachev <anatoly.myachev@intel.com>

…t.cpp` file (#4657) Signed-off-by: Anatoly Myachev <anatoly.myachev@intel.com>

…py` (#4643) Signed-off-by: Anatoly Myachev <anatoly.myachev@intel.com>

…rial (#4648) This adds a 4th variant to the persistent matmul tutorial that uses the device-side tensormap creation API. When running the tutorial I do see a small reduction in utilization, but I suppose this is to expected. The result is still superior to not using tma though: ``` ├─ 0.244 matmul_kernel [M=8192, N=8192, K=512] ├─ 0.285 matmul_kernel_device_tma_persistent [M=8192, N=8192, K=512] ├─ 0.259 matmul_kernel_persistent [M=8192, N=8192, K=512] ├─ 0.288 matmul_kernel_tma_persistent [M=8192, N=8192, K=512] ```

…4659) Mirrored from upstream change: llvm/llvm-project#107442

As per title

When running kernel compilation with`MLIR_ENABLE_DUMP=1`, the dump doesn't always work when the compiled kernel is the triton cache. Repro: ``` MLIR_ENABLE_DUMP=1 python3 python/triton/tools/compile.py \ --kernel-name add_kernel \ --signature "*fp32,*fp32,*fp32,i32,64" \ --grid=1024,1024,1024 \ python/tutorials/01-vector-add.py ``` prints MLIR, However, running it a second time results in no dump. Adding this to README since it was not obvious to me and took me some time to figure this out.

…pp` and `RewriteTensorPointer.cpp` (#4658) Signed-off-by: Anatoly Myachev <anatoly.myachev@intel.com>

The behavior for INTERPRET mode is different from the expected behavior in the test_reduce1d. For numpy <2.0 this difference does not matter, but for version >=2.0 an error occurs. This change fixes it. The core Triton is a small number of people, and we receive many PRs (thank you!). To help us review your code more quickly, **if you are a new contributor (less than 3 PRs merged) we ask that you complete the following tasks and include the filled-out checklist in your PR description.** Complete the following tasks before sending your PR, and replace `[ ]` with `[x]` to indicate you have done them. - [x] I am not making a trivial change, such as fixing a typo in a comment. - [x] I have written a PR description following these [rules](https://cbea.ms/git-commit/#why-not-how). - [x] I have run `pre-commit run --from-ref origin/main --to-ref HEAD`. - Select one of the following. - [ ] I have added tests. - `/test` for `lit` tests - `/unittest` for C++ tests - `/python/test` for end-to-end tests - [x] This PR does not need a test because `FILL THIS IN`. - Select one of the following. - [x] I have not added any `lit` tests. - [ ] The `lit` tests I have added follow these [best practices](https://mlir.llvm.org/getting_started/TestingGuide/#filecheck-best-practices), including the "tests should be minimal" section. (Usually running Python code and using the instructions it generates is not minimal.) Signed-off-by: Kirill Suvorov <kirill.suvorov@intel.com>

* Added predication for tt.store for tests * Moved ttng.warp_group_dot for tests

Bumping llvm to include a loop unroller fix: https://github.com/llvm/llvm-project/pull/106164/files . This is needed for subsequent loop unroller upstreaming work.

Explicit is better than implicit

…antics for scalars (#4613) The idea here is that if you have a tensor `t` of dtype `uint8` and you want to do `t << 2`, the result should be of dtype `uint8`, not `int32`! We do this for all dunder ops that don't output booleans. This follows roughly the semantics of PyTorch, JAX and NumPy 2.0. I would like to document this behaviour, but it's not clear to me where is the best place to say so. The PR has much more churn than I would like, as I had to move the `to_tensor` method to `semantic` (which is where it belongs anyway). For reviewers, the only two relevant changes are in `computation_type_impl` and in `bitwise_op_type_checking_impl`, where we say that we do perform casting for bitwise ops.

… (#4176)" (#4631) Turn back aggressive strategy by default to enable block merging given now we have upstream fixes for it brought in: triton-lang/triton#4619. This reverts commit cf2ad02.

as per title

triton-lang/triton#4655 sets the TargetMachine primarily for the AMD backend but we want it to remain nullptr when using LLVM IR level plugins as well. This PR just adds an additional check when adding TargetMachine value.

… (#4675) With block level kind of operations like TMA it is possible that some ops access the shared memory but don't require barriers. This adds a lambda that backends can pass to explicitly skip barriers in between some ops.

Before, the ===== sections were being shown in the sidebar of the docs. ![image](https://github.com/user-attachments/assets/deb9bef0-f8d2-4152-8d17-e009495d75e9)

…s (#4679) Incrementing refcount for Py_None in TMA descriptor runtime helpers to avoid deallocating the None object. This fixes a runtime failure like Fatal Python error: none_dealloc: deallocating None

There was a bug in how we were splitting the uniform/non-uniform offset contribution for addptr. Consider this IR (where U is a uniform value, e.g., , coming from a splat and NU is non-uniform, coming e.g., from a `make_range`). ``` %a = %U+%NU %b = %a + %NU %c = addptr %ptr, %b ``` It would have been rewritten to ``` %b = %NU+%NU %basePtr = addptr %basePtrOld, %U %c = addi %offset, %b ``` The main issue here is that `%b`'s operand #0 has changed, i.e., the scalar contribution has been removed. This is fine if `addptr` is the only operation that uses `%b`. If any other operation uses `%b`, they need the "old" `%b`. The solution is to accumulate both the uniform and non-uniform contributions in a separate IR and leave the original `%b` untouched. Possible duplications will be removed by the canonicalizer . Doing things in this way, I also could generalize the pass to all expressions of the form `(U+NU)*(U+NU)`. I tried enabling this pass and running all the suite and it is working fine.

@Jokeren

Also enable the special rules by default in lowering. Problem caught by @Jokeren during code review.

Move writing to LDS and reading from LDS right after the loading of a tensor from global memory. This PR does reordering by considering 2 possible patterns depending on whether writing to LDS is done using an optional local_alloc argument or a local_store instruction: 1) load -> local_alloc -> local_store -> local_load, 2) load -> local_alloc -> local_load. --------- Co-authored-by: Ognjen Plavsic <ognjen.plavsic@luxoft.com> Co-authored-by: Lei Zhang <antiagainst@gmail.com>

…ments (#5) [BUILD] Fix ext_sourcedir in editable_wheel mode [BUILD] Update requirements

Fix build and test workflow, merge build and test jobs to one.

This reverts commit cf34004.

…` in `do_bench` (#4733)" This reverts commit 566b63c.

This reverts commit fc879bd.

This reverts commit 23fefa6.

…ange (#4696)" This reverts commit 5fe38ff.

…tion on AMD gpus (#4569)" This reverts commit 757b6a6.

This reverts commit b514024.

…339)" This reverts commit 91f24d8.

This reverts commit 5cbb516.

…ng (#4061) (#4223)" This reverts commit 107465b.

This reverts commit 21eae95.

…164)" This reverts commit dedb7bd.

This reverts commit 5d77ae2.

This reverts commit dcc22b8.

…shape if it is not needed (#4161)" This reverts commit 3d1da66.

… 3.8 compatibility (#4160)" This reverts commit 45b0f6b.

This reverts commit 085b010.

* [init adaptation] init adapt ascend_backend for flagtree * [add triton-adapter-opt] * [Bug fix] fix import bug * Update README.md * [Bug fix] fix product so-package in editable mode * [Cache Tools] cache ascend tools --------- Co-authored-by: Galaxy1458 <55453380+Galaxy1458@users.noreply.github.com>

CLAassistant · 2025-06-12T07:44:07Z

Thank you for your submission! We really appreciate it. Like many open source projects, we ask that you all sign our Contributor License Agreement before we can accept your contribution.
2 out of 12 committers have signed the CLA.

✅ zhzhcookie
✅ StrongSpoon
❌ alexbaden
❌ binarman
❌ bertmaher
❌ peterbell10
❌ aakhundov
❌ ThomasRaoux
❌ htyu
❌ jataylo
❌ atalman
❌ davidberard98
_{You have signed the CLA already but the status is still pending? Let us recheck it.}

When running [convert_blocked1d_to_slice0](https://github.com/triton-lang/triton/blob/0ba5f0c3cd029d5c3d1f01b9bf29dac32c27345e/test/Conversion/tritongpu_to_llvm.mlir#L924) Triton ends up computing a rank of a matrix with 0 columns during linear layout lowering, which trips up f2reduce, and causes undefined behavior, detectable through [UBSAN](https://clang.llvm.org/docs/UndefinedBehaviorSanitizer.html). Fix this by returning the rank (0) early in these cases, without calling f2reduce. <details><summary>Stack trace</summary> <p> ``` third_party/triton/third_party/f2reduce/f2reduce.cpp:421:30: runtime error: shift exponent 18446744073709551615 is too large for 64-bit type 'unsigned long long' #0 0x556ee2fea3be in inplace_rref_small third_party/triton/third_party/f2reduce/f2reduce.cpp:421:30 flagos-ai#1 0x556ee2fea3be in f2reduce::inplace_rref_strided(unsigned long*, unsigned long, unsigned long, unsigned long) third_party/triton/third_party/f2reduce/f2reduce.cpp:470:9 flagos-ai#2 0x556ee2ea70da in getMatrixRank third_party/triton/lib/Tools/LinearLayout.cpp:125:3 flagos-ai#3 0x556ee2ea70da in mlir::triton::LinearLayout::checkInvariants(bool) third_party/triton/lib/Tools/LinearLayout.cpp:299:7 flagos-ai#4 0x556ee2ea656d in mlir::triton::LinearLayout::tryCreate(llvm::MapVector<mlir::StringAttr, std::__u::vector<std::__u::vector<int, std::__u::allocator<int>>, std::__u::allocator<std::__u::vector<int, std::__u::allocator<int>>>>, llvm::DenseMap<mlir::StringAttr, unsigned int, llvm::DenseMapInfo<mlir::StringAttr, void>, llvm::detail::DenseMapPair<mlir::StringAttr, unsigned int>>, llvm::SmallVector<std::__u::pair<mlir::StringAttr, std::__u::vector<std::__u::vector<int, std::__u::allocator<int>>, std::__u::allocator<std::__u::vector<int, std::__u::allocator<int>>>>>, 0u>>, llvm::ArrayRef<std::__u::pair<mlir::StringAttr, int>>, bool) third_party/triton/lib/Tools/LinearLayout.cpp:190:41 flagos-ai#5 0x556ee2eb2150 in mlir::triton::LinearLayout::divideRight(mlir::triton::LinearLayout const&) third_party/triton/lib/Tools/LinearLayout.cpp:654:51 flagos-ai#6 0x556ee2ee1c39 in mlir::cvtNeedsSharedMemory(mlir::RankedTensorType, mlir::RankedTensorType) third_party/triton/lib/Analysis/Utility.cpp:652:14 flagos-ai#7 0x556ee2cf38fd in mlir::triton::getRepShapeForCvtLayout(mlir::triton::gpu::ConvertLayoutOp) third_party/triton/lib/Analysis/Allocation.cpp:66:8 flagos-ai#8 0x556ee2cf3efa in mlir::triton::getScratchConfigForCvtLayout(mlir::triton::gpu::ConvertLayoutOp, unsigned int&, unsigned int&) third_party/triton/lib/Analysis/Allocation.cpp:95:19 flagos-ai#9 0x556ee2cf6057 in mlir::triton::AllocationAnalysis::getScratchValueSize(mlir::Operation*) third_party/triton/lib/Analysis/Allocation.cpp:272:24 flagos-ai#10 0x556ee2cf5499 in operator() third_party/triton/lib/Analysis/Allocation.cpp:343:7 flagos-ai#11 0x556ee2cf5499 in void llvm::function_ref<void (mlir::Operation*)>::callback_fn<mlir::triton::AllocationAnalysis::getValuesAndSizes()::'lambda'(mlir::Operation*)>(long, mlir::Operation*) third_party/llvm/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:45:12 flagos-ai#12 0x556edeeee7a9 in operator() third_party/llvm/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:68:12 flagos-ai#13 0x556edeeee7a9 in void mlir::detail::walk<mlir::ForwardIterator>(mlir::Operation*, llvm::function_ref<void (mlir::Operation*)>, mlir::WalkOrder) third_party/llvm/llvm-project/mlir/include/mlir/IR/Visitors.h:174:5 flagos-ai#14 0x556edeeee87c in void mlir::detail::walk<mlir::ForwardIterator>(mlir::Operation*, llvm::function_ref<void (mlir::Operation*)>, mlir::WalkOrder) third_party/llvm/llvm-project/mlir/include/mlir/IR/Visitors.h:182:9 flagos-ai#15 0x556ee2cf49e7 in walk<(mlir::WalkOrder)0, mlir::ForwardIterator, (lambda at third_party/triton/lib/Analysis/Allocation.cpp:341:42), mlir::Operation *, void> third_party/llvm/llvm-project/mlir/include/mlir/IR/Visitors.h:313:10 flagos-ai#16 0x556ee2cf49e7 in walk<(mlir::WalkOrder)0, mlir::ForwardIterator, (lambda at third_party/triton/lib/Analysis/Allocation.cpp:341:42), void> third_party/llvm/llvm-project/mlir/include/mlir/IR/Operation.h:794:12 flagos-ai#17 0x556ee2cf49e7 in mlir::triton::AllocationAnalysis::getValuesAndSizes() third_party/triton/lib/Analysis/Allocation.cpp:341:16 flagos-ai#18 0x556ee2cf4852 in run third_party/triton/lib/Analysis/Allocation.cpp:182:5 flagos-ai#19 0x556ee2cf4852 in AllocationAnalysis third_party/triton/lib/Analysis/Allocation.cpp:169:5 flagos-ai#20 0x556ee2cf4852 in mlir::Allocation::run(llvm::DenseMap<mlir::FunctionOpInterface, mlir::Allocation, llvm::DenseMapInfo<mlir::FunctionOpInterface, void>, llvm::detail::DenseMapPair<mlir::FunctionOpInterface, mlir::Allocation>>&) third_party/triton/lib/Analysis/Allocation.cpp:627:3 flagos-ai#21 0x556ee1677402 in operator() third_party/triton/include/triton/Analysis/Allocation.h:227:26 flagos-ai#22 0x556ee1677402 in void mlir::CallGraph<mlir::Allocation>::doWalk<(mlir::WalkOrder)0, (mlir::WalkOrder)1, mlir::ModuleAllocation::ModuleAllocation(mlir::ModuleOp)::'lambda'(mlir::CallOpInterface, mlir::FunctionOpInterface), mlir::ModuleAllocation::ModuleAllocation(mlir::ModuleOp)::'lambda'(mlir::FunctionOpInterface)>(mlir::FunctionOpInterface, llvm::DenseSet<mlir::FunctionOpInterface, llvm::DenseMapInfo<mlir::FunctionOpInterface, void>>&, mlir::ModuleAllocation::ModuleAllocation(mlir::ModuleOp)::'lambda'(mlir::CallOpInterface, mlir::FunctionOpInterface), mlir::ModuleAllocation::ModuleAllocation(mlir::ModuleOp)::'lambda'(mlir::FunctionOpInterface)) third_party/triton/include/triton/Analysis/Utility.h:350:7 flagos-ai#23 0x556ee16756b3 in walk<(mlir::WalkOrder)0, (mlir::WalkOrder)1, (lambda at third_party/triton/include/triton/Analysis/Allocation.h:222:9), (lambda at third_party/triton/include/triton/Analysis/Allocation.h:224:9)> third_party/triton/include/triton/Analysis/Utility.h:242:7 flagos-ai#24 0x556ee16756b3 in mlir::ModuleAllocation::ModuleAllocation(mlir::ModuleOp) third_party/triton/include/triton/Analysis/Allocation.h:220:5 flagos-ai#25 0x556ee2c2bf18 in (anonymous namespace)::AllocateSharedMemory::runOnOperation() third_party/triton/lib/Conversion/TritonGPUToLLVM/AllocateSharedMemory.cpp:26:22 ... UndefinedBehaviorSanitizer: invalid-shift-exponent third_party/triton/third_party/f2reduce/f2reduce.cpp:421:30 ``` </p> </details>

…leaveTMem.cpp (#7924) `TritonNvidiaGPU/interleave_tmem.mlir` fails under address sanitizer. The `ConstantIntOp` operations were created without attachment to any block in http://github.com/triton-lang/triton/pull/7622, which caused a memory leak. This change addresses the problem by adding an insertion point. <details open> <summary>Full log</summary> ================================================================= ==3831==ERROR: LeakSanitizer: detected memory leaks Direct leak of 576 byte(s) in 6 object(s) allocated from: #0 0x55c3eca39164 in malloc [third_party/llvm/llvm-project/compiler-rt/lib/asan/asan_malloc_linux.cpp:67](https://cs.corp.google.com/piper///depot/google3/third_party/llvm/llvm-project/compiler-rt/lib/asan/asan_malloc_linux.cpp?l=67&ws=tap-presubmit-server/421956858&snapshot=2):3 #1 0x55c3f176afb3 in mlir::Operation::create(mlir::Location, mlir::OperationName, mlir::TypeRange, mlir::ValueRange, mlir::DictionaryAttr, mlir::OpaqueProperties, mlir::BlockRange, unsigned int) [third_party/llvm/llvm-project/mlir/lib/IR/Operation.cpp:113](https://cs.corp.google.com/piper///depot/google3/third_party/llvm/llvm-project/mlir/lib/IR/Operation.cpp?l=113&ws=tap-presubmit-server/421956858&snapshot=2):46 #2 0x55c3f176a90c in create [third_party/llvm/llvm-project/mlir/lib/IR/Operation.cpp:74](https://cs.corp.google.com/piper///depot/google3/third_party/llvm/llvm-project/mlir/lib/IR/Operation.cpp?l=74&ws=tap-presubmit-server/421956858&snapshot=2):10 #3 0x55c3f176a90c in mlir::Operation::create(mlir::Location, mlir::OperationName, mlir::TypeRange, mlir::ValueRange, mlir::NamedAttrList&&, mlir::OpaqueProperties, mlir::BlockRange, mlir::RegionRange) [third_party/llvm/llvm-project/mlir/lib/IR/Operation.cpp:57](https://cs.corp.google.com/piper///depot/google3/third_party/llvm/llvm-project/mlir/lib/IR/Operation.cpp?l=57&ws=tap-presubmit-server/421956858&snapshot=2):7 #4 0x55c3f176a61b in mlir::Operation::create(mlir::OperationState const&) [third_party/llvm/llvm-project/mlir/lib/IR/Operation.cpp:35](https://cs.corp.google.com/piper///depot/google3/third_party/llvm/llvm-project/mlir/lib/IR/Operation.cpp?l=35&ws=tap-presubmit-server/421956858&snapshot=2):7 #5 0x55c3f1678a78 in mlir::OpBuilder::create(mlir::OperationState const&) [third_party/llvm/llvm-project/mlir/lib/IR/Builders.cpp:453](https://cs.corp.google.com/piper///depot/google3/third_party/llvm/llvm-project/mlir/lib/IR/Builders.cpp?l=453&ws=tap-presubmit-server/421956858&snapshot=2):17 #6 0x55c3ecf3668f in mlir::arith::ConstantIntOp mlir::OpBuilder::create<mlir::arith::ConstantIntOp, int, int>(mlir::Location, int&&, int&&) [third_party/llvm/llvm-project/mlir/include/mlir/IR/Builders.h:507](https://cs.corp.google.com/piper///depot/google3/third_party/llvm/llvm-project/mlir/include/mlir/IR/Builders.h?l=507&ws=tap-presubmit-server/421956858&snapshot=2):16 #7 0x55c3eefa690a in findBufferAccessMemdescSubview [third_party/triton/lib/Dialect/TritonNvidiaGPU/Transforms/InterleaveTMem.cpp:75](https://cs.corp.google.com/piper///depot/google3/third_party/triton/lib/Dialect/TritonNvidiaGPU/Transforms/InterleaveTMem.cpp?l=75&ws=tap-presubmit-server/421956858&snapshot=2):33 #8 0x55c3eefa690a in mlir::triton::nvidia_gpu::(anonymous namespace)::findBufferAccess(mlir::Value) [third_party/triton/lib/Dialect/TritonNvidiaGPU/Transforms/InterleaveTMem.cpp:151](https://cs.corp.google.com/piper///depot/google3/third_party/triton/lib/Dialect/TritonNvidiaGPU/Transforms/InterleaveTMem.cpp?l=151&ws=tap-presubmit-server/421956858&snapshot=2):12 #9 0x55c3eefa70e7 in mlir::triton::nvidia_gpu::(anonymous namespace)::findBufferAccess(mlir::Value) [third_party/triton/lib/Dialect/TritonNvidiaGPU/Transforms/InterleaveTMem.cpp:156](https://cs.corp.google.com/piper///depot/google3/third_party/triton/lib/Dialect/TritonNvidiaGPU/Transforms/InterleaveTMem.cpp?l=156&ws=tap-presubmit-server/421956858&snapshot=2):34 #10 0x55c3eefa4c0c in tmemMayAlias [third_party/triton/lib/Dialect/TritonNvidiaGPU/Transforms/InterleaveTMem.cpp:173](https://cs.corp.google.com/piper///depot/google3/third_party/triton/lib/Dialect/TritonNvidiaGPU/Transforms/InterleaveTMem.cpp?l=173&ws=tap-presubmit-server/421956858&snapshot=2):28 #11 0x55c3eefa4c0c in sinkOps [third_party/triton/lib/Dialect/TritonNvidiaGPU/Transforms/InterleaveTMem.cpp:227](https://cs.corp.google.com/piper///depot/google3/third_party/triton/lib/Dialect/TritonNvidiaGPU/Transforms/InterleaveTMem.cpp?l=227&ws=tap-presubmit-server/421956858&snapshot=2):36 #12 0x55c3eefa4c0c in trySinkOp [third_party/triton/lib/Dialect/TritonNvidiaGPU/Transforms/InterleaveTMem.cpp:253](https://cs.corp.google.com/piper///depot/google3/third_party/triton/lib/Dialect/TritonNvidiaGPU/Transforms/InterleaveTMem.cpp?l=253&ws=tap-presubmit-server/421956858&snapshot=2):10 #13 0x55c3eefa4c0c in mlir::triton::nvidia_gpu::TritonNvidiaGPUInterleaveTMemPass::runOnOperation() [third_party/triton/lib/Dialect/TritonNvidiaGPU/Transforms/InterleaveTMem.cpp:275](https://cs.corp.google.com/piper///depot/google3/third_party/triton/lib/Dialect/TritonNvidiaGPU/Transforms/InterleaveTMem.cpp?l=275&ws=tap-presubmit-server/421956858&snapshot=2):14 #14 0x55c3f1560ad1 in operator() [third_party/llvm/llvm-project/mlir/lib/Pass/Pass.cpp:553](https://cs.corp.google.com/piper///depot/google3/third_party/llvm/llvm-project/mlir/lib/Pass/Pass.cpp?l=553&ws=tap-presubmit-server/421956858&snapshot=2):17 #15 0x55c3f1560ad1 in void llvm::function_ref<void ()>::callback_fn<mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int)::$_1>(long) [third_party/llvm/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:46](https://cs.corp.google.com/piper///depot/google3/third_party/llvm/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h?l=46&ws=tap-presubmit-server/421956858&snapshot=2):12 #16 0x55c3f1559920 in operator() [third_party/llvm/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:69](https://cs.corp.google.com/piper///depot/google3/third_party/llvm/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h?l=69&ws=tap-presubmit-server/421956858&snapshot=2):12 #17 0x55c3f1559920 in executeAction<mlir::PassExecutionAction, mlir::Pass &> [third_party/llvm/llvm-project/mlir/include/mlir/IR/MLIRContext.h:280](https://cs.corp.google.com/piper///depot/google3/third_party/llvm/llvm-project/mlir/include/mlir/IR/MLIRContext.h?l=280&ws=tap-presubmit-server/421956858&snapshot=2):7 #18 0x55c3f1559920 in mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) [third_party/llvm/llvm-project/mlir/lib/Pass/Pass.cpp:547](https://cs.corp.google.com/piper///depot/google3/third_party/llvm/llvm-project/mlir/lib/Pass/Pass.cpp?l=547&ws=tap-presubmit-server/421956858&snapshot=2):21 #19 0x55c3f155d46f in runPipeline [third_party/llvm/llvm-project/mlir/lib/Pass/Pass.cpp:619](https://cs.corp.google.com/piper///depot/google3/third_party/llvm/llvm-project/mlir/lib/Pass/Pass.cpp?l=619&ws=tap-presubmit-server/421956858&snapshot=2):16 #20 0x55c3f155d46f in mlir::PassManager::runPasses(mlir::Operation*, mlir::AnalysisManager) [third_party/llvm/llvm-project/mlir/lib/Pass/Pass.cpp:933](https://cs.corp.google.com/piper///depot/google3/third_party/llvm/llvm-project/mlir/lib/Pass/Pass.cpp?l=933&ws=tap-presubmit-server/421956858&snapshot=2):10 #21 0x55c3f155d15b in mlir::PassManager::run(mlir::Operation*) [third_party/llvm/llvm-project/mlir/lib/Pass/Pass.cpp:913](https://cs.corp.google.com/piper///depot/google3/third_party/llvm/llvm-project/mlir/lib/Pass/Pass.cpp?l=913&ws=tap-presubmit-server/421956858&snapshot=2):60 #22 0x55c3ed0a8b20 in performActions(llvm::raw_ostream&, std::__u::shared_ptr<llvm::SourceMgr> const&, mlir::MLIRContext*, mlir::MlirOptMainConfig const&) [third_party/llvm/llvm-project/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:477](https://cs.corp.google.com/piper///depot/google3/third_party/llvm/llvm-project/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp?l=477&ws=tap-presubmit-server/421956858&snapshot=2):17 #23 0x55c3ed0a8363 in processBuffer [third_party/llvm/llvm-project/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:553](https://cs.corp.google.com/piper///depot/google3/third_party/llvm/llvm-project/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp?l=553&ws=tap-presubmit-server/421956858&snapshot=2):12 #24 0x55c3ed0a8363 in operator() [third_party/llvm/llvm-project/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:642](https://cs.corp.google.com/piper///depot/google3/third_party/llvm/llvm-project/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp?l=642&ws=tap-presubmit-server/421956858&snapshot=2):12 #25 0x55c3ed0a8363 in llvm::LogicalResult llvm::function_ref<llvm::LogicalResult (std::__u::unique_ptr<llvm::MemoryBuffer, std::__u::default_delete<llvm::MemoryBuffer>>, llvm::MemoryBufferRef const&, llvm::raw_ostream&)>::callback_fn<mlir::MlirOptMain(llvm::raw_ostream&, std::__u::unique_ptr<llvm::MemoryBuffer, std::__u::default_delete<llvm::MemoryBuffer>>, mlir::DialectRegistry&, mlir::MlirOptMainConfig const&)::$_0>(long, std::__u::unique_ptr<llvm::MemoryBuffer, std::__u::default_delete<llvm::MemoryBuffer>>, llvm::MemoryBufferRef const&, llvm::raw_ostream&) [third_party/llvm/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:46](https://cs.corp.google.com/piper///depot/google3/third_party/llvm/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h?l=46&ws=tap-presubmit-server/421956858&snapshot=2):12 #26 0x55c3f17bd34f in operator() [third_party/llvm/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:69](https://cs.corp.google.com/piper///depot/google3/third_party/llvm/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h?l=69&ws=tap-presubmit-server/421956858&snapshot=2):12 #27 0x55c3f17bd34f in mlir::splitAndProcessBuffer(std::__u::unique_ptr<llvm::MemoryBuffer, std::__u::default_delete<llvm::MemoryBuffer>>, llvm::function_ref<llvm::LogicalResult (std::__u::unique_ptr<llvm::MemoryBuffer, std::__u::default_delete<llvm::MemoryBuffer>>, llvm::MemoryBufferRef const&, llvm::raw_ostream&)>, llvm::raw_ostream&, llvm::StringRef, llvm::StringRef) [third_party/llvm/llvm-project/mlir/lib/Support/ToolUtilities.cpp:30](https://cs.corp.google.com/piper///depot/google3/third_party/llvm/llvm-project/mlir/lib/Support/ToolUtilities.cpp?l=30&ws=tap-presubmit-server/421956858&snapshot=2):12 #28 0x55c3ed09d0c6 in mlir::MlirOptMain(llvm::raw_ostream&, std::__u::unique_ptr<llvm::MemoryBuffer, std::__u::default_delete<llvm::MemoryBuffer>>, mlir::DialectRegistry&, mlir::MlirOptMainConfig const&) [third_party/llvm/llvm-project/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:647](https://cs.corp.google.com/piper///depot/google3/third_party/llvm/llvm-project/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp?l=647&ws=tap-presubmit-server/421956858&snapshot=2):26 #29 0x55c3ed09d67f in mlir::MlirOptMain(int, char**, llvm::StringRef, llvm::StringRef, mlir::DialectRegistry&) [third_party/llvm/llvm-project/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:693](https://cs.corp.google.com/piper///depot/google3/third_party/llvm/llvm-project/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp?l=693&ws=tap-presubmit-server/421956858&snapshot=2):14 #30 0x55c3ed09dc59 in mlir::MlirOptMain(int, char**, llvm::StringRef, mlir::DialectRegistry&) [third_party/llvm/llvm-project/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:709](https://cs.corp.google.com/piper///depot/google3/third_party/llvm/llvm-project/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp?l=709&ws=tap-presubmit-server/421956858&snapshot=2):10 #31 0x55c3eca74a70 in main [third_party/triton/bin/triton-opt.cpp:14](https://cs.corp.google.com/piper///depot/google3/third_party/triton/bin/triton-opt.cpp?l=14&ws=tap-presubmit-server/421956858&snapshot=2):33 #32 0x7f1fd58613d3 in __libc_start_main (/usr/grte/v5/lib64/libc.so.6+0x613d3) (BuildId: 9a996398ce14a94560b0c642eb4f6e94) #33 0x55c3ec995aa9 in _start /usr/grte/v5/debug-src/src/csu/../sysdeps/x86_64/start.S:120 </details> --------- Co-authored-by: Thomas Raoux <thomas.raoux@openai.com>

When running [convert_blocked1d_to_slice0](https://github.com/triton-lang/triton/blob/0ba5f0c3cd029d5c3d1f01b9bf29dac32c27345e/test/Conversion/tritongpu_to_llvm.mlir#L924) Triton ends up computing a rank of a matrix with 0 columns during linear layout lowering, which trips up f2reduce, and causes undefined behavior, detectable through [UBSAN](https://clang.llvm.org/docs/UndefinedBehaviorSanitizer.html). Fix this by returning the rank (0) early in these cases, without calling f2reduce. <details><summary>Stack trace</summary> <p> ``` third_party/triton/third_party/f2reduce/f2reduce.cpp:421:30: runtime error: shift exponent 18446744073709551615 is too large for 64-bit type 'unsigned long long' #0 0x556ee2fea3be in inplace_rref_small third_party/triton/third_party/f2reduce/f2reduce.cpp:421:30 #1 0x556ee2fea3be in f2reduce::inplace_rref_strided(unsigned long*, unsigned long, unsigned long, unsigned long) third_party/triton/third_party/f2reduce/f2reduce.cpp:470:9 #2 0x556ee2ea70da in getMatrixRank third_party/triton/lib/Tools/LinearLayout.cpp:125:3 #3 0x556ee2ea70da in mlir::triton::LinearLayout::checkInvariants(bool) third_party/triton/lib/Tools/LinearLayout.cpp:299:7 #4 0x556ee2ea656d in mlir::triton::LinearLayout::tryCreate(llvm::MapVector<mlir::StringAttr, std::__u::vector<std::__u::vector<int, std::__u::allocator<int>>, std::__u::allocator<std::__u::vector<int, std::__u::allocator<int>>>>, llvm::DenseMap<mlir::StringAttr, unsigned int, llvm::DenseMapInfo<mlir::StringAttr, void>, llvm::detail::DenseMapPair<mlir::StringAttr, unsigned int>>, llvm::SmallVector<std::__u::pair<mlir::StringAttr, std::__u::vector<std::__u::vector<int, std::__u::allocator<int>>, std::__u::allocator<std::__u::vector<int, std::__u::allocator<int>>>>>, 0u>>, llvm::ArrayRef<std::__u::pair<mlir::StringAttr, int>>, bool) third_party/triton/lib/Tools/LinearLayout.cpp:190:41 #5 0x556ee2eb2150 in mlir::triton::LinearLayout::divideRight(mlir::triton::LinearLayout const&) third_party/triton/lib/Tools/LinearLayout.cpp:654:51 #6 0x556ee2ee1c39 in mlir::cvtNeedsSharedMemory(mlir::RankedTensorType, mlir::RankedTensorType) third_party/triton/lib/Analysis/Utility.cpp:652:14 #7 0x556ee2cf38fd in mlir::triton::getRepShapeForCvtLayout(mlir::triton::gpu::ConvertLayoutOp) third_party/triton/lib/Analysis/Allocation.cpp:66:8 #8 0x556ee2cf3efa in mlir::triton::getScratchConfigForCvtLayout(mlir::triton::gpu::ConvertLayoutOp, unsigned int&, unsigned int&) third_party/triton/lib/Analysis/Allocation.cpp:95:19 #9 0x556ee2cf6057 in mlir::triton::AllocationAnalysis::getScratchValueSize(mlir::Operation*) third_party/triton/lib/Analysis/Allocation.cpp:272:24 #10 0x556ee2cf5499 in operator() third_party/triton/lib/Analysis/Allocation.cpp:343:7 #11 0x556ee2cf5499 in void llvm::function_ref<void (mlir::Operation*)>::callback_fn<mlir::triton::AllocationAnalysis::getValuesAndSizes()::'lambda'(mlir::Operation*)>(long, mlir::Operation*) third_party/llvm/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:45:12 #12 0x556edeeee7a9 in operator() third_party/llvm/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:68:12 #13 0x556edeeee7a9 in void mlir::detail::walk<mlir::ForwardIterator>(mlir::Operation*, llvm::function_ref<void (mlir::Operation*)>, mlir::WalkOrder) third_party/llvm/llvm-project/mlir/include/mlir/IR/Visitors.h:174:5 #14 0x556edeeee87c in void mlir::detail::walk<mlir::ForwardIterator>(mlir::Operation*, llvm::function_ref<void (mlir::Operation*)>, mlir::WalkOrder) third_party/llvm/llvm-project/mlir/include/mlir/IR/Visitors.h:182:9 #15 0x556ee2cf49e7 in walk<(mlir::WalkOrder)0, mlir::ForwardIterator, (lambda at third_party/triton/lib/Analysis/Allocation.cpp:341:42), mlir::Operation *, void> third_party/llvm/llvm-project/mlir/include/mlir/IR/Visitors.h:313:10 #16 0x556ee2cf49e7 in walk<(mlir::WalkOrder)0, mlir::ForwardIterator, (lambda at third_party/triton/lib/Analysis/Allocation.cpp:341:42), void> third_party/llvm/llvm-project/mlir/include/mlir/IR/Operation.h:794:12 #17 0x556ee2cf49e7 in mlir::triton::AllocationAnalysis::getValuesAndSizes() third_party/triton/lib/Analysis/Allocation.cpp:341:16 #18 0x556ee2cf4852 in run third_party/triton/lib/Analysis/Allocation.cpp:182:5 #19 0x556ee2cf4852 in AllocationAnalysis third_party/triton/lib/Analysis/Allocation.cpp:169:5 #20 0x556ee2cf4852 in mlir::Allocation::run(llvm::DenseMap<mlir::FunctionOpInterface, mlir::Allocation, llvm::DenseMapInfo<mlir::FunctionOpInterface, void>, llvm::detail::DenseMapPair<mlir::FunctionOpInterface, mlir::Allocation>>&) third_party/triton/lib/Analysis/Allocation.cpp:627:3 #21 0x556ee1677402 in operator() third_party/triton/include/triton/Analysis/Allocation.h:227:26 #22 0x556ee1677402 in void mlir::CallGraph<mlir::Allocation>::doWalk<(mlir::WalkOrder)0, (mlir::WalkOrder)1, mlir::ModuleAllocation::ModuleAllocation(mlir::ModuleOp)::'lambda'(mlir::CallOpInterface, mlir::FunctionOpInterface), mlir::ModuleAllocation::ModuleAllocation(mlir::ModuleOp)::'lambda'(mlir::FunctionOpInterface)>(mlir::FunctionOpInterface, llvm::DenseSet<mlir::FunctionOpInterface, llvm::DenseMapInfo<mlir::FunctionOpInterface, void>>&, mlir::ModuleAllocation::ModuleAllocation(mlir::ModuleOp)::'lambda'(mlir::CallOpInterface, mlir::FunctionOpInterface), mlir::ModuleAllocation::ModuleAllocation(mlir::ModuleOp)::'lambda'(mlir::FunctionOpInterface)) third_party/triton/include/triton/Analysis/Utility.h:350:7 #23 0x556ee16756b3 in walk<(mlir::WalkOrder)0, (mlir::WalkOrder)1, (lambda at third_party/triton/include/triton/Analysis/Allocation.h:222:9), (lambda at third_party/triton/include/triton/Analysis/Allocation.h:224:9)> third_party/triton/include/triton/Analysis/Utility.h:242:7 #24 0x556ee16756b3 in mlir::ModuleAllocation::ModuleAllocation(mlir::ModuleOp) third_party/triton/include/triton/Analysis/Allocation.h:220:5 #25 0x556ee2c2bf18 in (anonymous namespace)::AllocateSharedMemory::runOnOperation() third_party/triton/lib/Conversion/TritonGPUToLLVM/AllocateSharedMemory.cpp:26:22 ... UndefinedBehaviorSanitizer: invalid-shift-exponent third_party/triton/third_party/f2reduce/f2reduce.cpp:421:30 ``` </p> </details>

lezcano and others added 30 commits September 4, 2024 18:56

Use device fixture for runtime/test_cache.py and `runtime/test_la…

ca3fb5f

…unch.py` (#4641) Signed-off-by: Anatoly Myachev <anatoly.myachev@intel.com>

[BACKEND] Relax layout supported by SplitOp (#4653)

2f6cec9

Remove the restriction that the split dim needs to be the fastest moving one. As long as all the registers are within a thread we can implement splitOp as a no-op. This allows more layout propagation.

[NFC] Remove dead code for hex==True branch in getFormatSubstr fu…

fc68eae

…nction (#4656) Signed-off-by: Anatoly Myachev <anatoly.myachev@intel.com>

[NFC] Use const auto& instead of auto in `bin/triton-tensor-layou…

4a56c0e

…t.cpp` file (#4657) Signed-off-by: Anatoly Myachev <anatoly.myachev@intel.com>

[TEST] Use device fixture for assert_helper.py and `print_helper.…

c902145

…py` (#4643) Signed-off-by: Anatoly Myachev <anatoly.myachev@intel.com>

[Pipeliner] Properly fail instead of assert if cannot predicate op (#…

96f091d

…4659) Mirrored from upstream change: llvm/llvm-project#107442

[frontend] Warn on usage of fp8e4b15 on Hopper rather than error (#4660)

759275d

As per title

[NFC] Use const auto& instead of auto in `ConvertLayoutOpToLLVM.c…

a4a7043

…pp` and `RewriteTensorPointer.cpp` (#4658) Signed-off-by: Anatoly Myachev <anatoly.myachev@intel.com>

[AMD] Enable dynamic peeling for stream pipeliner (#4650)

79174e7

* Added predication for tt.store for tests * Moved ttng.warp_group_dot for tests

[BACKEND] Update llvm to llvm/llvm-project@c08c6a71cfc5 (#4639)

9414165

Bumping llvm to include a loop unroller fix: https://github.com/llvm/llvm-project/pull/106164/files . This is needed for subsequent loop unroller upstreaming work.

[Frontend] Warn on implicit casting in the condition in tl.where (#4646)

b933f0f

Explicit is better than implicit

Revert "[AMD] Disable block merging to avoid block argument explosion…

3104056

… (#4176)" (#4631) Turn back aggressive strategy by default to enable block merging given now we have upstream fixes for it brought in: triton-lang/triton#4619. This reverts commit cf2ad02.

[Frontend] Improve the error when doing tensor[uint32] * -3 (#4667)

0557336

as per title

[DOC] Move sections in triton-semantics.rst from ==== to ---- (#4677)

8e2d90a

Before, the ===== sections were being shown in the sidebar of the docs. ![image](https://github.com/user-attachments/assets/deb9bef0-f8d2-4152-8d17-e009495d75e9)

[TMA] Increment refcount for Py_None in TMA descriptor runtime helper…

4b34a20

…s (#4679) Incrementing refcount for Py_None in TMA descriptor runtime helpers to avoid deallocating the None object. This fixes a runtime failure like Fatal Python error: none_dealloc: deallocating None

[BACKEND] Fix potential bug in membar TMA rules (#4681)

a63a477

Also enable the special rules by default in lowering. Problem caught by @Jokeren during code review.

zhzhcookie and others added 25 commits March 26, 2025 17:21

[BUILD] Fix ext_sourcedir in editable_wheel mode, and update require…

23d2857

…ments (#5) [BUILD] Fix ext_sourcedir in editable_wheel mode [BUILD] Update requirements

[CI] Fix build and test workflow (#6)

b382dd8

Fix build and test workflow, merge build and test jobs to one.

Revert "AMD requested cherry-picks for release/3.1.x (#4794)"

a27316f

This reverts commit cf34004.

Revert "[RELEASE] Cherry-pick use of device-agnostic `DeviceInterface…

a69557d

…` in `do_bench` (#4733)" This reverts commit 566b63c.

Revert "Cherry Pick of #4247 to release/3.1.x branch (#4706)"

fc9e223

This reverts commit fc879bd.

Revert "Cherry Pick of #4311 to release/3.1.x branch (#4705)"

7a9e367

This reverts commit 23fefa6.

Revert "[Release 3.1.0] Bump version to 3.1.0 and remove temporary ch…

1c3ee97

…ange (#4696)" This reverts commit 5fe38ff.

Revert "[AMD] Cherry-pick commits from mainline to support Flex atten…

af32f8a

…tion on AMD gpus (#4569)" This reverts commit 757b6a6.

Revert "Repack wheels with build-number 1 (#4354)"

7b5fffc

This reverts commit b514024.

Revert "[Runtime] Dynamically load cuTensorMapEncodeTiled (#4330) (#4…

6ea9e5f

…339)" This reverts commit 91f24d8.

Revert "[VERSION] Bumped to 3.0.0-post1 (#4340)"

13b1932

This reverts commit 5cbb516.

Revert "[RELEASE][BACKEND] Fix getThreadsPerWarp for MD sliced encodi…

13a139b

…ng (#4061) (#4223)" This reverts commit 107465b.

Revert "[RELEASE] [AMD] Additional AMD cherry-picks (#4175)"

d7e6ce2

This reverts commit 21eae95.

Revert "[BACKEND] Prevent for/yield argument number drift (#4097) (#4…

e0cd80d

…164)" This reverts commit dedb7bd.

Revert "[AMD] Add more math functions in libdevice (#4086) (#4163)"

a0ac7d8

This reverts commit 5d77ae2.

Revert "Remove redundant options from passes (#4015) (#4162)"

efea94b

This reverts commit dcc22b8.

Revert "[AMD] Move MFMA shortcut check to not compute scratch buffer …

d5c6092

…shape if it is not needed (#4161)" This reverts commit 3d1da66.

Revert "Use dictionary unpacking instead of merge operator for python…

3e3afef

… 3.8 compatibility (#4160)" This reverts commit 45b0f6b.

Revert "[BUILD] Change hatchet as an extra dependency (#4138) (#4158)"

e693ac5

This reverts commit 085b010.

Merge remote-tracking branch 'upstream/release/3.2.x' into triton_v3.2.x

b65638a

[BUILD] Fix merge Triton 3.2.x, disable triton_shared temporarily

72053b3

[BUILD] Update build ascend backend (#10)

3a72726

[DOC] Update readme (#14)

97e653a

[init refactor] refactor setup

ae40993

Galaxy1458 closed this Jun 12, 2025

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

Code refactoring #18

Code refactoring #18

Uh oh!

Galaxy1458 commented Jun 12, 2025

Uh oh!

CLAassistant commented Jun 12, 2025 •

edited

Loading

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

63 participants

Code refactoring #18

Code refactoring #18

Uh oh!

Conversation

Galaxy1458 commented Jun 12, 2025

Uh oh!

CLAassistant commented Jun 12, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

63 participants

CLAassistant commented Jun 12, 2025 •

edited

Loading