Index: include/clang/Basic/Builtins.def
===================================================================
--- include/clang/Basic/Builtins.def
+++ include/clang/Basic/Builtins.def
@@ -107,9 +107,9 @@
 BUILTIN(__builtin_copysign, "ddd", "ncF")
 BUILTIN(__builtin_copysignf, "fff", "ncF")
 BUILTIN(__builtin_copysignl, "LdLdLd", "ncF")
-BUILTIN(__builtin_fabs , "dd"  , "ncF")
-BUILTIN(__builtin_fabsf, "ff"  , "ncF")
-BUILTIN(__builtin_fabsl, "LdLd", "ncF")
+BUILTIN(__builtin_fabs , "dd"  , "nc")
+BUILTIN(__builtin_fabsf, "ff"  , "nc")
+BUILTIN(__builtin_fabsl, "LdLd", "nc")
 BUILTIN(__builtin_fmod , "ddd"  , "Fnc")
 BUILTIN(__builtin_fmodf, "fff"  , "Fnc")
 BUILTIN(__builtin_fmodl, "LdLdLd", "Fnc")
@@ -165,9 +165,9 @@
 BUILTIN(__builtin_cbrt , "dd", "Fnc")
 BUILTIN(__builtin_cbrtf, "ff", "Fnc")
 BUILTIN(__builtin_cbrtl, "LdLd", "Fnc")
-BUILTIN(__builtin_ceil , "dd"  , "Fnc")
-BUILTIN(__builtin_ceilf, "ff"  , "Fnc")
-BUILTIN(__builtin_ceill, "LdLd", "Fnc")
+BUILTIN(__builtin_ceil , "dd"  , "nc")
+BUILTIN(__builtin_ceilf, "ff"  , "nc")
+BUILTIN(__builtin_ceill, "LdLd", "nc")
 BUILTIN(__builtin_cos , "dd"  , "Fnc")
 BUILTIN(__builtin_cosf, "ff"  , "Fnc")
 BUILTIN(__builtin_cosh , "dd"  , "Fnc")
@@ -192,18 +192,18 @@
 BUILTIN(__builtin_fdim, "ddd", "Fnc")
 BUILTIN(__builtin_fdimf, "fff", "Fnc")
 BUILTIN(__builtin_fdiml, "LdLdLd", "Fnc")
-BUILTIN(__builtin_floor , "dd"  , "Fnc")
-BUILTIN(__builtin_floorf, "ff"  , "Fnc")
-BUILTIN(__builtin_floorl, "LdLd", "Fnc")
+BUILTIN(__builtin_floor , "dd"  , "nc")
+BUILTIN(__builtin_floorf, "ff"  , "nc")
+BUILTIN(__builtin_floorl, "LdLd", "nc")
 BUILTIN(__builtin_fma, "dddd", "Fnc")
 BUILTIN(__builtin_fmaf, "ffff", "Fnc")
 BUILTIN(__builtin_fmal, "LdLdLdLd", "Fnc")
-BUILTIN(__builtin_fmax, "ddd", "Fnc")
-BUILTIN(__builtin_fmaxf, "fff", "Fnc")
-BUILTIN(__builtin_fmaxl, "LdLdLd", "Fnc")
-BUILTIN(__builtin_fmin, "ddd", "Fnc")
-BUILTIN(__builtin_fminf, "fff", "Fnc")
-BUILTIN(__builtin_fminl, "LdLdLd", "Fnc")
+BUILTIN(__builtin_fmax, "ddd", "nc")
+BUILTIN(__builtin_fmaxf, "fff", "nc")
+BUILTIN(__builtin_fmaxl, "LdLdLd", "nc")
+BUILTIN(__builtin_fmin, "ddd", "nc")
+BUILTIN(__builtin_fminf, "fff", "nc")
+BUILTIN(__builtin_fminl, "LdLdLd", "nc")
 BUILTIN(__builtin_hypot , "ddd"  , "Fnc")
 BUILTIN(__builtin_hypotf, "fff"  , "Fnc")
 BUILTIN(__builtin_hypotl, "LdLdLd", "Fnc")
@@ -240,9 +240,9 @@
 BUILTIN(__builtin_lround , "Lid", "Fnc")
 BUILTIN(__builtin_lroundf, "Lif", "Fnc")
 BUILTIN(__builtin_lroundl, "LiLd", "Fnc")
-BUILTIN(__builtin_nearbyint , "dd", "Fnc")
-BUILTIN(__builtin_nearbyintf, "ff", "Fnc")
-BUILTIN(__builtin_nearbyintl, "LdLd", "Fnc")
+BUILTIN(__builtin_nearbyint , "dd", "nc")
+BUILTIN(__builtin_nearbyintf, "ff", "nc")
+BUILTIN(__builtin_nearbyintl, "LdLd", "nc")
 BUILTIN(__builtin_nextafter , "ddd", "Fnc")
 BUILTIN(__builtin_nextafterf, "fff", "Fnc")
 BUILTIN(__builtin_nextafterl, "LdLdLd", "Fnc")
@@ -255,12 +255,12 @@
 BUILTIN(__builtin_remquo , "dddi*", "Fn")
 BUILTIN(__builtin_remquof, "fffi*", "Fn")
 BUILTIN(__builtin_remquol, "LdLdLdi*", "Fn")
-BUILTIN(__builtin_rint , "dd", "Fnc")
-BUILTIN(__builtin_rintf, "ff", "Fnc")
-BUILTIN(__builtin_rintl, "LdLd", "Fnc")
-BUILTIN(__builtin_round, "dd"  , "Fnc")
-BUILTIN(__builtin_roundf, "ff"  , "Fnc")
-BUILTIN(__builtin_roundl, "LdLd"  , "Fnc")
+BUILTIN(__builtin_rint , "dd", "nc")
+BUILTIN(__builtin_rintf, "ff", "nc")
+BUILTIN(__builtin_rintl, "LdLd", "nc")
+BUILTIN(__builtin_round, "dd"  , "nc")
+BUILTIN(__builtin_roundf, "ff"  , "nc")
+BUILTIN(__builtin_roundl, "LdLd"  , "nc")
 BUILTIN(__builtin_scalbln , "ddLi", "Fnc")
 BUILTIN(__builtin_scalblnf, "ffLi", "Fnc")
 BUILTIN(__builtin_scalblnl, "LdLdLi", "Fnc")
@@ -285,9 +285,9 @@
 BUILTIN(__builtin_tgamma , "dd", "Fnc")
 BUILTIN(__builtin_tgammaf, "ff", "Fnc")
 BUILTIN(__builtin_tgammal, "LdLd", "Fnc")
-BUILTIN(__builtin_trunc , "dd", "Fnc")
-BUILTIN(__builtin_truncf, "ff", "Fnc")
-BUILTIN(__builtin_truncl, "LdLd", "Fnc")
+BUILTIN(__builtin_trunc , "dd", "nc")
+BUILTIN(__builtin_truncf, "ff", "nc")
+BUILTIN(__builtin_truncl, "LdLd", "nc")
 
 // C99 complex builtins
 BUILTIN(__builtin_cabs, "dXd", "Fnc")
Index: lib/CodeGen/CGBuiltin.cpp
===================================================================
--- lib/CodeGen/CGBuiltin.cpp
+++ lib/CodeGen/CGBuiltin.cpp
@@ -217,6 +217,51 @@
                        ValueType);
 }
 
+// Emit a simple mangled intrinsic that has 1 argument and a return type
+// matching the argument type.
+static Value *emitUnaryBuiltin(CodeGenFunction &CGF,
+                               const CallExpr *E,
+                               unsigned IntrinsicID) {
+  llvm::Value *Src0 = CGF.EmitScalarExpr(E->getArg(0));
+
+  Value *F = CGF.CGM.getIntrinsic(IntrinsicID, Src0->getType());
+  return CGF.Builder.CreateCall(F, Src0);
+}
+
+// Emit an intrinsic that has 2 operands of the same type as its result.
+static Value *emitBinaryBuiltin(CodeGenFunction &CGF,
+                                const CallExpr *E,
+                                unsigned IntrinsicID) {
+  llvm::Value *Src0 = CGF.EmitScalarExpr(E->getArg(0));
+  llvm::Value *Src1 = CGF.EmitScalarExpr(E->getArg(1));
+
+  Value *F = CGF.CGM.getIntrinsic(IntrinsicID, Src0->getType());
+  return CGF.Builder.CreateCall(F, { Src0, Src1 });
+}
+
+// Emit an intrinsic that has 3 operands of the same type as its result.
+static Value *emitTernaryBuiltin(CodeGenFunction &CGF,
+                                 const CallExpr *E,
+                                 unsigned IntrinsicID) {
+  llvm::Value *Src0 = CGF.EmitScalarExpr(E->getArg(0));
+  llvm::Value *Src1 = CGF.EmitScalarExpr(E->getArg(1));
+  llvm::Value *Src2 = CGF.EmitScalarExpr(E->getArg(2));
+
+  Value *F = CGF.CGM.getIntrinsic(IntrinsicID, Src0->getType());
+  return CGF.Builder.CreateCall(F, { Src0, Src1, Src2 });
+}
+
+// Emit an intrinsic that has 1 float or double operand, and 1 integer.
+static Value *emitFPIntBuiltin(CodeGenFunction &CGF,
+                               const CallExpr *E,
+                               unsigned IntrinsicID) {
+  llvm::Value *Src0 = CGF.EmitScalarExpr(E->getArg(0));
+  llvm::Value *Src1 = CGF.EmitScalarExpr(E->getArg(1));
+
+  Value *F = CGF.CGM.getIntrinsic(IntrinsicID, Src0->getType());
+  return CGF.Builder.CreateCall(F, {Src0, Src1});
+}
+
 /// EmitFAbs - Emit a call to @llvm.fabs().
 static Value *EmitFAbs(CodeGenFunction &CGF, Value *V) {
   Value *F = CGF.CGM.getIntrinsic(Intrinsic::fabs, V->getType());
@@ -286,40 +331,6 @@
   return CGF.Builder.CreateExtractValue(Tmp, 0);
 }
 
-// Emit a simple mangled intrinsic that has 1 argument and a return type
-// matching the argument type.
-static Value *emitUnaryBuiltin(CodeGenFunction &CGF,
-                               const CallExpr *E,
-                               unsigned IntrinsicID) {
-  llvm::Value *Src0 = CGF.EmitScalarExpr(E->getArg(0));
-
-  Value *F = CGF.CGM.getIntrinsic(IntrinsicID, Src0->getType());
-  return CGF.Builder.CreateCall(F, Src0);
-}
-
-// Emit an intrinsic that has 3 float or double operands.
-static Value *emitTernaryFPBuiltin(CodeGenFunction &CGF,
-                                   const CallExpr *E,
-                                   unsigned IntrinsicID) {
-  llvm::Value *Src0 = CGF.EmitScalarExpr(E->getArg(0));
-  llvm::Value *Src1 = CGF.EmitScalarExpr(E->getArg(1));
-  llvm::Value *Src2 = CGF.EmitScalarExpr(E->getArg(2));
-
-  Value *F = CGF.CGM.getIntrinsic(IntrinsicID, Src0->getType());
-  return CGF.Builder.CreateCall(F, {Src0, Src1, Src2});
-}
-
-// Emit an intrinsic that has 1 float or double operand, and 1 integer.
-static Value *emitFPIntBuiltin(CodeGenFunction &CGF,
-                               const CallExpr *E,
-                               unsigned IntrinsicID) {
-  llvm::Value *Src0 = CGF.EmitScalarExpr(E->getArg(0));
-  llvm::Value *Src1 = CGF.EmitScalarExpr(E->getArg(1));
-
-  Value *F = CGF.CGM.getIntrinsic(IntrinsicID, Src0->getType());
-  return CGF.Builder.CreateCall(F, {Src0, Src1});
-}
-
 namespace {
   struct WidthAndSignedness {
     unsigned Width;
@@ -497,9 +508,7 @@
   case Builtin::BI__builtin_fabs:
   case Builtin::BI__builtin_fabsf:
   case Builtin::BI__builtin_fabsl: {
-    Value *Arg1 = EmitScalarExpr(E->getArg(0));
-    Value *Result = EmitFAbs(*this, Arg1);
-    return RValue::get(Result);
+    return RValue::get(emitUnaryBuiltin(*this, E, Intrinsic::fabs));
   }
   case Builtin::BI__builtin_fmod:
   case Builtin::BI__builtin_fmodf:
@@ -509,7 +518,51 @@
     Value *Result = Builder.CreateFRem(Arg1, Arg2, "fmod");
     return RValue::get(Result);
   }
-
+  case Builtin::BI__builtin_copysign:
+  case Builtin::BI__builtin_copysignf:
+  case Builtin::BI__builtin_copysignl: {
+    return RValue::get(emitBinaryBuiltin(*this, E, Intrinsic::copysign));
+  }
+  case Builtin::BI__builtin_ceil:
+  case Builtin::BI__builtin_ceilf:
+  case Builtin::BI__builtin_ceill: {
+    return RValue::get(emitUnaryBuiltin(*this, E, Intrinsic::ceil));
+  }
+  case Builtin::BI__builtin_floor:
+  case Builtin::BI__builtin_floorf:
+  case Builtin::BI__builtin_floorl: {
+    return RValue::get(emitUnaryBuiltin(*this, E, Intrinsic::floor));
+  }
+  case Builtin::BI__builtin_trunc:
+  case Builtin::BI__builtin_truncf:
+  case Builtin::BI__builtin_truncl: {
+    return RValue::get(emitUnaryBuiltin(*this, E, Intrinsic::trunc));
+  }
+  case Builtin::BI__builtin_rint:
+  case Builtin::BI__builtin_rintf:
+  case Builtin::BI__builtin_rintl: {
+    return RValue::get(emitUnaryBuiltin(*this, E, Intrinsic::rint));
+  }
+  case Builtin::BI__builtin_nearbyint:
+  case Builtin::BI__builtin_nearbyintf:
+  case Builtin::BI__builtin_nearbyintl: {
+    return RValue::get(emitUnaryBuiltin(*this, E, Intrinsic::nearbyint));
+  }
+  case Builtin::BI__builtin_round:
+  case Builtin::BI__builtin_roundf:
+  case Builtin::BI__builtin_roundl: {
+    return RValue::get(emitUnaryBuiltin(*this, E, Intrinsic::round));
+  }
+  case Builtin::BI__builtin_fmin:
+  case Builtin::BI__builtin_fminf:
+  case Builtin::BI__builtin_fminl: {
+    return RValue::get(emitBinaryBuiltin(*this, E, Intrinsic::minnum));
+  }
+  case Builtin::BI__builtin_fmax:
+  case Builtin::BI__builtin_fmaxf:
+  case Builtin::BI__builtin_fmaxl: {
+    return RValue::get(emitBinaryBuiltin(*this, E, Intrinsic::maxnum));
+  }
   case Builtin::BI__builtin_conj:
   case Builtin::BI__builtin_conjf:
   case Builtin::BI__builtin_conjl: {
@@ -7326,7 +7379,7 @@
   }
   case AMDGPU::BI__builtin_amdgcn_div_fixup:
   case AMDGPU::BI__builtin_amdgcn_div_fixupf:
-    return emitTernaryFPBuiltin(*this, E, Intrinsic::amdgcn_div_fixup);
+    return emitTernaryBuiltin(*this, E, Intrinsic::amdgcn_div_fixup);
   case AMDGPU::BI__builtin_amdgcn_trig_preop:
   case AMDGPU::BI__builtin_amdgcn_trig_preopf:
     return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_trig_preop);
Index: test/CodeGen/builtins.c
===================================================================
--- test/CodeGen/builtins.c
+++ test/CodeGen/builtins.c
@@ -257,6 +257,98 @@
   // CHECK: call float @llvm.canonicalize.f32(float
   // CHECK: call double @llvm.canonicalize.f64(double
   // CHECK: call x86_fp80 @llvm.canonicalize.f80(x86_fp80
+
+  resf = __builtin_fminf(F, F);
+  // CHECK: call float @llvm.minnum.f32
+
+  resd = __builtin_fmin(D, D);
+  // CHECK: call double @llvm.minnum.f64
+
+  resld = __builtin_fminl(LD, LD);
+  // CHECK: call x86_fp80 @llvm.minnum.f80
+
+  resf = __builtin_fmaxf(F, F);
+  // CHECK: call float @llvm.maxnum.f32
+
+  resd = __builtin_fmax(D, D);
+  // CHECK: call double @llvm.maxnum.f64
+
+  resld = __builtin_fmaxl(LD, LD);
+  // CHECK: call x86_fp80 @llvm.maxnum.f80
+
+  resf = __builtin_fabsf(F);
+  // CHECK: call float @llvm.fabs.f32
+
+  resd = __builtin_fabs(D);
+  // CHECK: call double @llvm.fabs.f64
+
+  resld = __builtin_fabsl(LD);
+  // CHECK: call x86_fp80 @llvm.fabs.f80
+
+  resf = __builtin_copysignf(F, F);
+  // CHECK: call float @llvm.copysign.f32
+
+  resd = __builtin_copysign(D, D);
+  // CHECK: call double @llvm.copysign.f64
+
+  resld = __builtin_copysignl(LD, LD);
+  // CHECK: call x86_fp80 @llvm.copysign.f80
+
+
+  resf = __builtin_ceilf(F);
+  // CHECK: call float @llvm.ceil.f32
+
+  resd = __builtin_ceil(D);
+  // CHECK: call double @llvm.ceil.f64
+
+  resld = __builtin_ceill(LD);
+  // CHECK: call x86_fp80 @llvm.ceil.f80
+
+  resf = __builtin_floorf(F);
+  // CHECK: call float @llvm.floor.f32
+
+  resd = __builtin_floor(D);
+  // CHECK: call double @llvm.floor.f64
+
+  resld = __builtin_floorl(LD);
+  // CHECK: call x86_fp80 @llvm.floor.f80
+
+  resf = __builtin_truncf(F);
+  // CHECK: call float @llvm.trunc.f32
+
+  resd = __builtin_trunc(D);
+  // CHECK: call double @llvm.trunc.f64
+
+  resld = __builtin_truncl(LD);
+  // CHECK: call x86_fp80 @llvm.trunc.f80
+
+  resf = __builtin_rintf(F);
+  // CHECK: call float @llvm.rint.f32
+
+  resd = __builtin_rint(D);
+  // CHECK: call double @llvm.rint.f64
+
+  resld = __builtin_rintl(LD);
+  // CHECK: call x86_fp80 @llvm.rint.f80
+
+  resf = __builtin_nearbyintf(F);
+  // CHECK: call float @llvm.nearbyint.f32
+
+  resd = __builtin_nearbyint(D);
+  // CHECK: call double @llvm.nearbyint.f64
+
+  resld = __builtin_nearbyintl(LD);
+  // CHECK: call x86_fp80 @llvm.nearbyint.f80
+
+  resf = __builtin_roundf(F);
+  // CHECK: call float @llvm.round.f32
+
+  resd = __builtin_round(D);
+  // CHECK: call double @llvm.round.f64
+
+  resld = __builtin_roundl(LD);
+  // CHECK: call x86_fp80 @llvm.round.f80
+
 }
 
 // __builtin_longjmp isn't supported on all platforms, so only test it on X86.