[amdgpu] Add `llvm.amdgcn.endpgm` support.

[hliao]

Summary

• llvm.amdgcn.endpgm is added to enable "abort" support.

[rampitec]

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

1581

Mayby also IntrCold?

[hliao]

Add IntrCold.

[hliao]

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

1581

Good point!

[b-sumner]

Should this also be IntrConvergent?

[rampitec]

In D90809#2376994, @b-sumner wrote:

Should this also be IntrConvergent?

Probably yes... This is control flow after all.

[hliao]

In D90809#2377083, @rampitec wrote:

In D90809#2376994, @b-sumner wrote:

Should this also be IntrConvergent?

Probably yes... This is control flow after all.

The real effect of this intrinsic seems not changed if it's moved somewhere. For example,

convert this

if (a || b)
  endpgm();

to

if (a)
  endpgm();
else if (b)
  endpgm();

no matter a or b are divergent or not, it works the same as the original one as s_endpgm is a scalar instruction. Ofc, if we really doubt bad things may happen, I could add that for safety as we definitely will revise this later.

[rampitec]

In D90809#2377221, @hliao wrote:

In D90809#2377083, @rampitec wrote:

In D90809#2376994, @b-sumner wrote:

Should this also be IntrConvergent?

Probably yes... This is control flow after all.

The real effect of this intrinsic seems not changed if it's moved somewhere. For example,

convert this

if (a || b)
  endpgm();

to

if (a)
  endpgm();
else if (b)
  endpgm();

no matter a or b are divergent or not, it works the same as the original one as s_endpgm is a scalar instruction. Ofc, if we really doubt bad things may happen, I could add that for safety as we definitely will revise this later.

It is probably OK to clone it, we shall insert skips around it anyway. LGTM.

[foad]

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

1580

The intrinsic def needs a comment. Is it supposed to literally just generate an s_endpgm instruction, or is it something more high-level?

clang/include/clang/Basic/BuiltinsAMDGPU.def

 
 //===----------------------------------------------------------------------===//
 // Special builtins.
 //===----------------------------------------------------------------------===//
 BUILTIN(__builtin_amdgcn_read_exec, "LUi", "nc")
 BUILTIN(__builtin_amdgcn_read_exec_lo, "Ui", "nc")
 BUILTIN(__builtin_amdgcn_read_exec_hi, "Ui", "nc")
 
+BUILTIN(__builtin_amdgcn_endpgm, "v", "nr")
+
 //===----------------------------------------------------------------------===//
 // R600-NI only builtins.
 //===----------------------------------------------------------------------===//
 
 BUILTIN(__builtin_r600_implicitarg_ptr, "Uc*7", "nc")
 
 BUILTIN(__builtin_r600_read_tgid_x, "Ui", "nc")
 BUILTIN(__builtin_r600_read_tgid_y, "Ui", "nc")

clang/test/CodeGenCUDA/builtins-amdgcn.cu

 
 // CHECK-LABEL: @_Z12test_ds_fmaxf(
 // CHECK: call contract float @llvm.amdgcn.ds.fmax.f32(float addrspace(3)* @_ZZ12test_ds_fmaxfE6shared, float %{{[^,]*}}, i32 0, i32 0, i1 false)
 __global__
 void test_ds_fmax(float src) {
   __shared__ float shared;
   volatile float x = __builtin_amdgcn_ds_fmaxf(&shared, src, 0, 0, false);
 }
+
+// CHECK-LABEL: @_Z6endpgmv(
+// CHECK: call void @llvm.amdgcn.endpgm()
+__global__ void endpgm() {
+  __builtin_amdgcn_endpgm();
+}

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

 def int_amdgcn_wqm_vote : Intrinsic<[llvm_i1_ty],
   [llvm_i1_ty], [IntrNoMem, IntrConvergent, IntrWillReturn]
 >;
 
 // If false, set EXEC=0 for the current thread until the end of program.
 // FIXME: Should this be IntrNoMem, IntrHasSideEffects, or IntrWillReturn?
 def int_amdgcn_kill : Intrinsic<[], [llvm_i1_ty], []>;
 
+def int_amdgcn_endpgm : GCCBuiltin<"__builtin_amdgcn_endpgm">,
+  Intrinsic<[], [], [IntrNoReturn, IntrCold, IntrNoMem, IntrHasSideEffects]
+>;
+
 // Copies the active channels of the source value to the destination value,
 // with the guarantee that the source value is computed as if the entire
 // program were executed in Whole Wavefront Mode, i.e. with all channels
 // enabled, with a few exceptions: - Phi nodes with require WWM return an
 // undefined value.
 def int_amdgcn_wwm : Intrinsic<[llvm_any_ty],
   [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable,
                        IntrConvergent, IntrWillReturn]

llvm/lib/Target/AMDGPU/SOPInstructions.td

 }
 
 def S_NOP : SOPP_Pseudo<"s_nop" , (ins i16imm:$simm16), "$simm16">;
 
 let isTerminator = 1 in {
 def S_ENDPGM : SOPP_Pseudo<"s_endpgm", (ins EndpgmImm:$simm16), "$simm16"> {
   let isBarrier = 1;
   let isReturn = 1;
+  let hasSideEffects = 1;
 }
 
 def S_ENDPGM_SAVED : SOPP_Pseudo<"s_endpgm_saved", (ins)> {
   let SubtargetPredicate = isGFX8Plus;
   let simm16 = 0;
   let fixed_imm = 1;
   let isBarrier = 1;
   let isReturn = 1;

 //===----------------------------------------------------------------------===//
 
 def : GCNPat <
   (AMDGPUendpgm),
     (S_ENDPGM (i16 0))
 >;
 
 def : GCNPat <
+  (int_amdgcn_endpgm),
+    (S_ENDPGM (i16 0))
+>;
+
+def : GCNPat <
   (i64 (ctpop i64:$src)),
     (i64 (REG_SEQUENCE SReg_64,
      (i32 (COPY_TO_REGCLASS (S_BCNT1_I32_B64 $src), SReg_32)), sub0,
      (S_MOV_B32 (i32 0)), sub1))
 >;
 
 def : GCNPat <
   (i32 (smax i32:$x, (i32 (ineg i32:$x)))),

llvm/test/CodeGen/AMDGPU/amd.endpgm.ll

+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
+; RUN: llc -march=amdgcn -mcpu=gfx900 -verify-machineinstrs < %s | FileCheck %s
+
+define amdgpu_kernel void @test0() {
+; CHECK-LABEL: test0:
+; CHECK:       ; %bb.0:
+; CHECK-NEXT:    s_endpgm
+  tail call void @llvm.amdgcn.endpgm()
+  unreachable
+}
+
+define void @test1() {
+; CHECK-LABEL: test1:
+; CHECK:       ; %bb.0:
+; CHECK-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT:    s_endpgm
+  tail call void @llvm.amdgcn.endpgm()
+  unreachable
+}
+
+define amdgpu_kernel void @test2(i32* %p, i32 %x) {
+; CHECK-LABEL: test2:
+; CHECK:       ; %bb.0:
+; CHECK-NEXT:    s_load_dword s2, s[0:1], 0x2c
+; CHECK-NEXT:    s_waitcnt lgkmcnt(0)
+; CHECK-NEXT:    s_cmp_lt_i32 s2, 1
+; CHECK-NEXT:    s_cbranch_scc0 BB2_2
+; CHECK-NEXT:  ; %bb.1: ; %else
+; CHECK-NEXT:    s_load_dwordx2 s[0:1], s[0:1], 0x24
+; CHECK-NEXT:    v_mov_b32_e32 v2, s2
+; CHECK-NEXT:    s_waitcnt lgkmcnt(0)
+; CHECK-NEXT:    v_mov_b32_e32 v0, s0
+; CHECK-NEXT:    v_mov_b32_e32 v1, s1
+; CHECK-NEXT:    flat_store_dword v[0:1], v2
+; CHECK-NEXT:    s_endpgm
+; CHECK-NEXT:  BB2_2: ; %then
+; CHECK-NEXT:    s_endpgm
+  %cond = icmp sgt i32 %x, 0
+  br i1 %cond, label %then, label %else
+
+then:
+  tail call void @llvm.amdgcn.endpgm()
+  unreachable
+
+else:
+  store i32 %x, i32* %p
+  ret void
+}
+
+declare void @llvm.amdgcn.endpgm()