================
@@ -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
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
