[clang] d917e35 - [CIR][AMDGPU] Use nsw add and ordered compare in logb lowering (#228004)

[via cfe-commits]

Author: Arseniy Obolenskiy
Date: 2026-10-02T09:49:22Z
New Revision: d917e358c84b694c9e4250a270f4833acbfb8f41

URL: https://github.com/llvm/llvm-project/commit/d917e358c84b694c9e4250a270f4833acbfb8f41
DIFF: https://github.com/llvm/llvm-project/commit/d917e358c84b694c9e4250a270f4833acbfb8f41.diff

LOG: [CIR][AMDGPU] Use nsw add and ordered compare in logb lowering (#228004)

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

Added: 
    

Modified: 
    clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
    clang/test/CIR/CodeGenHIP/builtins-amdgcn-logb-scalbn.hip

Removed: 
    


################################################################################
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);
 }