[clang] [llvm] [AMDGPU] Add builtins for wave reduction intrinsics (PR #127013)

[Matt Arsenault via cfe-commits]

================
@@ -20212,6 +20212,59 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
     llvm::Value *Env = EmitScalarExpr(E->getArg(0));
     return Builder.CreateCall(F, {Env});
   }
+  case AMDGPU::BI__builtin_amdgcn_wave_reduce_add_i32:
+  case AMDGPU::BI__builtin_amdgcn_wave_reduce_uadd_i32:
+  case AMDGPU::BI__builtin_amdgcn_wave_reduce_sub_i32:
+  case AMDGPU::BI__builtin_amdgcn_wave_reduce_usub_i32:
+  case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_i32:
+  case AMDGPU::BI__builtin_amdgcn_wave_reduce_umin_i32:
+  case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_i32:
+  case AMDGPU::BI__builtin_amdgcn_wave_reduce_umax_i32:
+  case AMDGPU::BI__builtin_amdgcn_wave_reduce_and_i32:
+  case AMDGPU::BI__builtin_amdgcn_wave_reduce_or_i32:
+  case AMDGPU::BI__builtin_amdgcn_wave_reduce_xor_i32: {
+    Intrinsic::ID IID;
+    switch (BuiltinID) {
----------------
arsenm wrote:

We should have a builtin -> intrinsic switch function somewhere 

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