Index: include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- include/clang/Basic/DiagnosticSemaKinds.td
+++ include/clang/Basic/DiagnosticSemaKinds.td
@@ -3362,6 +3362,9 @@
   InGroup<IgnoredAttributes>;
 def warn_attribute_unknown_visibility : Warning<"unknown visibility %0">,
   InGroup<IgnoredAttributes>;
+def warn_attribute_hidden_visibility :
+  Warning<"'hidden' visibility on %select{function|variable}0 with incompatible language attribute will be ignored">,
+  InGroup<IgnoredAttributes>;
 def warn_attribute_protected_visibility :
   Warning<"target does not support 'protected' visibility; using 'default'">,
   InGroup<DiagGroup<"unsupported-visibility">>;
Index: include/clang/Basic/Visibility.h
===================================================================
--- include/clang/Basic/Visibility.h
+++ include/clang/Basic/Visibility.h
@@ -53,8 +53,6 @@
   uint8_t linkage_    : 3;
   uint8_t visibility_ : 2;
   uint8_t explicit_   : 1;
-
-  void setVisibility(Visibility V, bool E) { visibility_ = V; explicit_ = E; }
 public:
   LinkageInfo() : linkage_(ExternalLinkage), visibility_(DefaultVisibility),
                   explicit_(false) {}
@@ -86,6 +84,8 @@
 
   void setLinkage(Linkage L) { linkage_ = L; }
 
+  void setVisibility(Visibility V, bool E) { visibility_ = V; explicit_ = E; }
+
   void mergeLinkage(Linkage L) {
     setLinkage(minLinkage(getLinkage(), L));
   }
Index: lib/AST/Decl.cpp
===================================================================
--- lib/AST/Decl.cpp
+++ lib/AST/Decl.cpp
@@ -731,6 +731,19 @@
     }
   }
 
+  // We consider OpenCL kernels, __global__ Cuda functions, and __device__ Cuda
+  // variables to have explicit visibility of "non-hidden".
+  if (D->hasAttr<OpenCLKernelAttr>() ||
+      (isa<FunctionDecl>(D) && D->hasAttr<CUDAGlobalAttr>()) ||
+      (isa<VarDecl>(D) && D->hasAttr<CUDADeviceAttr>())) {
+    Visibility Vis = LV.getVisibility();
+    if (LV.getVisibility() == HiddenVisibility)
+      Vis = Context.getTargetInfo().hasProtectedVisibility()
+                ? ProtectedVisibility
+                : DefaultVisibility;
+    LV.setVisibility(Vis, true);
+  }
+
   // C++ [basic.link]p4:
 
   //   A name having namespace scope has external linkage if it is the
Index: lib/CodeGen/TargetInfo.cpp
===================================================================
--- lib/CodeGen/TargetInfo.cpp
+++ lib/CodeGen/TargetInfo.cpp
@@ -7840,23 +7840,8 @@
 };
 }
 
-static bool requiresAMDGPUProtectedVisibility(const Decl *D,
-                                              llvm::GlobalValue *GV) {
-  if (GV->getVisibility() != llvm::GlobalValue::HiddenVisibility)
-    return false;
-
-  return D->hasAttr<OpenCLKernelAttr>() ||
-         (isa<FunctionDecl>(D) && D->hasAttr<CUDAGlobalAttr>()) ||
-         (isa<VarDecl>(D) && D->hasAttr<CUDADeviceAttr>());
-}
-
 void AMDGPUTargetCodeGenInfo::setTargetAttributes(
     const Decl *D, llvm::GlobalValue *GV, CodeGen::CodeGenModule &M) const {
-  if (requiresAMDGPUProtectedVisibility(D, GV)) {
-    GV->setVisibility(llvm::GlobalValue::ProtectedVisibility);
-    GV->setDSOLocal(true);
-  }
-
   if (GV->isDeclaration())
     return;
   const FunctionDecl *FD = dyn_cast_or_null<FunctionDecl>(D);
Index: lib/Sema/SemaDeclAttr.cpp
===================================================================
--- lib/Sema/SemaDeclAttr.cpp
+++ lib/Sema/SemaDeclAttr.cpp
@@ -7375,6 +7375,16 @@
     Diag(D->getLocation(), diag::err_designated_init_attr_non_init);
     D->dropAttr<ObjCDesignatedInitializerAttr>();
   }
+
+  if ((D->hasAttr<VisibilityAttr>() &&
+       D->getAttr<VisibilityAttr>()->getVisibility() ==
+           VisibilityAttr::Hidden) &&
+      (D->hasAttr<OpenCLKernelAttr>() ||
+       (isa<FunctionDecl>(D) && D->hasAttr<CUDAGlobalAttr>()) ||
+       (isa<VarDecl>(D) && D->hasAttr<CUDADeviceAttr>()))) {
+    Diag(D->getLocation(), diag::warn_attribute_hidden_visibility)
+        << (isa<FunctionDecl>(D) ? 0 : 1);
+  }
 }
 
 // Helper for delayed processing TransparentUnion attribute.
Index: test/SemaCUDA/visibility-diagnostics.cu
===================================================================
--- /dev/null
+++ test/SemaCUDA/visibility-diagnostics.cu
@@ -0,0 +1,13 @@
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify %s
+
+#include "Inputs/cuda.h"
+
+__attribute__((visibility("hidden"))) __global__ void global_func_hidden() {} // expected-warning {{'hidden' visibility on function with incompatible language attribute will be ignored}}
+__attribute__((visibility("protected"))) __global__ void global_func_protected() {}
+__attribute__((visibility("default"))) __global__ void global_func_default() {}
+__global__ void global_func() {}
+
+__attribute__((visibility("hidden"))) __device__ int device_var_hidden; // expected-warning {{'hidden' visibility on variable with incompatible language attribute will be ignored}}
+__attribute__((visibility("protected"))) __device__ int device_var_protected;
+__attribute__((visibility("default"))) __device__ int device_var_default;
+__device__ int device_var;
Index: test/SemaOpenCL/visibility-diagnostics.cl
===================================================================
--- /dev/null
+++ test/SemaOpenCL/visibility-diagnostics.cl
@@ -0,0 +1,11 @@
+// RUN: %clang_cc1 -std=cl2.0 -verify -pedantic -fsyntax-only %s
+
+__attribute__((visibility("hidden"))) kernel void kern_hidden() {} // expected-warning {{'hidden' visibility on function with incompatible language attribute will be ignored}}
+__attribute__((visibility("protected"))) kernel void kern_protected();
+__attribute__((visibility("default"))) kernel void kern_default();
+kernel void kern();
+
+__attribute__((visibility("hidden"))) extern kernel void ext_kern_hidden(); // expected-warning {{'hidden' visibility on function with incompatible language attribute will be ignored}}
+__attribute__((visibility("protected"))) extern kernel void ext_kern_protected();
+__attribute__((visibility("default"))) extern kernel void ext_kern_default();
+extern kernel void ext_kern();