Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

[ROCM] Include rocm Triton build in build system. #7

Closed
wants to merge 7 commits into from
Closed

[ROCM] Include rocm Triton build in build system. #7

wants to merge 7 commits into from

Conversation

zoranjovanovic-ns
Copy link

Add Bazel files for ROCm part of Triton.
Just a minimal set of changes that should be part of the "OpenXLA-specific changes" commit.

Moerafaat and others added 7 commits March 25, 2024 10:11
@i-chaochen
Copy link

i-chaochen commented Apr 19, 2024 •
edited
Loading

To whom it may concern. This is from the series of ROCm triton enablement.

  1. [ROCm] Triton in XLA for ROCm - gemm rewriter triton related changes. xla#10510
  2. [ROCm] Triton in XLA for ROCm - ir_emitter_triton related changes. xla#10649
  3. [ROCm] Fix compilation issues related to aa08925e2dd1351c68e7838e074fe35908e3 xla#11331
  4. [ROCm] Triton in XLA for ROCm - changed ir_emitter_triton_test  xla#11352

Thanks!

@chsigg
Copy link
Member

chsigg commented Apr 22, 2024

There are a few errors in these BUILD files. I will get them sorted out and land the change. Stay tuned.

@zoranjovanovic-ns
Copy link
Author

Let me know if I need to do something in order to get this merged.

@chsigg chsigg force-pushed the llvm-head branch 3 times, most recently from 37b598d to 10c56aa Compare May 14, 2024 18:43
@khasanovaa khasanovaa force-pushed the llvm-head branch 2 times, most recently from c7feaed to 38c82bd Compare May 22, 2024 15:48
@chsigg
Copy link
Member

chsigg commented May 29, 2024

Oh dear, I dropped the ball here. Sorry about that.
Good news is that it landed internally now and will get pushed with the integrate from upstream Triton next week.

copybara-service bot pushed a commit to openxla/xla that referenced this pull request May 30, 2024
@chsigg @copybara-github
Update OpenXLA's Triton dependency to include the AMD backend.
0edfa1f 
- openxla/triton#7
- openxla/triton@25e4e02
- triton-lang/triton#4031

PiperOrigin-RevId: 638583630
@copybara-service copybara-service bot mentioned this pull request May 30, 2024
Merged
copybara-service bot pushed a commit to tensorflow/tensorflow that referenced this pull request May 30, 2024
@chsigg @tensorflower-gardener
Update OpenXLA's Triton dependency to include the AMD backend.
b4451fc 
- openxla/triton#7
- openxla/triton@25e4e02
- triton-lang/triton#4031

PiperOrigin-RevId: 638583630
@copybara-service copybara-service bot mentioned this pull request May 30, 2024
Merged
copybara-service bot pushed a commit to openxla/xla that referenced this pull request May 30, 2024
@chsigg @copybara-github
Update OpenXLA's Triton dependency to include the AMD backend.
c8a8c77 
- openxla/triton#7
- openxla/triton@25e4e02
- triton-lang/triton#4031

PiperOrigin-RevId: 638583630
copybara-service bot pushed a commit to tensorflow/tensorflow that referenced this pull request May 30, 2024
@chsigg @tensorflower-gardener
Update OpenXLA's Triton dependency to include the AMD backend.
ea157c0 
- openxla/triton#7
- openxla/triton@25e4e02
- triton-lang/triton#4031

PiperOrigin-RevId: 638583630
copybara-service bot pushed a commit to openxla/xla that referenced this pull request May 30, 2024
@chsigg @copybara-github
Update OpenXLA's Triton dependency to include the AMD backend.
75a141a 
- openxla/triton#7
- openxla/triton@25e4e02
- triton-lang/triton#4031

PiperOrigin-RevId: 638615931
copybara-service bot pushed a commit to tensorflow/tensorflow that referenced this pull request May 30, 2024
@chsigg @tensorflower-gardener
Update OpenXLA's Triton dependency to include the AMD backend.
de59727 
- openxla/triton#7
- openxla/triton@25e4e02
- triton-lang/triton#4031

PiperOrigin-RevId: 638615931
@chsigg
Copy link
Member

chsigg commented May 30, 2024

With 25e4e02, (and openxla/xla@75a141a picking it up), all relevant parts of this PR have landed.

I'm closing it. If I missed anything, please let me know. Thanks for your patience, and sorry again for dropping the ball.

@chsigg chsigg closed this May 30, 2024
gflegar added a commit that referenced this pull request Jun 24, 2024
@gflegar
Avoid underflow in f2reduce (triton-lang#4168)
7ca6d12 
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
    triton-lang#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
    triton-lang#21 0x556ee1677402 in operator() third_party/triton/include/triton/Analysis/Allocation.h:227:26
    triton-lang#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
    triton-lang#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
    triton-lang#24 0x556ee16756b3 in mlir::ModuleAllocation::ModuleAllocation(mlir::ModuleOp) third_party/triton/include/triton/Analysis/Allocation.h:220:5
    triton-lang#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>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Reviewers
No reviews
Assignees
No one assigned
Labels
None yet
Projects
None yet
Milestone
No milestone
Development

Successfully merging this pull request may close these issues.
8 participants
@zoranjovanovic-ns @i-chaochen @chsigg @Moerafaat @gflegar @khasanovaa @vwbaker @karupayun