[AMDGPU] Intrinsic to expose s_wait_event for export ready

[foad]

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

2070

Should this be int_amdgcn_s_wait_event_export_ready? I'm not sure how consistent we are about including the "s" or "v" prefix in intrinsic names.

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

8375 ↗

(On Diff #476134)

Why do you need this? Can't you do it with a tablegen selection pattern?

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.wait.event.ll

1 ↗

(On Diff #476134)

Should test globalisel as well.

[dstuttard]

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

2070

I'm in two minds about this since the intrinsic doesn't strictly map to the instruction. I think I slightly prefer this one.

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

8375 ↗

(On Diff #476134)

Yes, I'll update it. There are potentially some complexities that may have to be handled in a later patch, but I think they can also be handled in tablegen too.

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.wait.event.ll

1 ↗

(On Diff #476134)

Yes, will do.

[dstuttard]

Updating based on review comments

[arsenm]

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

2071

Why is the immediate missing from the intrinsic? Should this also have the ClangBuiltin name declared here?

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.wait.event.ll

1 ↗

(On Diff #476149)

Should use an explicit -global-isel=0

9 ↗

(On Diff #476149)

All the other scalar intrinsics have the s in the name

[dstuttard]

Address review comments

[dstuttard]

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

2071

Added the ClangBuiltin
Updated the name to include the "s" since that's 2 people who've asked now.
The intrinsic doesn't map to s_wait_event, just to the s_wait_event export_ready - that's why it doesn't take the immediate argument.

[Joe_Nash]

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

2071

If you are adding the clang builtin perhaps there should there be an OpenCL test.

[dstuttard]

Updated - adding test for builtin

[Joe_Nash]

LGTM

clang/include/clang/Basic/BuiltinsAMDGPU.def

 
 
 //===----------------------------------------------------------------------===//
 // GFX11+ only builtins.
 //===----------------------------------------------------------------------===//
 
 // TODO: This is a no-op in wave32. Should the builtin require wavefrontsize64?
 TARGET_BUILTIN(__builtin_amdgcn_permlane64, "UiUi", "nc", "gfx11-insts")
+TARGET_BUILTIN(__builtin_amdgcn_s_wait_event_export_ready, "v", "n", "gfx11-insts")
 
 //===----------------------------------------------------------------------===//
 // WMMA builtins.
 // Postfix w32 indicates the builtin requires wavefront size of 32.
 // Postfix w64 indicates the builtin requires wavefront size of 64.
 //===----------------------------------------------------------------------===//
 TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_f16_w32, "V8fV16hV16hV8f", "nc", "gfx11-insts")
 TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_bf16_w32, "V8fV16sV16sV8f", "nc", "gfx11-insts")

clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11.cl

void test_ds_bvh_stack_rtn(global uint2* out, uint addr, uint data, uint4 data1)
   *out = __builtin_amdgcn_ds_bvh_stack_rtn(addr, data, data1, 128);
 }
 
 // CHECK-LABEL: @test_permlane64(
 // CHECK: call i32 @llvm.amdgcn.permlane64(i32 %a)
 void test_permlane64(global uint* out, uint a) {
   *out = __builtin_amdgcn_permlane64(a);
 }
+
+// CHECK-LABEL: @test_s_wait_event_export_ready
+// CHECK: call void @llvm.amdgcn.s.wait.event.export.ready
+void test_s_wait_event_export_ready() {
+  __builtin_amdgcn_s_wait_event_export_ready();
+}

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

 
 def int_amdgcn_wmma_f32_16x16x16_f16   : AMDGPUWmmaIntrinsic<llvm_v16f16_ty, llvm_anyfloat_ty>;
 def int_amdgcn_wmma_f32_16x16x16_bf16  : AMDGPUWmmaIntrinsic<llvm_v16i16_ty, llvm_anyfloat_ty>;
 def int_amdgcn_wmma_f16_16x16x16_f16   : AMDGPUWmmaIntrinsicOPSEL<llvm_v16f16_ty, llvm_anyfloat_ty>;
 def int_amdgcn_wmma_bf16_16x16x16_bf16 : AMDGPUWmmaIntrinsicOPSEL<llvm_v16i16_ty, llvm_anyint_ty>;
 def int_amdgcn_wmma_i32_16x16x16_iu8   : AMDGPUWmmaIntrinsicIU<llvm_v4i32_ty, llvm_anyint_ty>;
 def int_amdgcn_wmma_i32_16x16x16_iu4   : AMDGPUWmmaIntrinsicIU<llvm_v2i32_ty, llvm_anyint_ty>;
 
+def int_amdgcn_s_wait_event_export_ready :
+  ClangBuiltin<"__builtin_amdgcn_s_wait_event_export_ready">,
+  Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects, IntrWillReturn]
+>;
 
 //===----------------------------------------------------------------------===//
 // Deep learning intrinsics.
 //===----------------------------------------------------------------------===//
 
 // f32 %r = llvm.amdgcn.fdot2(v2f16 %a, v2f16 %b, f32 %c, i1 %clamp)
 //   %r = %a[0] * %b[0] + %a[1] * %b[1] + %c
 def int_amdgcn_fdot2 :

llvm/lib/Target/AMDGPU/SOPInstructions.td

let SubtargetPredicate = isGFX10Plus in {
   }
 
   def S_TTRACEDATA_IMM :
     SOPP_Pseudo<"s_ttracedata_imm", (ins s16imm:$simm16), "$simm16">;
 } // End SubtargetPredicate = isGFX10Plus
 
 let SubtargetPredicate = isGFX11Plus in {
   def S_WAIT_EVENT : SOPP_Pseudo<"s_wait_event", (ins s16imm:$simm16),
-                                 "$simm16">;
+                                 "$simm16"> {
+                                   let hasSideEffects = 1;
+                                 }
   def S_DELAY_ALU : SOPP_Pseudo<"s_delay_alu", (ins DELAY_FLAG:$simm16),
                                 "$simm16">;
 } // End SubtargetPredicate = isGFX11Plus
 
 //===----------------------------------------------------------------------===//
 // SOP1 Patterns
 //===----------------------------------------------------------------------===//
 

 >;
 
 // Same as a 32-bit inreg
 def : GCNPat<
   (i32 (UniformUnaryFrag<sext> i16:$src)),
   (S_SEXT_I32_I16 $src)
 >;
 
+def : GCNPat <
+  (int_amdgcn_s_wait_event_export_ready),
+    (S_WAIT_EVENT (i16 0))
+>;
 
 //===----------------------------------------------------------------------===//
 // SOP2 Patterns
 //===----------------------------------------------------------------------===//
 
 // V_ADD_I32_e32/S_ADD_U32 produces carry in VCC/SCC. For the vector
 // case, the sgpr-copies pass will fix this to use the vector version.
 def : GCNPat <

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.wait.event.ll

+; RUN: llc -global-isel=0 -march=amdgcn -verify-machineinstrs -mcpu=gfx1100 < %s | FileCheck -check-prefix=GCN %s
+; RUN: llc -global-isel -march=amdgcn -verify-machineinstrs -mcpu=gfx1100 < %s | FileCheck -check-prefix=GCN %s
+
+; GCN-LABEL: {{^}}test_wait_event:
+; GCN: s_wait_event 0x0
+
+define amdgpu_ps void @test_wait_event() #0 {
+entry:
+  call void @llvm.amdgcn.s.wait.event.export.ready() #0
+  ret void
+}
+
+declare void @llvm.amdgcn.s.wait.event.export.ready() #0
+
+attributes #0 = { nounwind }