Skip to content

Commit 5529b87

Browse files
Workaround distribute-to-warps.mlir failure
Signed-off-by: Whitney Tsang <[email protected]>
1 parent 2beff56 commit 5529b87

File tree

2 files changed

+6
-1
lines changed

2 files changed

+6
-1
lines changed

lib/Dialect/TritonGPU/IR/Dialect.cpp

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1046,6 +1046,10 @@ SmallVector<unsigned> DotOperandEncodingAttr::getCTASplitNum() const {
10461046
SmallVector<unsigned> DotOperandEncodingAttr::getWarpsPerCTA() const {
10471047
auto distributedLayout = mlir::cast<DistributedEncodingTrait>(getParent());
10481048
auto warps = distributedLayout.getWarpsPerCTA();
1049+
// FIXME: This is a temporary solution to avoid distribute-to-warps.mlir
1050+
// failure.
1051+
if (mlir::triton::tools::getBoolEnv("TRITON_INTEL_ADVANCED_PATH"))
1052+
return warps;
10491053
auto rank = warps.size();
10501054
auto kDim = getOpIdx() == 0 ? rank - 1 : rank - 2;
10511055
warps[kDim] = 1;

test/TritonIntelGPU/distribute-to-warps.mlir

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,5 @@
1-
// RUN: triton-opt %s -split-input-file -tritonintelgpu-distribute-to-warps | FileCheck %s
1+
// FIXME: Remove the env. variable once issue #2529 is fixed.
2+
// RUN: env TRITON_INTEL_ADVANCED_PATH=1 triton-opt %s -split-input-file -tritonintelgpu-distribute-to-warps | FileCheck %s
23

34
#blocked1 = #triton_gpu.blocked<{sizePerThread = [32, 32], threadsPerWarp = [1, 1], warpsPerCTA = [4, 1], order = [1, 0], CTAsPerCGA = [1, 1], CTASplitNum = [1, 1], CTAOrder = [0, 1]}>
45
#blocked2 = #triton_gpu.blocked<{sizePerThread = [32, 32], threadsPerWarp = [1, 1], warpsPerCTA = [1, 4], order = [1, 0], CTAsPerCGA = [1, 1], CTASplitNum = [1, 1], CTAOrder = [0, 1]}>

0 commit comments

Comments
 (0)