Skip to content

Commit fdac594

Browse files
antiagainstantiagainst
authored
This updates LLVM to pull in two fixes we need for AMD: * llvm/llvm-project#110553 * llvm/llvm-project#104743 Fixed `LLVM::CallOp` and `LLVM::CallIntrinsicOp` builder API after * llvm/llvm-project#108933
1 parent 5f9bb95 commit fdac594

File tree

12 files changed

+61
-40
lines changed

12 files changed

+61
-40
lines changed

‎cmake/llvm-hash.txt‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1 +1 @@
1-
df0864e761107b07e38f5503e0cbee0cebb4c5e8
1+
61f8a7f618901797ee8663389a29722f29216a96

‎include/triton/Conversion/TritonGPUToLLVM/Utility.h‎

Lines changed: 7 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -101,7 +101,7 @@ using namespace mlir::triton;
101101
#define barrier() rewriter.create<mlir::gpu::BarrierOp>(loc)
102102
#define undef(...) rewriter.create<LLVM::UndefOp>(loc, __VA_ARGS__)
103103
#define null(...) rewriter.create<LLVM::ZeroOp>(loc, __VA_ARGS__)
104-
#define call(...) rewriter.create<LLVM::CallOp>(loc, __VA_ARGS__)
104+
#define call(...) LLVM::createLLVMCallOp(rewriter, loc, __VA_ARGS__)
105105

106106
// Types
107107
#define int_ty(width) rewriter.getIntegerType(width)
@@ -228,6 +228,12 @@ Value createIndexConstant(OpBuilder &builder, Location loc,
228228
Value createLLVMIntegerConstant(OpBuilder &builder, Location loc, short width,
229229
int64_t value);
230230

231+
LLVM::CallOp createLLVMCallOp(OpBuilder &builder, Location loc,
232+
LLVMFuncOp funcOp, ValueRange args);
233+
LLVM::CallIntrinsicOp
234+
createLLVMIntrinsicCallOp(OpBuilder &builder, Location loc, StringRef intrinsic,
235+
TypeRange types, ValueRange args);
236+
231237
// Is v an integer or floating-point scalar constant equal to 0?
232238
bool isConstantZero(Value v);
233239

‎lib/Conversion/TritonGPUToLLVM/ControlFlowOpToLLVM.cpp‎

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -109,6 +109,10 @@ struct CallOpConversion : public ConvertOpToLLVMPattern<triton::CallOp> {
109109
auto newCallOp = rewriter.create<LLVM::CallOp>(
110110
callOp.getLoc(), packedResult ? TypeRange(packedResult) : TypeRange(),
111111
promotedOperands, callOp->getAttrs());
112+
newCallOp.getProperties().setOpBundleSizes(
113+
rewriter.getDenseI32ArrayAttr({}));
114+
newCallOp.getProperties().setOperandSegmentSizes(
115+
{static_cast<int>(promotedOperands.size()), 0});
112116
return newCallOp;
113117
}
114118

‎lib/Conversion/TritonGPUToLLVM/ElementwiseOpToLLVM.cpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -299,7 +299,7 @@ struct MulhiUIOpConversion
299299
LLVM::LLVMFuncOp funcOp =
300300
appendOrGetExternFuncOp(rewriter, op, funcName, funcType);
301301
return {
302-
rewriter.create<LLVM::CallOp>(loc, funcOp, operands[0]).getResult()};
302+
LLVM::createLLVMCallOp(rewriter, loc, funcOp, operands[0]).getResult()};
303303
}
304304

305305
protected:
@@ -327,7 +327,7 @@ struct ExternElementwiseOpConversion
327327
LLVM::LLVMFuncOp funcOp = appendOrGetExternFuncOp(
328328
rewriter, op, funcName, funcType, op.getLibname(), op.getLibpath());
329329
return {
330-
rewriter.create<LLVM::CallOp>(loc, funcOp, operands[0]).getResult()};
330+
LLVM::createLLVMCallOp(rewriter, loc, funcOp, operands[0]).getResult()};
331331
}
332332
};
333333

‎lib/Conversion/TritonGPUToLLVM/Utility.cpp‎

Lines changed: 19 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,8 +1,7 @@
11
#include "triton/Conversion/TritonGPUToLLVM/Utility.h"
22
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
3-
#include "mlir/Dialect/LLVMIR/NVVMDialect.h"
3+
#include "mlir/IR/Attributes.h"
44
#include "triton/Conversion/TritonGPUToLLVM/TargetInfoBase.h"
5-
#include "triton/Conversion/TritonGPUToLLVM/TypeConverter.h"
65
#include "triton/Dialect/TritonGPU/IR/Attributes.h"
76
#include "triton/Dialect/TritonGPU/IR/LinearLayoutConversions.h"
87
#include "llvm/ADT/STLExtras.h"
@@ -518,6 +517,24 @@ Value createLLVMIntegerConstant(OpBuilder &builder, Location loc, short width,
518517
builder.getIntegerAttr(ty, value));
519518
}
520519

520+
LLVM::CallOp createLLVMCallOp(OpBuilder &builder, Location loc,
521+
LLVMFuncOp funcOp, ValueRange args) {
522+
auto op = builder.create<LLVM::CallOp>(loc, funcOp, args);
523+
op.getProperties().setOpBundleSizes(builder.getDenseI32ArrayAttr({}));
524+
op.getProperties().setOperandSegmentSizes({static_cast<int>(args.size()), 0});
525+
return op;
526+
}
527+
528+
LLVM::CallIntrinsicOp
529+
createLLVMIntrinsicCallOp(OpBuilder &builder, Location loc, StringRef intrinsic,
530+
TypeRange types, ValueRange args) {
531+
auto op = builder.create<LLVM::CallIntrinsicOp>(loc, types, args);
532+
op.getProperties().setIntrin(builder.getStringAttr(intrinsic));
533+
op.getProperties().setOpBundleSizes(builder.getDenseI32ArrayAttr({}));
534+
op.getProperties().setOperandSegmentSizes({static_cast<int>(args.size()), 0});
535+
return op;
536+
}
537+
521538
bool isConstantZero(Value v) {
522539
if (auto constantOp = v.getDefiningOp<arith::ConstantOp>()) {
523540
if (auto attr = dyn_cast<IntegerAttr>(constantOp.getValue())) {

‎third_party/amd/lib/TritonAMDGPUToLLVM/BuiltinFuncToLLVM.cpp‎

Lines changed: 5 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -4,6 +4,7 @@
44
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
55
#include "mlir/Pass/Pass.h"
66
#include "mlir/Transforms/GreedyPatternRewriteDriver.h"
7+
#include "triton/Conversion/TritonGPUToLLVM/Utility.h"
78

89
namespace mlir {
910
namespace triton {
@@ -187,11 +188,11 @@ class CallOpConversion : public mlir::RewritePattern {
187188
rewriter.create<LLVM::FPToSIOp>(loc, returnType, op->getResult(0));
188189
} else if (calleeName == "__triton_hip_fast_fdividef") {
189190
assert(operands.size() == 2);
190-
auto name = StringAttr::get(callOp.getContext(), "llvm.amdgcn.rcp.f32");
191-
LLVM::FastmathFlagsAttr defaultFlags{};
192-
auto rcpOp = rewriter.create<LLVM::CallIntrinsicOp>(
193-
loc, returnType, name, operands[1], defaultFlags);
191+
const char *intrinsic = "llvm.amdgcn.rcp.f32";
192+
auto rcpOp = LLVM::createLLVMIntrinsicCallOp(rewriter, loc, intrinsic,
193+
returnType, operands[1]);
194194

195+
LLVM::FastmathFlagsAttr defaultFlags{};
195196
replacementOp = rewriter.create<LLVM::FMulOp>(
196197
loc, returnType, operands[0], rcpOp->getResult(0), defaultFlags);
197198
}

‎third_party/amd/lib/TritonAMDGPUToLLVM/DotOpToLLVM/WMMA.cpp‎

Lines changed: 3 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -24,6 +24,7 @@
2424
#include "../PatternTritonGPUOpToLLVM.h"
2525
#include "Utility.h"
2626
#include "mlir/Dialect/LLVMIR/ROCDLDialect.h"
27+
#include "triton/Conversion/TritonGPUToLLVM/Utility.h"
2728

2829
namespace mlir::triton::AMD {
2930
namespace {
@@ -219,10 +220,8 @@ Value generateWMMAIntrinsic(ConversionPatternRewriter &rewriter, Location loc,
219220
if (32 / dElType.getIntOrFloatBitWidth() > 1 || dElType.isInteger(32)) {
220221
operands.push_back(int_val(1, false));
221222
}
222-
auto wmmaIntrinsic = rewriter.create<mlir::LLVM::CallIntrinsicOp>(
223-
loc, TypeRange{valC.getType()}, StringAttr::get(loc.getContext(), name),
224-
operands, defaultFlags);
225-
223+
auto wmmaIntrinsic = LLVM::createLLVMIntrinsicCallOp(
224+
rewriter, loc, name, valC.getType(), operands);
226225
return wmmaIntrinsic.getResult(0);
227226
}
228227

‎third_party/amd/lib/TritonAMDGPUToLLVM/ElementwiseOpToLLVM.cpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1243,7 +1243,7 @@ struct ExpOpConversionApprox
12431243
LLVM::LLVMFuncOp funcOp =
12441244
appendOrGetExternFuncOp(rewriter, op, funcName, funcType);
12451245

1246-
return {rewriter.create<LLVM::CallOp>(loc, funcOp, prod).getResult()};
1246+
return {LLVM::createLLVMCallOp(rewriter, loc, funcOp, prod).getResult()};
12471247
}
12481248
};
12491249

@@ -1276,7 +1276,7 @@ struct Exp2OpConversion
12761276
appendOrGetExternFuncOp(rewriter, op, funcName, funcType);
12771277

12781278
return {
1279-
rewriter.create<LLVM::CallOp>(loc, funcOp, operands[0]).getResult()};
1279+
LLVM::createLLVMCallOp(rewriter, loc, funcOp, operands[0]).getResult()};
12801280
}
12811281

12821282
private:

‎third_party/amd/lib/TritonAMDGPUToLLVM/SchedInstructions.cpp‎

Lines changed: 9 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -38,7 +38,7 @@ void createSchedGroupBarrier(PatternRewriter &rewriter, Location loc,
3838
InstructionKindMask maskValue, int sizeValue,
3939
int groupIdValue) {
4040
MLIRContext *ctx = rewriter.getContext();
41-
auto intrinsicName = str_attr("llvm.amdgcn.sched.group.barrier");
41+
const char *intrinsicName = "llvm.amdgcn.sched.group.barrier";
4242

4343
Value mask =
4444
LLVM::createConstantI32(loc, rewriter, static_cast<int32_t>(maskValue));
@@ -47,36 +47,34 @@ void createSchedGroupBarrier(PatternRewriter &rewriter, Location loc,
4747
Value groupId = LLVM::createConstantI32(loc, rewriter,
4848
static_cast<int32_t>(groupIdValue));
4949

50-
LLVM::FastmathFlagsAttr defaultFlags{};
51-
rewriter.create<LLVM::CallIntrinsicOp>(loc, TypeRange{}, intrinsicName,
52-
ValueRange{mask, size, groupId},
53-
defaultFlags);
50+
LLVM::createLLVMIntrinsicCallOp(rewriter, loc, intrinsicName, TypeRange{},
51+
ValueRange{mask, size, groupId});
5452
}
5553

5654
// Insert intrinsic that controls the types of instructions that may be
5755
// allowed to cross the intrinsic during instruction scheduling
5856
Operation *createSchedBarrier(PatternRewriter &rewriter, Location loc,
5957
int64_t maskValue) {
6058
MLIRContext *ctx = rewriter.getContext();
61-
auto intrinsicName = str_attr("llvm.amdgcn.sched.barrier");
59+
const char *intrinsicName = "llvm.amdgcn.sched.barrier";
6260
LLVM::FastmathFlagsAttr defaultFlags{};
6361

6462
Value mask =
6563
LLVM::createConstantI32(loc, rewriter, static_cast<int32_t>(maskValue));
66-
return rewriter.create<LLVM::CallIntrinsicOp>(loc, TypeRange{}, intrinsicName,
67-
ValueRange{mask}, defaultFlags);
64+
return LLVM::createLLVMIntrinsicCallOp(rewriter, loc, intrinsicName,
65+
TypeRange{}, ValueRange{mask});
6866
}
6967

7068
// Insert an experimental intrinsic for instruction group level parallelism.
7169
// The intrinsic takes a value that specifies the strategy.
7270
Operation *createIglpOpt(PatternRewriter &rewriter, Location loc, int value) {
7371
MLIRContext *ctx = rewriter.getContext();
74-
auto intrinsicName = str_attr("llvm.amdgcn.iglp.opt");
72+
const char *intrinsicName = "llvm.amdgcn.iglp.opt";
7573
LLVM::FastmathFlagsAttr defaultFlags{};
7674
Value iglpValue =
7775
LLVM::createConstantI32(loc, rewriter, static_cast<int32_t>(value));
78-
return rewriter.create<LLVM::CallIntrinsicOp>(
79-
loc, TypeRange{}, intrinsicName, ValueRange{iglpValue}, defaultFlags);
76+
return LLVM::createLLVMIntrinsicCallOp(rewriter, loc, intrinsicName,
77+
TypeRange{}, ValueRange{iglpValue});
8078
}
8179

8280
struct InstructionSchedHintsRewriter

‎third_party/amd/lib/TritonAMDGPUToLLVM/TargetInfo.cpp‎

Lines changed: 3 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -69,12 +69,9 @@ Value TargetInfo::getClusterCTAId(RewriterBase &rewriter, Location loc) const {
6969

7070
Value TargetInfo::ballot(RewriterBase &rewriter, Location loc, Type type,
7171
Value cmp) const {
72-
auto stringAttr = rewriter.getStringAttr("llvm.amdgcn.ballot");
73-
SmallVector<Value> operands = {cmp};
74-
Value asmResult =
75-
rewriter.create<LLVM::CallIntrinsicOp>(loc, type, stringAttr, operands)
76-
->getResult(0);
77-
return asmResult;
72+
return LLVM::createLLVMIntrinsicCallOp(rewriter, loc, "llvm.amdgcn.ballot",
73+
type, cmp)
74+
->getResult(0);
7875
}
7976

8077
void TargetInfo::storeDShared(RewriterBase &rewriter, Location loc, Value ptr,

0 commit comments

Comments
 (0)