Index: lib/Sema/SemaDeclAttr.cpp
===================================================================
--- lib/Sema/SemaDeclAttr.cpp
+++ lib/Sema/SemaDeclAttr.cpp
@@ -348,6 +348,26 @@
                                         Attr.getAttributeSpellingListIndex()));
 }
 
+template <typename AttrType>
+static void handleSimpleAttributeWithExclusions(Sema &S, Decl *D,
+                                                const AttributeList &Attr) {
+  handleSimpleAttribute<AttrType>(S, D, Attr);
+}
+
+/// \brief Applies the given attribute to the Decl so long as the Decl doesn't
+/// already have one of the given incompatible attributes.
+template <typename AttrType, typename IncompatibleAttrType,
+          typename... IncompatibleAttrTypes>
+static void handleSimpleAttributeWithExclusions(Sema &S, Decl *D,
+                                                const AttributeList &Attr) {
+  if (checkAttrMutualExclusion<IncompatibleAttrType>(S, D, Attr.getRange(),
+                                                     Attr.getName())) {
+    return;
+  }
+  handleSimpleAttributeWithExclusions<AttrType, IncompatibleAttrTypes...>(S, D,
+                                                                          Attr);
+}
+
 /// \brief Check if the passed-in expression is of type int or bool.
 static bool isIntOrBool(Expr *Exp) {
   QualType QT = Exp->getType();
@@ -3588,6 +3608,12 @@
 }
 
 static void handleGlobalAttr(Sema &S, Decl *D, const AttributeList &Attr) {
+  if (checkAttrMutualExclusion<CUDADeviceAttr>(S, D, Attr.getRange(),
+                                               Attr.getName()) ||
+      checkAttrMutualExclusion<CUDAHostAttr>(S, D, Attr.getRange(),
+                                             Attr.getName())) {
+    return;
+  }
   FunctionDecl *FD = cast<FunctionDecl>(D);
   if (!FD->getReturnType()->isVoidType()) {
     SourceRange RTRange = FD->getReturnTypeSourceRange();
@@ -4558,14 +4584,6 @@
     handleARMInterruptAttr(S, D, Attr);
 }
 
-static void handleMips16Attribute(Sema &S, Decl *D, const AttributeList &Attr) {
-  if (checkAttrMutualExclusion<MipsInterruptAttr>(S, D, Attr.getRange(),
-                                                  Attr.getName()))
-    return;
-
-  handleSimpleAttribute<Mips16Attr>(S, D, Attr);
-}
-
 static void handleAMDGPUNumVGPRAttr(Sema &S, Decl *D,
                                     const AttributeList &Attr) {
   uint32_t NumRegs;
@@ -4955,7 +4973,8 @@
     handleDLLAttr(S, D, Attr);
     break;
   case AttributeList::AT_Mips16:
-    handleMips16Attribute(S, D, Attr);
+    handleSimpleAttributeWithExclusions<Mips16Attr, MipsInterruptAttr>(S, D,
+                                                                       Attr);
     break;
   case AttributeList::AT_NoMips16:
     handleSimpleAttribute<NoMips16Attr>(S, D, Attr);
@@ -5006,7 +5025,8 @@
     handleCommonAttr(S, D, Attr);
     break;
   case AttributeList::AT_CUDAConstant:
-    handleSimpleAttribute<CUDAConstantAttr>(S, D, Attr);
+    handleSimpleAttributeWithExclusions<CUDAConstantAttr, CUDASharedAttr>(S, D,
+                                                                          Attr);
     break;
   case AttributeList::AT_PassObjectSize:
     handlePassObjectSizeAttr(S, D, Attr);
@@ -5051,10 +5071,12 @@
     handleGlobalAttr(S, D, Attr);
     break;
   case AttributeList::AT_CUDADevice:
-    handleSimpleAttribute<CUDADeviceAttr>(S, D, Attr);
+    handleSimpleAttributeWithExclusions<CUDADeviceAttr, CUDAGlobalAttr>(S, D,
+                                                                        Attr);
     break;
   case AttributeList::AT_CUDAHost:
-    handleSimpleAttribute<CUDAHostAttr>(S, D, Attr);
+    handleSimpleAttributeWithExclusions<CUDAHostAttr, CUDAGlobalAttr>(S, D,
+                                                                      Attr);
     break;
   case AttributeList::AT_GNUInline:
     handleGNUInlineAttr(S, D, Attr);
@@ -5114,7 +5136,8 @@
     handleSimpleAttribute<NoThrowAttr>(S, D, Attr);
     break;
   case AttributeList::AT_CUDAShared:
-    handleSimpleAttribute<CUDASharedAttr>(S, D, Attr);
+    handleSimpleAttributeWithExclusions<CUDASharedAttr, CUDAConstantAttr>(S, D,
+                                                                          Attr);
     break;
   case AttributeList::AT_VecReturn:
     handleVecReturnAttr(S, D, Attr);
Index: test/SemaCUDA/attributes-on-non-cuda.cu
===================================================================
--- test/SemaCUDA/attributes-on-non-cuda.cu
+++ test/SemaCUDA/attributes-on-non-cuda.cu
@@ -1,4 +1,5 @@
-// Tests handling of CUDA attributes.
+// Tests that CUDA attributes are warnings when compiling C files, but not when
+// compiling CUDA files.
 //
 // RUN: %clang_cc1 -fsyntax-only -verify %s
 // RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s
Index: test/SemaCUDA/bad-attributes.cu
===================================================================
--- /dev/null
+++ test/SemaCUDA/bad-attributes.cu
@@ -0,0 +1,65 @@
+// Tests handling of CUDA attributes that are bad either because they're
+// applied to the wrong sort of thing, or because they're given in illegal
+// combinations.
+//
+// RUN: %clang_cc1 -fsyntax-only -verify %s
+// RUN: %clang_cc1 -fcuda-is-device -fsyntax-only -verify %s
+
+// TODO: Check __cudart_builtin__, __nv_weak__,
+// __device_builtin_surface_type__, __device_builtin_texture_type__.
+
+// Make this file work with nvcc, for testing compatibility.
+#ifndef __NVCC__
+#define __device__ __attribute__((device))
+#define __host__ __attribute__((host))
+#define __constant__ __attribute__((constant))
+#define __shared__ __attribute__((shared))
+#define __global__ __attribute__((global))
+#define __device_builtin__ __attribute__((device_builtin))
+#endif
+
+// Try applying attributes to functions and variables.  Some should generate
+// warnings; others not.
+__device__ int a1;
+__device__ void a2();
+__host__ int b1; // expected-warning {{attribute only applies to functions}}
+__host__ void b2();
+__constant__ int c1;
+__constant__ void c2(); // expected-warning {{attribute only applies to variables}}
+__shared__ int d1;
+__shared__ void d2(); // expected-warning {{attribute only applies to variables}}
+__global__ int e1; // expected-warning {{attribute only applies to functions}}
+__global__ void e2();
+__device_builtin__ int f1;
+__device_builtin__ void f2();
+
+// Try all pairs of attributes which can be present on a function or a
+// variable.  Check both orderings of the attributes, as that can matter in
+// clang.
+__device__ __host__ void z1();
+__device__ __constant__ int z2;
+__device__ __shared__ int z3;
+__device__ __global__ void z4(); // expected-error {{attributes are not compatible}}
+// expected-note@-1 {{conflicting attribute is here}}
+__device__ __device_builtin__ void z5();
+
+__host__ __device__ void z6();
+__host__ __global__ void z7();  // expected-error {{attributes are not compatible}}
+// expected-note@-1 {{conflicting attribute is here}}
+__host__ __device_builtin__ void z8();
+
+__constant__ __device__ int z9;
+__constant__ __shared__ int z10;  // expected-error {{attributes are not compatible}}
+// expected-note@-1 {{conflicting attribute is here}}
+__constant__ __device_builtin__ int z11;
+
+__shared__ __device__ int z12;
+__shared__ __constant__ int z13;  // expected-error {{attributes are not compatible}}
+// expected-note@-1 {{conflicting attribute is here}}
+__shared__ __device_builtin__ int z14;
+
+__global__ __device__ void z15();  // expected-error {{attributes are not compatible}}
+// expected-note@-1 {{conflicting attribute is here}}
+__global__ __host__ void z16();  // expected-error {{attributes are not compatible}}
+// expected-note@-1 {{conflicting attribute is here}}
+__global__ __device_builtin__ void z17();