https://github.com/RyanRio created
https://github.com/llvm/llvm-project/pull/212929
```
unroll_full for (u32_t i = 1; i <= CONSTANT; ++i) {
__builtin_amdgcn_wait_asyncmark(CONSTANT - i);
// some other stuff
}
```
Often we'll want to do something like this, where it's very time consuming to
rewrite the source and manually unroll.
>From 3446b90431723e5d6b8845bb4117e36f5a259a81 Mon Sep 17 00:00:00 2001
From: Ryan Mitchell <[email protected]>
Date: Wed, 29 Jul 2026 21:52:49 -0700
Subject: [PATCH] Allow non const asyncmark operands until isel
---
clang/include/clang/Basic/BuiltinsAMDGPU.td | 2 +-
.../builtins-amdgcn-asyncmark.cl | 14 +++++++++++
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 3 ++-
llvm/lib/Target/AMDGPU/SOPInstructions.td | 2 +-
.../AMDGPU/asyncmark-non-constant-err.ll | 18 +++++++++++++
.../CodeGen/AMDGPU/asyncmark-unrolled-loop.ll | 25 +++++++++++++++++++
6 files changed, 61 insertions(+), 3 deletions(-)
create mode 100644 llvm/test/CodeGen/AMDGPU/asyncmark-non-constant-err.ll
create mode 100644 llvm/test/CodeGen/AMDGPU/asyncmark-unrolled-loop.ll
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td
b/clang/include/clang/Basic/BuiltinsAMDGPU.td
index 73b27a2b5c5cd..18811fc44ca72 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.td
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.td
@@ -331,7 +331,7 @@ def __builtin_amdgcn_av_store_b128
// Mirrors GCNSubtarget::hasAsyncMark()
def __builtin_amdgcn_asyncmark : AMDGPUBuiltin<"void()", [],
"vmem-to-lds-load-insts|asynccnt">;
-def __builtin_amdgcn_wait_asyncmark : AMDGPUBuiltin<"void(_Constant unsigned
short)", [], "vmem-to-lds-load-insts|asynccnt">;
+def __builtin_amdgcn_wait_asyncmark : AMDGPUBuiltin<"void(unsigned short)",
[], "vmem-to-lds-load-insts|asynccnt">;
//===----------------------------------------------------------------------===//
// Ballot builtins.
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-asyncmark.cl
b/clang/test/CodeGenOpenCL/builtins-amdgcn-asyncmark.cl
index 838bfdbdcb8f1..85843590b14de 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-asyncmark.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-asyncmark.cl
@@ -15,3 +15,17 @@ void test_invocation() {
__builtin_amdgcn_asyncmark();
__builtin_amdgcn_wait_asyncmark(0);
}
+
+// The argument is no longer required to be an integer constant expression in
+// the frontend; it must only fold to a constant by instruction selection.
+// CHECK-LABEL: @test_non_ice(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[N_ADDR:%.*]] = alloca i16, align 2, addrspace(5)
+// CHECK-NEXT: store i16 [[N:%.*]], ptr addrspace(5) [[N_ADDR]], align 2
+// CHECK-NEXT: [[TMP0:%.*]] = load i16, ptr addrspace(5) [[N_ADDR]], align 2
+// CHECK-NEXT: call void @llvm.amdgcn.wait.asyncmark(i16 [[TMP0]])
+// CHECK-NEXT: ret void
+//
+void test_non_ice(unsigned short n) {
+ __builtin_amdgcn_wait_asyncmark(n);
+}
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index 33203de2d19a3..cfdeed160a86b 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -2880,9 +2880,10 @@ def int_amdgcn_asyncmark :
ClangBuiltin<"__builtin_amdgcn_asyncmark">,
Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects]>;
// Waits until the Nth previous marker is completed, if it exists.
+// N must fold to a constant by instruction selection.
def int_amdgcn_wait_asyncmark :
ClangBuiltin<"__builtin_amdgcn_wait_asyncmark">,
- Intrinsic<[], [llvm_i16_ty], [ImmArg<ArgIndex<0>>, IntrNoMem,
IntrHasSideEffects]>;
+ Intrinsic<[], [llvm_i16_ty], [IntrNoMem, IntrHasSideEffects]>;
//===----------------------------------------------------------------------===//
// GFX10 Intrinsics
diff --git a/llvm/lib/Target/AMDGPU/SOPInstructions.td
b/llvm/lib/Target/AMDGPU/SOPInstructions.td
index d18e96b7621b1..bc0ccb84e115d 100644
--- a/llvm/lib/Target/AMDGPU/SOPInstructions.td
+++ b/llvm/lib/Target/AMDGPU/SOPInstructions.td
@@ -1729,7 +1729,7 @@ def ASYNCMARK : SPseudoInstSI<(outs), (ins),
let isMeta = 1;
}
def WAIT_ASYNCMARK : SOPP_Pseudo <"", (ins s16imm:$simm16), "$simm16",
- [(int_amdgcn_wait_asyncmark timm:$simm16)]> {
+ [(int_amdgcn_wait_asyncmark imm:$simm16)]> {
let maybeAtomic = 0;
let Size = 0;
let isMeta = 1;
diff --git a/llvm/test/CodeGen/AMDGPU/asyncmark-non-constant-err.ll
b/llvm/test/CodeGen/AMDGPU/asyncmark-non-constant-err.ll
new file mode 100644
index 0000000000000..c0e9e1190c4a2
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/asyncmark-non-constant-err.ll
@@ -0,0 +1,18 @@
+; RUN: not llc -global-isel=0 -mtriple=amdgpu12.50 -filetype=null %s 2>&1 |
FileCheck %s
+; RUN: not llc -global-isel=1 -mtriple=amdgpu12.50 -filetype=null %s 2>&1 |
FileCheck %s -check-prefix=CHECK-GISEL
+
+; The wait_asyncmark operand indexes compiler-tracked async mark state, so it
+; must fold to a constant by instruction selection.
+
+; CHECK: LLVM ERROR: Cannot select: intrinsic %llvm.amdgcn.wait.asyncmark
+; CHECK-GISEL: LLVM ERROR: cannot select: G_INTRINSIC_W_SIDE_EFFECTS
intrinsic(@llvm.amdgcn.wait.asyncmark), %10:sgpr(i16) (in function:
non_constant_wait)
+
+define amdgpu_kernel void @non_constant_wait(i16 %n) {
+entry:
+ call void @llvm.amdgcn.asyncmark()
+ call void @llvm.amdgcn.wait.asyncmark(i16 %n)
+ ret void
+}
+
+declare void @llvm.amdgcn.asyncmark()
+declare void @llvm.amdgcn.wait.asyncmark(i16)
diff --git a/llvm/test/CodeGen/AMDGPU/asyncmark-unrolled-loop.ll
b/llvm/test/CodeGen/AMDGPU/asyncmark-unrolled-loop.ll
new file mode 100644
index 0000000000000..735c6cafa3482
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/asyncmark-unrolled-loop.ll
@@ -0,0 +1,25 @@
+; RUN: opt -O2 -S %s | llc -global-isel=0 -mtriple=amdgpu12.50 -o - |
FileCheck %s
+; RUN: opt -O2 -S %s | llc -global-isel=1 -mtriple=amdgpu12.50 -o - |
FileCheck %s
+
+; CHECK: ; wait_asyncmark(2)
+; CHECK: ; wait_asyncmark(1)
+
+define amdgpu_kernel void @unrolled_loop() {
+entry:
+ br label %loop
+
+loop:
+ %i = phi i16 [ 0, %entry ], [ %inc, %loop ]
+ call void @llvm.amdgcn.asyncmark()
+ %n = sub i16 2, %i
+ call void @llvm.amdgcn.wait.asyncmark(i16 %n)
+ %inc = add i16 %i, 1
+ %cmp = icmp ult i16 %inc, 2
+ br i1 %cmp, label %loop, label %exit
+
+exit:
+ ret void
+}
+
+declare void @llvm.amdgcn.asyncmark()
+declare void @llvm.amdgcn.wait.asyncmark(i16)
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
