llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang @llvm/pr-subscribers-backend-amdgpu Author: Arseniy Obolenskiy (aobolensk) <details> <summary>Changes</summary> Align with classic codegen, which emits add nsw and fcmp one Classic codegen code we are aligning with: https://github.com/llvm/llvm-project/blob/8995a358fa77a669679594aa9857ba0e0b224ab1/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp#L2233-L2239 --- Full diff: https://github.com/llvm/llvm-project/pull/228004.diff 2 Files Affected: - (modified) clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp (+2-2) - (modified) clang/test/CIR/CodeGenHIP/builtins-amdgcn-logb-scalbn.hip (+4) ``````````diff diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index b2aac0d376fa1..e0fc28f278d1b 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -89,7 +89,7 @@ static mlir::Value emitLogbBuiltin(CIRGenFunction &cgf, const CallExpr *e, mlir::Value negativeOne = builder.getConstant(loc, cir::IntAttr::get(int32Ty, -1)); - mlir::Value expMinus1 = builder.createAdd(loc, exp, negativeOne); + mlir::Value expMinus1 = builder.createNSWAdd(loc, exp, negativeOne); mlir::Value siToFp = cir::CastOp::create( builder, loc, srcTy, cir::CastKind::int_to_float, expMinus1); @@ -100,7 +100,7 @@ static mlir::Value emitLogbBuiltin(CIRGenFunction &cgf, const CallExpr *e, mlir::Value inf = builder.getConstant(loc, cir::FPAttr::get(srcTy, infVal)); mlir::Value fabsNegInf = - builder.createCompare(loc, cir::CmpOpKind::ne, fabs, inf); + builder.createCompare(loc, cir::CmpOpKind::one, fabs, inf); mlir::Value sel = builder.createSelect(loc, fabsNegInf, siToFp, fabs); diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-logb-scalbn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-logb-scalbn.hip index b4d28a8bb7487..d5f31ff5cf019 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-logb-scalbn.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-logb-scalbn.hip @@ -17,6 +17,8 @@ // CIR: cir.call_llvm_intrinsic "frexp" {{.*}} : (!cir.float) -> {{.*}} // LLVM: define{{.*}} float @_Z11test_logbfff( // LLVM: call { float, i32 } @llvm.frexp.f32.i32(float {{.*}}) +// LLVM: add nsw i32 %{{.*}}, -1 +// LLVM: fcmp{{.*}} one float %{{.*}}, +inf __device__ float test_logbff(float a) { return __builtin_logbf(a); } @@ -25,6 +27,8 @@ __device__ float test_logbff(float a) { // CIR: cir.call_llvm_intrinsic "frexp" {{.*}} : (!cir.double) -> {{.*}} // LLVM: define{{.*}} double @_Z11test_logbddd( // LLVM: call { double, i32 } @llvm.frexp.f64.i32(double {{.*}}) +// LLVM: add nsw i32 %{{.*}}, -1 +// LLVM: fcmp{{.*}} one double %{{.*}}, +inf __device__ double test_logbdd(double a) { return __builtin_logb(a); } `````````` </details> https://github.com/llvm/llvm-project/pull/228004 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
