Skip to content

Commit 282c825

Browse files
authored
Revert "[Backend] Bump to llvm/llvm-project@979132a (#9431) (#9477)" (#9942)
Please see issue: pytorch/pytorch#178554 (comment)
1 parent 9c288bc commit 282c825

3 files changed

Lines changed: 3 additions & 4 deletions

File tree

cmake/llvm-hash.txt

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1 +1 @@
1-
979132a02d146ec79e2f046e31877516d7f32d20
1+
ac5dc54d509169d387fcfd495d71853d81c46484

python/src/llvm.cc

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -133,6 +133,7 @@ createTargetMachine(llvm::Module *module, std::string proc,
133133
bool disableLLVMOpt = mlir::triton::tools::getBoolEnv("DISABLE_LLVM_OPT");
134134
if (enable_fp_fusion)
135135
opt.AllowFPOpFusion = llvm::FPOpFusion::Fast;
136+
opt.NoInfsFPMath = false;
136137
opt.NoNaNsFPMath = true;
137138
opt.TrapUnreachable = true;
138139
opt.MCOptions.AsmVerbose = true;

python/test/unit/language/test_core.py

Lines changed: 1 addition & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1382,9 +1382,7 @@ def kernel(X, Z):
13821382
# atom.add.bf16 is unsupported prior to Hopper so instead we generate an
13831383
# atom.cas add loop on Ampere and prior
13841384
if dst_type == 'bfloat16' and torch.cuda.get_device_capability()[0] < 9:
1385-
assert "atom.relaxed.gpu.global.cas" in h.asm["ptx"]
1386-
if sem_str != "relaxed":
1387-
assert "fence.acq_rel.gpu" in h.asm["ptx"]
1385+
assert f"atom.{sem_str}.gpu.global.cas" in h.asm["ptx"]
13881386
return
13891387

13901388
assert f"atom.global.gpu.{sem_str}" in h.asm["ptx"]

0 commit comments

Comments
 (0)