diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
index 52166ba2efb1..e9afbdf9d544 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
@@ -517,15 +517,11 @@ bool CIRGenModule::shouldEmitCUDAGlobalVar(const VarDecl *global) const {
   // device-side variables because the CUDA runtime needs their
   // size and host-side address in order to provide access to
   // their device-side incarnations.
-
-  if (global->hasAttr<CUDAConstantAttr>() ||
-      global->hasAttr<CUDASharedAttr>() ||
-      global->getType()->isCUDADeviceBuiltinSurfaceType() ||
-      global->getType()->isCUDADeviceBuiltinTextureType()) {
-    llvm_unreachable("NYI");
-  }
-
-  return !langOpts.CUDAIsDevice || global->hasAttr<CUDADeviceAttr>();
+  return !langOpts.CUDAIsDevice || global->hasAttr<CUDADeviceAttr>() ||
+          global->hasAttr<CUDAConstantAttr>() ||
+          global->hasAttr<CUDASharedAttr>() ||
+          global->getType()->isCUDADeviceBuiltinSurfaceType() ||
+          global->getType()->isCUDADeviceBuiltinTextureType();
 }
 
 void CIRGenModule::emitGlobal(GlobalDecl gd) {
@@ -549,7 +545,10 @@ void CIRGenModule::emitGlobal(GlobalDecl gd) {
 
   if (langOpts.CUDA || langOpts.HIP) {
     // clang uses the same flag when building HIP code
-    if (langOpts.CUDAIsDevice) {
+    if (const auto *vd = dyn_cast<VarDecl>(global)) {
+      if (!shouldEmitCUDAGlobalVar(vd))
+        return;
+    } else if (langOpts.CUDAIsDevice) {
       // This will implicitly mark templates and their
       // specializations as __host__ __device__.
       if (langOpts.OffloadImplicitHostDeviceTemplates)
@@ -571,11 +570,6 @@ void CIRGenModule::emitGlobal(GlobalDecl gd) {
         return;
       }
     }
-
-    if (const auto *vd = dyn_cast<VarDecl>(global)) {
-      if (!shouldEmitCUDAGlobalVar(vd))
-        return;
-    }
   }
 
   if (langOpts.OpenMP) {
@@ -1452,8 +1446,9 @@ void CIRGenModule::emitGlobalVarDefinition(const clang::VarDecl *d,
     emitter->finalize(gv);
 
   // TODO(cir): If it is safe to mark the global 'constant', do so now.
-  gv.setConstant(!needsGlobalCtor && !needsGlobalDtor &&
-                 isTypeConstant(d->getType(), true, true));
+  gv.setConstant((d->hasAttr<CUDAConstantAttr>() && langOpts.CUDAIsDevice) ||
+                 (!needsGlobalCtor && !needsGlobalDtor &&
+                 isTypeConstant(d->getType(), true, true)));
 
   // If it is in a read-only section, mark it 'constant'.
   if (const SectionAttr *sa = d->getAttr<SectionAttr>())
diff --git a/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp b/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp
index eaa5ffd151b0..96acf8d84ae8 100644
--- a/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp
+++ b/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp
@@ -652,6 +652,7 @@ AddressSpaceAttr::getValueFromLangAS(clang::LangAS langAS) {
   case LangAS::cuda_device:
     return Kind::offload_global;
   case LangAS::opencl_constant:
+  case LangAS::cuda_constant:
     return Kind::offload_constant;
   case LangAS::opencl_private:
     return Kind::offload_private;
@@ -660,7 +661,6 @@ AddressSpaceAttr::getValueFromLangAS(clang::LangAS langAS) {
 
   case LangAS::opencl_global_device:
   case LangAS::opencl_global_host:
-  case LangAS::cuda_constant:
   case LangAS::sycl_global:
   case LangAS::sycl_global_device:
   case LangAS::sycl_global_host:
diff --git a/clang/test/CIR/CodeGen/CUDA/global-vars.cu b/clang/test/CIR/CodeGen/CUDA/global-vars.cu
index 5b1374c085eb..f8babab01c45 100644
--- a/clang/test/CIR/CodeGen/CUDA/global-vars.cu
+++ b/clang/test/CIR/CodeGen/CUDA/global-vars.cu
@@ -8,4 +8,8 @@
 
 __device__ int a;
 
-// CIR-DEVICE: cir.global external addrspace(offload_global) @a = #cir.int<0> : !s32i {alignment = 4 : i64} loc(#loc3)
\ No newline at end of file
+// CIR-DEVICE: cir.global external addrspace(offload_global) @a ={{.*}}
+
+__constant__ int b;
+
+// CIR-DEVICE: cir.global constant external addrspace(offload_constant) @b ={{.*}}