[clang] [CIR][CUDA] Handle local, __device__, and __shared__ variables (PR #184248)

[Andy Kaylor via cfe-commits]

================
@@ -1563,6 +1575,39 @@ CIRGenModule::getAddrOfConstantStringFromLiteral(const StringLiteral *s,
   return builder.getGlobalViewAttr(ptrTy, gv);
 }
 
+LangAS CIRGenModule::getGlobalVarAddressSpace(const VarDecl *d) {
+  if (langOpts.OpenCL) {
+    LangAS as = d ? d->getType().getAddressSpace() : LangAS::opencl_global;
+    assert(as == LangAS::opencl_global || as == LangAS::opencl_global_device ||
+           as == LangAS::opencl_global_host || as == LangAS::opencl_constant ||
+           as == LangAS::opencl_local || as >= LangAS::FirstTargetAddressSpace);
+    return as;
+  }
+
+  if (langOpts.SYCLIsDevice &&
+      (!d || d->getType().getAddressSpace() == LangAS::Default))
+    llvm_unreachable("NYI");
----------------
andykaylor wrote:

```suggestion
    errorNYI(g->getSourceRange(), "global as for SYCL device");
```

https://github.com/llvm/llvm-project/pull/184248