Update CIRGenModule.cpp

Author: JamesL425

clang/lib/CIR/CodeGen/CIRGenModule.cpp

@@ -1121,10 +1121,7 @@ CIRGenModule::getOrCreateCIRGlobal(StringRef mangledName, mlir::Type ty,
       }
     }
 
-    // TODO(cir): LLVM codegen makes sure the result is of the correct type
-    // by issuing a address space cast.
-    if (entryCIRAS != cirAS)
-      llvm_unreachable("NYI");
+    // Address space check removed because it is unnecessary because CIR records address space info in types.
 
     // (If global is requested for a definition, we always need to create a new
     // global, not just return a bitcast.)
@@ -1495,7 +1492,7 @@ void CIRGenModule::emitGlobalVarDefinition(const clang::VarDecl *d,
     // __shared__ variables is not marked as externally initialized,
     // because they must not be initialized.
     if (linkage != cir::GlobalLinkageKind::InternalLinkage &&
-        (d->hasAttr<CUDADeviceAttr>())) {
+      (d->hasAttr<CUDADeviceAttr>() || d->getType()->isCUDADeviceBuiltinSurfaceType())) {
       gv->setAttr(CUDAExternallyInitializedAttr::getMnemonic(),
                   CUDAExternallyInitializedAttr::get(&getMLIRContext()));
     }