1
1
#include " intel/include/Dialect/TritonIntelGPU/IR/Dialect.h"
2
2
#include " intel/include/Dialect/TritonIntelGPU/Transforms/Passes.h"
3
3
#include " triton/Dialect/Triton/IR/Utility.h"
4
+ #include " triton/Dialect/TritonGPU/Transforms/Utility.h"
4
5
#include " llvm/ADT/PriorityWorklist.h"
5
6
6
7
namespace ttg = mlir::triton::gpu;
@@ -16,38 +17,6 @@ namespace gpu::intel {
16
17
17
18
namespace {
18
19
19
- SmallVector<Value> getTiedArgs (Operation *op, int resultIdx) {
20
- if (auto forOp = dyn_cast<scf::ForOp>(op)) {
21
- auto iterArg = forOp.getRegionIterArg (resultIdx);
22
- auto result = forOp.getResult (resultIdx);
23
- auto yieldVal = forOp.getBody ()->getTerminator ()->getOperand (resultIdx);
24
- auto initVal = forOp.getInitArgs ()[resultIdx];
25
- return {iterArg, result, yieldVal, initVal};
26
- } else if (auto whileOp = dyn_cast<scf::WhileOp>(op)) {
27
- auto iterArg = whileOp.getBeforeArguments ()[resultIdx];
28
- auto result = whileOp.getResults ()[resultIdx];
29
- auto yieldVal = whileOp.getConditionOp ().getArgs ()[resultIdx];
30
- auto initVal = whileOp.getOperands ()[resultIdx];
31
- auto bodyArg = whileOp.getAfterArguments ()[resultIdx];
32
- return {iterArg, result, yieldVal, initVal, bodyArg};
33
- } else if (auto ifOp = dyn_cast<scf::IfOp>(op)) {
34
- SmallVector<Value> values;
35
- for (auto &block : ifOp.getThenRegion ().getBlocks ()) {
36
- auto terminator = block.getTerminator ();
37
- if (isa<scf::YieldOp>(terminator))
38
- values.push_back (terminator->getOperands ()[resultIdx]);
39
- }
40
- for (auto &block : ifOp.getElseRegion ().getBlocks ()) {
41
- auto terminator = block.getTerminator ();
42
- if (isa<scf::YieldOp>(terminator))
43
- values.push_back (terminator->getOperands ()[resultIdx]);
44
- }
45
- values.push_back (ifOp->getResults ()[resultIdx]);
46
- return values;
47
- }
48
- return {};
49
- }
50
-
51
20
struct EncodingInfo {
52
21
Attribute desiredEncoding;
53
22
bool requiresConvert = false ;
0 commit comments