[llvm] [AMDGPU] Guard named barrier intrinsics with gfx12-insts (PR #227607)
Vikram Hegde via llvm-commits
llvm-commits at lists.llvm.org
Wed Sep 30 23:53:21 PDT 2026
https://github.com/vikramRH updated https://github.com/llvm/llvm-project/pull/227607
>From 3a3609760cb30b2b7bbe53612e269c4e08289633 Mon Sep 17 00:00:00 2001
From: vikhegde <vikram.hegde at amd.com>
Date: Wed, 30 Sep 2026 14:19:29 +0530
Subject: [PATCH] [AMDGPU] Guard named barrier intrinsics with gfx12-insts
---
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 8 ++
....amdgcn.s.barrier.unsupported-subtarget.ll | 90 +++++++++++++++++++
2 files changed, 98 insertions(+)
create mode 100644 llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.barrier.unsupported-subtarget.ll
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index 3ab24ef12db58..27cc0c864a4c3 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -286,6 +286,8 @@ def int_amdgcn_s_barrier : ClangBuiltin<"__builtin_amdgcn_s_barrier">,
// Lower-level split-barrier intrinsics
+let TargetFeatures = "gfx12-insts" in {
+
// void @llvm.amdgcn.s.barrier.signal(i32 %barrierType)
// only for non-named barrier
def int_amdgcn_s_barrier_signal : ClangBuiltin<"__builtin_amdgcn_s_barrier_signal">,
@@ -318,6 +320,8 @@ def int_amdgcn_s_barrier_join : ClangBuiltin<"__builtin_amdgcn_s_barrier_join">,
Intrinsic<[], [barrier_ptr_ty], [IntrNoMem, IntrHasSideEffects, IntrConvergent, IntrWillReturn,
IntrNoCallback, IntrNoFree]>;
+} // End TargetFeatures = "gfx12-insts"
+
// void @llvm.amdgcn.s.wakeup.barrier(ptr addrspace(3) %barrier)
// The %barrier argument must be uniform, otherwise behavior is undefined.
let TargetFeatures = "s-wakeup-barrier-inst" in
@@ -325,6 +329,8 @@ def int_amdgcn_s_wakeup_barrier : ClangBuiltin<"__builtin_amdgcn_s_wakeup_barrie
Intrinsic<[], [barrier_ptr_ty], [IntrNoMem, IntrHasSideEffects, IntrConvergent, IntrWillReturn,
IntrNoCallback, IntrNoFree]>;
+let TargetFeatures = "gfx12-insts" in {
+
// void @llvm.amdgcn.s.barrier.wait(i16 %barrierType)
def int_amdgcn_s_barrier_wait : ClangBuiltin<"__builtin_amdgcn_s_barrier_wait">,
Intrinsic<[], [llvm_i16_ty], [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects, IntrConvergent,
@@ -348,6 +354,8 @@ def int_amdgcn_s_get_named_barrier_state : ClangBuiltin<"__builtin_amdgcn_s_get_
Intrinsic<[llvm_i32_ty], [barrier_ptr_ty], [IntrNoMem, IntrHasSideEffects, IntrConvergent, IntrWillReturn,
IntrNoCallback, IntrNoFree]>;
+} // End TargetFeatures = "gfx12-insts"
+
def int_amdgcn_wave_barrier : ClangBuiltin<"__builtin_amdgcn_wave_barrier">,
Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects, IntrConvergent, IntrWillReturn, IntrNoCallback, IntrNoFree]>;
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.barrier.unsupported-subtarget.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.barrier.unsupported-subtarget.ll
new file mode 100644
index 0000000000000..65436fa680f58
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.barrier.unsupported-subtarget.ll
@@ -0,0 +1,90 @@
+; RUN: split-file %s %t
+
+; RUN: not llc -global-isel=0 -mtriple=amdgpu9.42-amd-amdhsa -filetype=null %t/signal.ll 2>&1 | FileCheck -check-prefix=ERR-SIGNAL %s
+; RUN: not llc -global-isel=0 -mtriple=amdgpu9.42-amd-amdhsa -filetype=null %t/signal-var.ll 2>&1 | FileCheck -check-prefix=ERR-SIGNAL-VAR %s
+; RUN: not llc -global-isel=0 -mtriple=amdgpu9.42-amd-amdhsa -filetype=null %t/signal-isfirst.ll 2>&1 | FileCheck -check-prefix=ERR-SIGNAL-ISFIRST %s
+; RUN: not llc -global-isel=0 -mtriple=amdgpu9.42-amd-amdhsa -filetype=null %t/init.ll 2>&1 | FileCheck -check-prefix=ERR-INIT %s
+; RUN: not llc -global-isel=0 -mtriple=amdgpu9.42-amd-amdhsa -filetype=null %t/join.ll 2>&1 | FileCheck -check-prefix=ERR-JOIN %s
+; RUN: not llc -global-isel=0 -mtriple=amdgpu9.42-amd-amdhsa -filetype=null %t/wait.ll 2>&1 | FileCheck -check-prefix=ERR-WAIT %s
+; RUN: not llc -global-isel=0 -mtriple=amdgpu9.42-amd-amdhsa -filetype=null %t/leave.ll 2>&1 | FileCheck -check-prefix=ERR-LEAVE %s
+; RUN: not llc -global-isel=0 -mtriple=amdgpu9.42-amd-amdhsa -filetype=null %t/get-state.ll 2>&1 | FileCheck -check-prefix=ERR-GET-STATE %s
+
+; TODO: GIsel should not crash.
+; RUN: not --crash llc -global-isel=1 -mtriple=amdgpu9.42-amd-amdhsa -filetype=null %t/signal.ll 2>&1 | FileCheck -check-prefix=ERR-SIGNAL %s
+; RUN: not --crash llc -global-isel=1 -mtriple=amdgpu9.42-amd-amdhsa -filetype=null %t/signal-var.ll 2>&1 | FileCheck -check-prefix=ERR-SIGNAL-VAR %s
+; RUN: not --crash llc -global-isel=1 -mtriple=amdgpu9.42-amd-amdhsa -filetype=null %t/signal-isfirst.ll 2>&1 | FileCheck -check-prefix=ERR-SIGNAL-ISFIRST %s
+; RUN: not --crash llc -global-isel=1 -mtriple=amdgpu9.42-amd-amdhsa -filetype=null %t/init.ll 2>&1 | FileCheck -check-prefix=ERR-INIT %s
+; RUN: not --crash llc -global-isel=1 -mtriple=amdgpu9.42-amd-amdhsa -filetype=null %t/join.ll 2>&1 | FileCheck -check-prefix=ERR-JOIN %s
+; RUN: not --crash llc -global-isel=1 -mtriple=amdgpu9.42-amd-amdhsa -filetype=null %t/wait.ll 2>&1 | FileCheck -check-prefix=ERR-WAIT %s
+; RUN: not llc -global-isel=1 -mtriple=amdgpu9.42-amd-amdhsa -filetype=null %t/leave.ll 2>&1 | FileCheck -check-prefix=ERR-LEAVE %s
+; RUN: not --crash llc -global-isel=1 -mtriple=amdgpu9.42-amd-amdhsa -filetype=null %t/get-state.ll 2>&1 | FileCheck -check-prefix=ERR-GET-STATE %s
+
+;--- signal.ll
+; ERR-SIGNAL: error: {{.*}}in function @signal void (): llvm.amdgcn.s.barrier.signal requires target feature 'gfx12-insts'
+
+define amdgpu_kernel void @signal() {
+ call void @llvm.amdgcn.s.barrier.signal(i32 -1)
+ ret void
+}
+
+;--- signal-var.ll
+; ERR-SIGNAL-VAR: error: {{.*}}in function @signal_var void (): llvm.amdgcn.s.barrier.signal.var requires target feature 'gfx12-insts'
+
+ at bar_signal_var = addrspace(15) global target("amdgcn.named.barrier", 0) poison
+
+define amdgpu_kernel void @signal_var() {
+ call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(15) @bar_signal_var, i32 4)
+ ret void
+}
+
+;--- signal-isfirst.ll
+; ERR-SIGNAL-ISFIRST: error: {{.*}}in function @signal_isfirst void (): llvm.amdgcn.s.barrier.signal.isfirst requires target feature 'gfx12-insts'
+
+define amdgpu_kernel void @signal_isfirst() {
+ %r = call i1 @llvm.amdgcn.s.barrier.signal.isfirst(i32 -1)
+ ret void
+}
+
+;--- init.ll
+; ERR-INIT: error: {{.*}}in function @init void (): llvm.amdgcn.s.barrier.init requires target feature 'gfx12-insts'
+
+ at bar_init = addrspace(15) global target("amdgcn.named.barrier", 0) poison
+
+define amdgpu_kernel void @init() {
+ call void @llvm.amdgcn.s.barrier.init(ptr addrspace(15) @bar_init, i32 4)
+ ret void
+}
+
+;--- join.ll
+; ERR-JOIN: error: {{.*}}in function @join void (): llvm.amdgcn.s.barrier.join requires target feature 'gfx12-insts'
+
+ at bar_join = addrspace(15) global target("amdgcn.named.barrier", 0) poison
+
+define amdgpu_kernel void @join() {
+ call void @llvm.amdgcn.s.barrier.join(ptr addrspace(15) @bar_join)
+ ret void
+}
+
+;--- wait.ll
+; ERR-WAIT: error: {{.*}}in function @wait void (): llvm.amdgcn.s.barrier.wait requires target feature 'gfx12-insts'
+
+define amdgpu_kernel void @wait() {
+ call void @llvm.amdgcn.s.barrier.wait(i16 -1)
+ ret void
+}
+
+;--- leave.ll
+; ERR-LEAVE: error: {{.*}}in function @leave void (): llvm.amdgcn.s.barrier.leave requires target feature 'gfx12-insts'
+
+define amdgpu_kernel void @leave() {
+ call void @llvm.amdgcn.s.barrier.leave(i16 -1)
+ ret void
+}
+
+;--- get-state.ll
+; ERR-GET-STATE: error: {{.*}}in function @get_state void (): llvm.amdgcn.s.get.barrier.state requires target feature 'gfx12-insts'
+
+define amdgpu_kernel void @get_state() {
+ %r = call i32 @llvm.amdgcn.s.get.barrier.state(i32 -1)
+ ret void
+}
More information about the llvm-commits
mailing list