Skip to content

Commit a962596

Browse files
ThomasRaouxThomasRaoux
committed
This reverts commit fdac594.
1 parent ca70f08 commit a962596

File tree

12 files changed

+40
-61
lines changed

12 files changed

+40
-61
lines changed

‎cmake/llvm-hash.txt‎

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

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

Lines changed: 1 addition & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -102,7 +102,7 @@ using namespace mlir::triton;
102102
#define barrier() rewriter.create<mlir::gpu::BarrierOp>(loc)
103103
#define undef(...) rewriter.create<LLVM::UndefOp>(loc, __VA_ARGS__)
104104
#define null(...) rewriter.create<LLVM::ZeroOp>(loc, __VA_ARGS__)
105-
#define call(...) LLVM::createLLVMCallOp(rewriter, loc, __VA_ARGS__)
105+
#define call(...) rewriter.create<LLVM::CallOp>(loc, __VA_ARGS__)
106106

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

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

‎lib/Conversion/TritonGPUToLLVM/ControlFlowOpToLLVM.cpp‎

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -109,10 +109,6 @@ 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});
116112
return newCallOp;
117113
}
118114

‎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-
LLVM::createLLVMCallOp(rewriter, loc, funcOp, operands[0]).getResult()};
302+
rewriter.create<LLVM::CallOp>(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-
LLVM::createLLVMCallOp(rewriter, loc, funcOp, operands[0]).getResult()};
330+
rewriter.create<LLVM::CallOp>(loc, funcOp, operands[0]).getResult()};
331331
}
332332
};
333333

‎lib/Conversion/TritonGPUToLLVM/Utility.cpp‎

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

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-
538521
bool isConstantZero(Value v) {
539522
if (auto constantOp = v.getDefiningOp<arith::ConstantOp>()) {
540523
if (auto attr = dyn_cast<IntegerAttr>(constantOp.getValue())) {

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

Lines changed: 4 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -4,7 +4,6 @@
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"
87

98
namespace mlir {
109
namespace triton {
@@ -188,11 +187,11 @@ class CallOpConversion : public mlir::RewritePattern {
188187
rewriter.create<LLVM::FPToSIOp>(loc, returnType, op->getResult(0));
189188
} else if (calleeName == "__triton_hip_fast_fdividef") {
190189
assert(operands.size() == 2);
191-
const char *intrinsic = "llvm.amdgcn.rcp.f32";
192-
auto rcpOp = LLVM::createLLVMIntrinsicCallOp(rewriter, loc, intrinsic,
193-
returnType, operands[1]);
194-
190+
auto name = StringAttr::get(callOp.getContext(), "llvm.amdgcn.rcp.f32");
195191
LLVM::FastmathFlagsAttr defaultFlags{};
192+
auto rcpOp = rewriter.create<LLVM::CallIntrinsicOp>(
193+
loc, returnType, name, operands[1], defaultFlags);
194+
196195
replacementOp = rewriter.create<LLVM::FMulOp>(
197196
loc, returnType, operands[0], rcpOp->getResult(0), defaultFlags);
198197
}

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

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

2928
namespace mlir::triton::AMD {
3029
namespace {
@@ -220,8 +219,10 @@ Value generateWMMAIntrinsic(ConversionPatternRewriter &rewriter, Location loc,
220219
if (32 / dElType.getIntOrFloatBitWidth() > 1 || dElType.isInteger(32)) {
221220
operands.push_back(int_val(1, false));
222221
}
223-
auto wmmaIntrinsic = LLVM::createLLVMIntrinsicCallOp(
224-
rewriter, loc, name, valC.getType(), operands);
222+
auto wmmaIntrinsic = rewriter.create<mlir::LLVM::CallIntrinsicOp>(
223+
loc, TypeRange{valC.getType()}, StringAttr::get(loc.getContext(), name),
224+
operands, defaultFlags);
225+
225226
return wmmaIntrinsic.getResult(0);
226227
}
227228

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

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

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

12821282
private:

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

Lines changed: 11 additions & 9 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-
const char *intrinsicName = "llvm.amdgcn.sched.group.barrier";
41+
auto intrinsicName = str_attr("llvm.amdgcn.sched.group.barrier");
4242

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

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

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

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

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

8082
struct InstructionSchedHintsRewriter

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

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

7070
Value TargetInfo::ballot(RewriterBase &rewriter, Location loc, Type type,
7171
Value cmp) const {
72-
return LLVM::createLLVMIntrinsicCallOp(rewriter, loc, "llvm.amdgcn.ballot",
73-
type, cmp)
74-
->getResult(0);
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;
7578
}
7679

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

0 commit comments

Comments
 (0)