Skip to content

Commit 80700bf

Browse files
Dewei-Wang-shwhitneywhtsang
authored andcommitted
Revert "Workaround distribute-to-warps.mlir failure"
This reverts commit 5529b87.
1 parent 24bd5a6 commit 80700bf

File tree

2 files changed

+1
-6
lines changed

2 files changed

+1
-6
lines changed

lib/Dialect/TritonGPU/IR/Dialect.cpp

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1046,10 +1046,6 @@ 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;
10531049
auto rank = warps.size();
10541050
auto kDim = getOpIdx() == 0 ? rank - 1 : rank - 2;
10551051
warps[kDim] = 1;

test/TritonIntelGPU/distribute-to-warps.mlir

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,4 @@
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
1+
// RUN: triton-opt %s -split-input-file -tritonintelgpu-distribute-to-warps | FileCheck %s
32

43
#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]}>
54
#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)