[llvm] [AMDGPU] Guard named barrier intrinsics with gfx12-insts (PR #227607)

[via llvm-commits]

llvmorg-github-actions[bot] wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-backend-amdgpu

Author: Vikram Hegde (vikramRH)

<details>
<summary>Changes</summary>

GIsel cases still crash since IRTranslator continues after diagnosing. probably needs to be fixed ?

Fixes LCOMPILER-2875
Test assisted by: Opus 5

---
Full diff: https://github.com/llvm/llvm-project/pull/227607.diff


2 Files Affected:

- (modified) llvm/include/llvm/IR/IntrinsicsAMDGPU.td (+8) 
- (added) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.barrier.unsupported-subtarget.ll (+90) 


``````````diff
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index 3ab24ef12db582..27cc0c864a4c38 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 00000000000000..65436fa680f584
--- /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
+}

``````````

</details>


https://github.com/llvm/llvm-project/pull/227607