diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index 6a66ecf6f94c1..d744f9a8c5109 100644
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -792,6 +792,7 @@ void Sema::MaybeAddCUDAConstantAttr(VarDecl *VD) {
       (VD->isFileVarDecl() || VD->isStaticDataMember()) &&
       !IsDependentVar(VD) &&
       ((VD->isConstexpr() || VD->getType().isConstQualified()) &&
+       VD->getStorageClass() != SC_Extern &&
        HasAllowedCUDADeviceStaticInitializer(*this, VD,
                                              CICK_DeviceOrConstant))) {
     VD->addAttr(CUDAConstantAttr::CreateImplicit(getASTContext()));
diff --git a/clang/test/CodeGenCUDA/device-use-host-var.cu b/clang/test/CodeGenCUDA/device-use-host-var.cu
index 64de57e41b4b9..65aac2b7eca89 100644
--- a/clang/test/CodeGenCUDA/device-use-host-var.cu
+++ b/clang/test/CodeGenCUDA/device-use-host-var.cu
@@ -2,6 +2,8 @@
 // RUN:   -fcuda-is-device -emit-llvm -o - -x hip %s | FileCheck %s
 // RUN: %clang_cc1 -std=c++14 -triple amdgcn-amd-amdhsa \
 // RUN:   -fcuda-is-device -emit-llvm -o - -x hip %s | FileCheck -check-prefix=NEG %s
+// RUN: %clang_cc1 -std=c++14 -triple nvptx64-nvidia-cuda  \
+// RUN:   -fcuda-is-device -emit-llvm -o - -x hip %s | FileCheck -check-prefix=NEG -implicit-check-not=external_ %s
 
 #include "Inputs/cuda.h"
 
@@ -104,3 +106,14 @@ void fun() {
   (void) b<double>;
   (void) var_host_only;
 }
+
+extern __global__ void external_func();
+extern void* const external_dep[] = {
+  (void*)(external_func)
+};
+extern void* const external_arr[] = {};
+
+void* host_fun() {
+  (void) external_dep;
+  (void) external_arr;
+}