Skip to content

Commit d415da6

Browse files
whitneywhtsangwhitneywhtsang
authored
Merge OpenAI Triton commit acd8104 (#5591)
This PR changes the Triton base from 29009f1 to acd8104 (Nov 22). Pass rate: 95.42%
2 parents 0bf934b + 97ccbdb commit d415da6

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

53 files changed

+409
-613
lines changed

‎cmake/llvm-hash.txt‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1 +1 @@
1-
f6ded0be897e2878612dd903f7e8bb85448269e5
1+
49d5bb0ad0cb31410184c462801c5049ad671517

‎lib/Dialect/TritonGPU/IR/Ops.cpp‎

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -931,8 +931,9 @@ void WarpSpecializeOp::getSuccessorRegions(
931931
return;
932932
}
933933
// And the default region branches transparently back to the parent.
934-
assert(src.getRegionOrNull() == &getDefaultRegion());
935-
successors.push_back(RegionSuccessor(getResults()));
934+
assert(src.getTerminatorPredecessorOrNull()->getParentRegion() ==
935+
&getDefaultRegion());
936+
successors.push_back(RegionSuccessor(getOperation(), getResults()));
936937
}
937938

938939
LogicalResult WarpSpecializeOp::verify() {

‎lib/Dialect/TritonGPU/Transforms/WarpSpecialization/PartitionLoops.cpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -338,7 +338,7 @@ void cloneOpsInBlock(Block *block, SmallVector<WarpGroupBuilder> &builders,
338338
builder.mapping.lookupOrDefault(yieldOp.getOperand(i)));
339339
}
340340

341-
builder.create<scf::YieldOp>(op->getLoc(), newYieldOperands);
341+
scf::YieldOp::create(builder, op->getLoc(), newYieldOperands);
342342
}
343343
} else {
344344
assert(hasPartition(op));
@@ -449,7 +449,7 @@ LogicalResult triton::gpu::partitionLoop(scf::ForOp loop) {
449449
for (auto [b, region, partition] : llvm::zip(
450450
builders, wgOp.getPartitionRegions(), partitions.getPartitions())) {
451451
if (!llvm::is_contained(getPartitionIds(loop), b.partitionId)) {
452-
b.create<nvws::WarpGroupYieldOp>(wgOp.getLoc(), SmallVector<Value>{});
452+
nvws::WarpGroupYieldOp::create(b, wgOp.getLoc(), SmallVector<Value>{});
453453
continue;
454454
}
455455
auto newForOp = *region.front().getOps<scf::ForOp>().begin();

‎python/src/llvm.cc‎

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -57,7 +57,6 @@ createTargetMachine(llvm::Module *module, std::string proc,
5757
bool disableLLVMOpt = mlir::triton::tools::getBoolEnv("DISABLE_LLVM_OPT");
5858
if (enable_fp_fusion)
5959
opt.AllowFPOpFusion = llvm::FPOpFusion::Fast;
60-
opt.UnsafeFPMath = false;
6160
opt.NoInfsFPMath = false;
6261
opt.NoNaNsFPMath = true;
6362
opt.TrapUnreachable = true;

‎test/Conversion/tritonnvidiagpu_to_llvm.mlir‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -215,9 +215,9 @@ module attributes {"ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 4 : i32, ttg.targ
215215
// CHECK-LABEL: async_copy_mbarrier_arrive
216216
module attributes {"ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 4 : i32, ttg.target = "cuda:90", "ttg.threads-per-warp" = 32 : i32} {
217217
tt.func public @async_copy_mbarrier_arrive(%arg0: !ttg.memdesc<1xi64, #shared, #ttg.shared_memory>) attributes { noinline = false } {
218-
// CHECK: nvvm.cp.async.mbarrier.arrive.shared %{{.*}} : !llvm.ptr<3>
218+
// CHECK: nvvm.cp.async.mbarrier.arrive %{{.*}} : !llvm.ptr<3>
219219
ttng.async_copy_mbarrier_arrive %arg0 : !ttg.memdesc<1xi64, #shared, #ttg.shared_memory>
220-
// CHECK: nvvm.cp.async.mbarrier.arrive.shared %{{.*}} {noinc = true} : !llvm.ptr<3>
220+
// CHECK: nvvm.cp.async.mbarrier.arrive %{{.*}} {noinc = true} : !llvm.ptr<3>
221221
ttng.async_copy_mbarrier_arrive %arg0 { noIncrement } : !ttg.memdesc<1xi64, #shared, #ttg.shared_memory>
222222
tt.return
223223
}

‎third_party/amd/include/Analysis/RangeAnalysis.h‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -84,7 +84,7 @@ struct TritonIntegerRangeAnalysis : dataflow::IntegerRangeAnalysis {
8484
/// the loop operands and all users and all users of the results of the loop.
8585
void visitRegionSuccessors(
8686
ProgramPoint *point, RegionBranchOpInterface branch,
87-
RegionBranchPoint successor,
87+
RegionSuccessor successor,
8888
ArrayRef<dataflow::AbstractSparseLattice *> abstractLattices) override;
8989

9090
/// Collect all operands that participate in assumptions (see description of

‎third_party/amd/lib/Analysis/RangeAnalysis.cpp‎

Lines changed: 6 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -630,7 +630,7 @@ void TritonIntegerRangeAnalysis::initializeFuncOp(tt::FuncOp op) {
630630

631631
void TritonIntegerRangeAnalysis::visitRegionSuccessors(
632632
ProgramPoint *point, RegionBranchOpInterface branch,
633-
RegionBranchPoint successor,
633+
RegionSuccessor successor,
634634
ArrayRef<dataflow::AbstractSparseLattice *> abstractLattices) {
635635
LLVM_DEBUG({
636636
DBGS() << "Visit Region Succesors of ";
@@ -715,10 +715,11 @@ void TritonIntegerRangeAnalysis::visitRegionSuccessors(
715715
if (!inputs.empty()) {
716716
firstIndex = cast<OpResult>(inputs.front()).getResultNumber();
717717
}
718-
visitNonControlFlowArguments(branch,
719-
RegionSuccessor(branch->getResults().slice(
720-
firstIndex, inputs.size())),
721-
lattices, firstIndex);
718+
visitNonControlFlowArguments(
719+
branch,
720+
RegionSuccessor(
721+
branch, branch->getResults().slice(firstIndex, inputs.size())),
722+
lattices, firstIndex);
722723
} else {
723724
if (!inputs.empty()) {
724725
firstIndex = cast<BlockArgument>(inputs.front()).getArgNumber();

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -481,8 +481,8 @@ class LocalBarrierOpConversion
481481
// amdgpu::MemoryCounterWaitOp will lower s_waitcnt
482482
// - s_barrier syncronizes the execution for the CTA
483483
auto dsAttr = rewriter.getI32IntegerAttr(0);
484-
rewriter.create<amdgpu::MemoryCounterWaitOp>(
485-
op->getLoc(), /* load= */ nullptr, /* store= */ nullptr,
484+
amdgpu::MemoryCounterWaitOp::create(
485+
rewriter, op->getLoc(), /* load= */ nullptr, /* store= */ nullptr,
486486
/* ds= */ dsAttr);
487487
rewriter.replaceOpWithNewOp<ROCDL::SBarrierOp>(op);
488488

‎third_party/intel/cmake/FindSPIRVToLLVMTranslator.cmake‎

Lines changed: 0 additions & 70 deletions
Original file line numberDiff line numberDiff line change
@@ -75,76 +75,6 @@ if (NOT SPIRVToLLVMTranslator_FOUND)
7575
if(NOT PATCH_RESULT EQUAL 0)
7676
message(FATAL_ERROR "Failed to apply 3388.patch to SPIRV-LLVM-Translator")
7777
endif()
78-
79-
# FIXME: Don't apply patch when LLVM commit update to 6cba572.
80-
execute_process(
81-
COMMAND git apply --check ${CMAKE_CURRENT_LIST_DIR}/revert_3385.patch
82-
WORKING_DIRECTORY ${spirv-llvm-translator_SOURCE_DIR}
83-
ERROR_QUIET
84-
RESULT_VARIABLE PATCH_RESULT
85-
)
86-
if(PATCH_RESULT EQUAL 0)
87-
execute_process(
88-
COMMAND git apply ${CMAKE_CURRENT_LIST_DIR}/revert_3385.patch
89-
WORKING_DIRECTORY ${spirv-llvm-translator_SOURCE_DIR}
90-
RESULT_VARIABLE PATCH_RESULT
91-
)
92-
else()
93-
execute_process( # Check if the patch is already applied
94-
COMMAND git apply --reverse --check ${CMAKE_CURRENT_LIST_DIR}/revert_3385.patch
95-
WORKING_DIRECTORY ${spirv-llvm-translator_SOURCE_DIR}
96-
RESULT_VARIABLE PATCH_RESULT
97-
)
98-
endif()
99-
if(NOT PATCH_RESULT EQUAL 0)
100-
message(FATAL_ERROR "Failed to apply revert_3385.patch to SPIRV-LLVM-Translator")
101-
endif()
102-
103-
# FIXME: Don't apply patch when LLVM commit update to 573ca36.
104-
execute_process(
105-
COMMAND git apply --check ${CMAKE_CURRENT_LIST_DIR}/revert_3406.patch
106-
WORKING_DIRECTORY ${spirv-llvm-translator_SOURCE_DIR}
107-
ERROR_QUIET
108-
RESULT_VARIABLE PATCH_RESULT
109-
)
110-
if(PATCH_RESULT EQUAL 0)
111-
execute_process(
112-
COMMAND git apply ${CMAKE_CURRENT_LIST_DIR}/revert_3406.patch
113-
WORKING_DIRECTORY ${spirv-llvm-translator_SOURCE_DIR}
114-
RESULT_VARIABLE PATCH_RESULT
115-
)
116-
else()
117-
execute_process( # Check if the patch is already applied
118-
COMMAND git apply --reverse --check ${CMAKE_CURRENT_LIST_DIR}/revert_3406.patch
119-
WORKING_DIRECTORY ${spirv-llvm-translator_SOURCE_DIR}
120-
RESULT_VARIABLE PATCH_RESULT
121-
)
122-
endif()
123-
if(NOT PATCH_RESULT EQUAL 0)
124-
message(FATAL_ERROR "Failed to apply revert_3406.patch to SPIRV-LLVM-Translator")
125-
endif()
126-
execute_process(
127-
COMMAND git apply --check ${CMAKE_CURRENT_LIST_DIR}/revert_3407.patch
128-
WORKING_DIRECTORY ${spirv-llvm-translator_SOURCE_DIR}
129-
ERROR_QUIET
130-
RESULT_VARIABLE PATCH_RESULT
131-
)
132-
if(PATCH_RESULT EQUAL 0)
133-
execute_process(
134-
COMMAND git apply ${CMAKE_CURRENT_LIST_DIR}/revert_3407.patch
135-
WORKING_DIRECTORY ${spirv-llvm-translator_SOURCE_DIR}
136-
RESULT_VARIABLE PATCH_RESULT
137-
)
138-
else()
139-
execute_process( # Check if the patch is already applied
140-
COMMAND git apply --reverse --check ${CMAKE_CURRENT_LIST_DIR}/revert_3407.patch
141-
WORKING_DIRECTORY ${spirv-llvm-translator_SOURCE_DIR}
142-
RESULT_VARIABLE PATCH_RESULT
143-
)
144-
endif()
145-
if(NOT PATCH_RESULT EQUAL 0)
146-
message(FATAL_ERROR "Failed to apply revert_3407.patch to SPIRV-LLVM-Translator")
147-
endif()
14878
endif()
14979

15080
set(SPIRVToLLVMTranslator_INCLUDE_DIR "${SPIRVToLLVMTranslator_SOURCE_DIR}/include"

‎third_party/intel/cmake/revert_3385.patch‎

Lines changed: 0 additions & 15 deletions
This file was deleted.

0 commit comments

Comments
 (0)