11#include " intel/include/Dialect/TritonIntelGPU/IR/Dialect.h"
22#include " intel/include/Dialect/TritonIntelGPU/Transforms/Passes.h"
33#include " triton/Dialect/Triton/IR/Utility.h"
4+ #include " triton/Dialect/TritonGPU/Transforms/Utility.h"
45#include " llvm/ADT/PriorityWorklist.h"
56
67namespace ttg = mlir::triton::gpu;
@@ -16,38 +17,6 @@ namespace gpu::intel {
1617
1718namespace {
1819
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-
5120struct EncodingInfo {
5221 Attribute desiredEncoding;
5322 bool requiresConvert = false ;
0 commit comments