diff --git a/lib/CodeGen/CGExprScalar.cpp b/lib/CodeGen/CGExprScalar.cpp
index a75d777..ffd0eb5 100644
--- a/lib/CodeGen/CGExprScalar.cpp
+++ b/lib/CodeGen/CGExprScalar.cpp
@@ -583,19 +583,17 @@
       Check = Builder.CreateAnd(GE, LE);
     }
   } else {
-    // Floating-point to integer or floating-point to floating-point. This has
-    // undefined behavior if the source is +-Inf, NaN, or doesn't fit into the
-    // destination type (after truncation to an integer for float-to-integer).
     const llvm::fltSemantics &SrcSema =
       CGF.getContext().getFloatTypeSemantics(OrigSrcType);
-    APFloat MaxSrc(SrcSema, APFloat::uninitialized);
-    APFloat MinSrc(SrcSema, APFloat::uninitialized);
-
     if (isa<llvm::IntegerType>(DstTy)) {
+      // Floating-point to integer. This has undefined behavior if the source is
+      // +-Inf, NaN, or doesn't fit into the destination type (after truncation
+      // to an integer).
       unsigned Width = CGF.getContext().getIntWidth(DstType);
       bool Unsigned = DstType->isUnsignedIntegerOrEnumerationType();
 
       APSInt Min = APSInt::getMinValue(Width, Unsigned);
+      APFloat MinSrc(SrcSema, APFloat::uninitialized);
       if (MinSrc.convertFromAPInt(Min, !Unsigned, APFloat::rmTowardZero) &
           APFloat::opOverflow)
         // Don't need an overflow check for lower bound. Just check for
@@ -607,6 +605,7 @@
         MinSrc.subtract(APFloat(SrcSema, 1), APFloat::rmTowardNegative);
 
       APSInt Max = APSInt::getMaxValue(Width, Unsigned);
+      APFloat MaxSrc(SrcSema, APFloat::uninitialized);
       if (MaxSrc.convertFromAPInt(Max, !Unsigned, APFloat::rmTowardZero) &
           APFloat::opOverflow)
         // Don't need an overflow check for upper bound. Just check for
@@ -616,44 +615,60 @@
         // Find the smallest value which is too large to represent (before
         // truncation toward zero).
         MaxSrc.add(APFloat(SrcSema, 1), APFloat::rmTowardPositive);
-    } else {
-      const llvm::fltSemantics &DstSema =
-        CGF.getContext().getFloatTypeSemantics(DstType);
-      bool IsInexact;
 
-      MinSrc = APFloat::getLargest(DstSema, true);
-      if (MinSrc.convert(SrcSema, APFloat::rmTowardZero, &IsInexact) &
-          APFloat::opOverflow)
-        MinSrc = APFloat::getLargest(SrcSema, true);
+      // If we're converting from __half, convert the range to float to match
+      // the type of src.
+      if (OrigSrcType->isHalfType()) {
+        const llvm::fltSemantics &Sema =
+          CGF.getContext().getFloatTypeSemantics(SrcType);
+        bool IsInexact;
+        MinSrc.convert(Sema, APFloat::rmTowardZero, &IsInexact);
+        MaxSrc.convert(Sema, APFloat::rmTowardZero, &IsInexact);
+      }
 
-      MaxSrc = APFloat::getLargest(DstSema, false);
-      if (MaxSrc.convert(SrcSema, APFloat::rmTowardZero, &IsInexact) &
-          APFloat::opOverflow)
-        MaxSrc = APFloat::getLargest(SrcSema, false);
-    }
-
-    // If we're converting from __half, convert the range to float to match
-    // the type of src.
-    if (OrigSrcType->isHalfType()) {
-      const llvm::fltSemantics &Sema =
-        CGF.getContext().getFloatTypeSemantics(SrcType);
-      bool IsInexact;
-      MinSrc.convert(Sema, APFloat::rmTowardZero, &IsInexact);
-      MaxSrc.convert(Sema, APFloat::rmTowardZero, &IsInexact);
-    }
-
-    if (isa<llvm::IntegerType>(DstTy)) {
       llvm::Value *GE =
         Builder.CreateFCmpOGT(Src, llvm::ConstantFP::get(VMContext, MinSrc));
       llvm::Value *LE =
         Builder.CreateFCmpOLT(Src, llvm::ConstantFP::get(VMContext, MaxSrc));
       Check = Builder.CreateAnd(GE, LE);
     } else {
+      // FIXME: Maybe split this sanitizer out from float-cast-overflow.
+      //
+      // Floating-point to floating-point. This has undefined behavior if the
+      // source is not in the range of representable values of the destination
+      // type. The C and C++ standards are spectacularly unclear here. We
+      // diagnose finite out-of-range conversions, but allow infinities and NaNs
+      // to convert to the corresponding value in the smaller type.
+      //
+      // C11 Annex F gives all such conversions defined behavior for IEC 60559
+      // conforming implementations. Unfortunately, LLVM's fptrunc instruction
+      // does not.
+
+      // Converting from a lower rank to a higher rank can never have
+      // undefined behavior, since higher-rank types must have a superset
+      // of values of lower-rank types.
+      if (CGF.getContext().getFloatingTypeOrder(OrigSrcType, DstType) != 1)
+        return;
+
+      assert(!OrigSrcType->isHalfType() &&
+             "should not check conversion from __half, it has the lowest rank");
+
+      const llvm::fltSemantics &DstSema =
+        CGF.getContext().getFloatTypeSemantics(DstType);
+      APFloat MinBad = APFloat::getLargest(DstSema, false);
+      APFloat MaxBad = APFloat::getInf(DstSema, false);
+
+      bool IsInexact;
+      MinBad.convert(SrcSema, APFloat::rmTowardZero, &IsInexact);
+      MaxBad.convert(SrcSema, APFloat::rmTowardZero, &IsInexact);
+
+      Value *AbsSrc = CGF.EmitNounwindRuntimeCall(
+        CGF.CGM.getIntrinsic(llvm::Intrinsic::fabs, Src->getType()), Src);
       llvm::Value *GE =
-        Builder.CreateFCmpOGE(Src, llvm::ConstantFP::get(VMContext, MinSrc));
+        Builder.CreateFCmpOGT(AbsSrc, llvm::ConstantFP::get(VMContext, MinBad));
       llvm::Value *LE =
-        Builder.CreateFCmpOLE(Src, llvm::ConstantFP::get(VMContext, MaxSrc));
-      Check = Builder.CreateAnd(GE, LE);
+        Builder.CreateFCmpOLT(AbsSrc, llvm::ConstantFP::get(VMContext, MaxBad));
+      Check = Builder.CreateNot(Builder.CreateAnd(GE, LE));
     }
   }
 
diff --git a/test/CodeGen/catch-undef-behavior.c b/test/CodeGen/catch-undef-behavior.c
index 8e3388e..ebe39fe 100644
--- a/test/CodeGen/catch-undef-behavior.c
+++ b/test/CodeGen/catch-undef-behavior.c
@@ -364,14 +364,17 @@
 // CHECK: @float_float_overflow
 // CHECK-TRAP: @float_float_overflow
 float float_float_overflow(double f) {
-  // CHECK: %[[GE:.*]] = fcmp oge double %[[F:.*]], 0xC7EFFFFFE0000000
-  // CHECK: %[[LE:.*]] = fcmp ole double %[[F]], 0x47EFFFFFE0000000
+  // CHECK: %[[F:.*]] = call double @llvm.fabs.f64(
+  // CHECK: %[[GE:.*]] = fcmp ogt double %[[F]], 0x47EFFFFFE0000000
+  // CHECK: %[[LE:.*]] = fcmp olt double %[[F]], 0x7FF0000000000000
   // CHECK: and i1 %[[GE]], %[[LE]]
   // CHECK: call void @__ubsan_handle_float_cast_overflow(
 
-  // CHECK-TRAP: %[[GE:.*]] = fcmp oge double %[[F:.*]], 0xC7EFFFFFE0000000
-  // CHECK-TRAP: %[[LE:.*]] = fcmp ole double %[[F]], 0x47EFFFFFE0000000
-  // CHECK-TRAP: %[[INBOUNDS:.*]] = and i1 %[[GE]], %[[LE]]
+  // CHECK-TRAP: %[[F:.*]] = call double @llvm.fabs.f64(
+  // CHECK-TRAP: %[[GE:.*]] = fcmp ogt double %[[F]], 0x47EFFFFFE0000000
+  // CHECK-TRAP: %[[LE:.*]] = fcmp olt double %[[F]], 0x7FF0000000000000
+  // CHECK-TRAP: %[[OUTOFBOUNDS:.*]] = and i1 %[[GE]], %[[LE]]
+  // CHECK-TRAP: %[[INBOUNDS:.*]] = xor i1 %[[OUTOFBOUNDS]], true
   // CHECK-TRAP-NEXT: br i1 %[[INBOUNDS]]
 
   // CHECK-TRAP:      call void @llvm.trap() [[NR_NUW]]