[clang] [AMDGPU] Use generic builtins for `wave_reduce` ops (PR #179589)

[Alex Voicu via cfe-commits]

================
@@ -417,37 +399,17 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
   llvm::AtomicOrdering AO = llvm::AtomicOrdering::SequentiallyConsistent;
   llvm::SyncScope::ID SSID;
   switch (BuiltinID) {
-  case AMDGPU::BI__builtin_amdgcn_wave_reduce_add_u32:
-  case AMDGPU::BI__builtin_amdgcn_wave_reduce_fadd_f32:
-  case AMDGPU::BI__builtin_amdgcn_wave_reduce_fadd_f64:
-  case AMDGPU::BI__builtin_amdgcn_wave_reduce_sub_u32:
-  case AMDGPU::BI__builtin_amdgcn_wave_reduce_fsub_f32:
-  case AMDGPU::BI__builtin_amdgcn_wave_reduce_fsub_f64:
-  case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_i32:
-  case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_u32:
-  case AMDGPU::BI__builtin_amdgcn_wave_reduce_fmin_f32:
-  case AMDGPU::BI__builtin_amdgcn_wave_reduce_fmin_f64:
-  case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_i32:
-  case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_u32:
-  case AMDGPU::BI__builtin_amdgcn_wave_reduce_fmax_f32:
-  case AMDGPU::BI__builtin_amdgcn_wave_reduce_fmax_f64:
-  case AMDGPU::BI__builtin_amdgcn_wave_reduce_and_b32:
-  case AMDGPU::BI__builtin_amdgcn_wave_reduce_or_b32:
-  case AMDGPU::BI__builtin_amdgcn_wave_reduce_xor_b32:
-  case AMDGPU::BI__builtin_amdgcn_wave_reduce_add_u64:
-  case AMDGPU::BI__builtin_amdgcn_wave_reduce_sub_u64:
-  case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_i64:
-  case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_u64:
-  case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_i64:
-  case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_u64:
-  case AMDGPU::BI__builtin_amdgcn_wave_reduce_and_b64:
-  case AMDGPU::BI__builtin_amdgcn_wave_reduce_or_b64:
-  case AMDGPU::BI__builtin_amdgcn_wave_reduce_xor_b64: {
-    Intrinsic::ID IID = getIntrinsicIDforWaveReduction(BuiltinID);
-    llvm::Value *Value = EmitScalarExpr(E->getArg(0));
-    llvm::Value *Strategy = EmitScalarExpr(E->getArg(1));
-    llvm::Function *F = CGM.getIntrinsic(IID, {Value->getType()});
-    return Builder.CreateCall(F, {Value, Strategy});
+  case AMDGPU::BI__builtin_amdgcn_wave_reduce_add:
+  case AMDGPU::BI__builtin_amdgcn_wave_reduce_and:
----------------
AlexVlx wrote:

I had these in enum order, but if you think it improves readability we can change it.

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