Skip to content

Commit c061685

Browse files
committed
[AMD] Enable supportBitwidth{16|32}Elementwise in TargetInfo
This helps to optimize reduction code generation.
1 parent 917dbde commit c061685

2 files changed

Lines changed: 9 additions & 0 deletions

File tree

third_party/amd/lib/TritonAMDGPUToLLVM/TargetInfo.cpp

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -675,6 +675,12 @@ bool TargetInfo::supportVectorizedAtomics() const {
675675
return true;
676676
}
677677

678+
bool TargetInfo::supportBitwidth16Elementwise() const { return true; }
679+
680+
bool TargetInfo::supportBitwidth32Elementwise() const {
681+
return getISAFamily() == ISAFamily::GFX1250;
682+
}
683+
678684
bool TargetInfo::supportsDirectToLDSScattering() const {
679685
switch (getISAFamily()) {
680686
case ISAFamily::GFX1250:

third_party/amd/lib/TritonAMDGPUToLLVM/TargetInfo.h

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -98,6 +98,9 @@ class TargetInfo : public mlir::triton::TargetInfoBase {
9898

9999
bool supportVectorizedAtomics() const override;
100100

101+
bool supportBitwidth16Elementwise() const override;
102+
bool supportBitwidth32Elementwise() const override;
103+
101104
// Returns true if the target supports per lane addresses into LDS for
102105
// direct-to-lds loads. Some architectures (e.g. GFX9) do not support
103106
// scattering and instead have to write warp coalesced into LDS

0 commit comments

Comments
 (0)