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

Support tl.histogram for CPU. #12

Merged
merged 2 commits into from
May 29, 2024
Merged

Support tl.histogram for CPU. #12

merged 2 commits into from
May 29, 2024

Conversation

ienkovich
Copy link
Collaborator

Each histogram bin is just a single value, so the implementation compares each input value with a target range and sums all comparison results.

@ienkovich
Support tl.histogram for CPU.
fc61ee7 
Signed-off-by: Ilya Enkovich <[email protected]>
@ienkovich ienkovich requested a review from minjang May 29, 2024 20:00
@ienkovich ienkovich requested a review from ptillet as a code owner May 29, 2024 20:00
minjang
minjang approved these changes May 29, 2024
View reviewed changes
Copy link
Collaborator

@minjang minjang left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Looks legit

@minjang minjang merged commit cea59f2 into triton-lang:main May 29, 2024
@ienkovich ienkovich deleted the ienkovich/cpu/histc branch May 29, 2024 21:00
minjang added a commit to minjang/triton-cpu that referenced this pull request Jun 22, 2024
@ienkovich @minjang
Support tl.histogram for CPU. (triton-lang#12)
2d92db9 
Signed-off-by: Ilya Enkovich <[email protected]>
Co-authored-by: Minjang Kim <[email protected]>
minjang pushed 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
    #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>
minjang added a commit that referenced this pull request Jun 24, 2024
@ienkovich @minjang
Support tl.histogram for CPU. (#12)
fd49147 
Signed-off-by: Ilya Enkovich <[email protected]>
Co-authored-by: Minjang Kim <[email protected]>
@vivekvpandya vivekvpandya mentioned this pull request Aug 9, 2024
Closed
Devjiu pushed a commit to Devjiu/triton-cpu that referenced this pull request Aug 13, 2024
@ienkovich @minjang @Devjiu
Support tl.histogram for CPU. (triton-lang#12)
0fd55a1 
Signed-off-by: Ilya Enkovich <[email protected]>
Co-authored-by: Minjang Kim <[email protected]>
int3 pushed a commit that referenced this pull request Aug 29, 2024
@ienkovich @minjang @int3
Support tl.histogram for CPU. (#12)
725ec53 
Signed-off-by: Ilya Enkovich <[email protected]>
Co-authored-by: Minjang Kim <[email protected]>
minjang added a commit that referenced this pull request Sep 22, 2024
@ienkovich @minjang
Support tl.histogram for CPU. (#12)
adbeccf 
Signed-off-by: Ilya Enkovich <[email protected]>
Co-authored-by: Minjang Kim <[email protected]>
minjang added a commit that referenced this pull request Oct 22, 2024
@ienkovich @minjang
Support tl.histogram for CPU. (#12)
7765c97 
Signed-off-by: Ilya Enkovich <[email protected]>
Co-authored-by: Minjang Kim <[email protected]>
minjang added a commit that referenced this pull request Oct 24, 2024
@ienkovich @minjang
Support tl.histogram for CPU. (#12)
435e750 
Signed-off-by: Ilya Enkovich <[email protected]>
Co-authored-by: Minjang Kim <[email protected]>
Devjiu pushed a commit to Devjiu/triton-cpu that referenced this pull request Nov 13, 2024
@adam-smnk @Devjiu
Matmul tutorial - K padding and dynamic K block size (triton-lang#12)
44545c6 
Adds two new optional flags to the matmul tutorial:
- K dim padding - pads input matrices into multiple
  of chosen BLOCK_SIZE_K
- dynamic K blocking - overrides set BLOCK_SIZE_K
  and adjusts it based on the input K dimension;
  input is padded if needed

The main motivation is to allow testing with larger reduction
dimension blocks without kernel lossing support for various sizes.
Padding is required to meet triton's requirement for power-of-2 sizes.
Dynamic blocking can be used to decrease reduction dimension range or
completely eliminate it.

Allowing the kernel to work on larger K blocks is also helpful for
future rewriting of GEMM into BRGEMM to ensure larger batch dimension.
Devjiu pushed a commit to Devjiu/triton-cpu that referenced this pull request Nov 13, 2024
@adam-smnk @Devjiu
Matmul tutorial - K padding and dynamic K block size (triton-lang#12)
4a93f09 
Adds two new optional flags to the matmul tutorial:
- K dim padding - pads input matrices into multiple
  of chosen BLOCK_SIZE_K
- dynamic K blocking - overrides set BLOCK_SIZE_K
  and adjusts it based on the input K dimension;
  input is padded if needed

The main motivation is to allow testing with larger reduction
dimension blocks without kernel lossing support for various sizes.
Padding is required to meet triton's requirement for power-of-2 sizes.
Dynamic blocking can be used to decrease reduction dimension range or
completely eliminate it.

Allowing the kernel to work on larger K blocks is also helpful for
future rewriting of GEMM into BRGEMM to ensure larger batch dimension.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Reviewers

@minjang minjang minjang approved these changes

@ptillet ptillet Awaiting requested review from ptillet ptillet is a code owner

Assignees
No one assigned
Labels
None yet
Projects
None yet
Milestone
No milestone
Development

Successfully merging this pull request may close these issues.
2 participants
@ienkovich @minjang