[CIR][CUDA] support for device variables (#1394)

Author: JamesL425

clang/lib/CIR/CodeGen/CIRGenModule.cpp

@@ -510,6 +510,23 @@ const ABIInfo &CIRGenModule::getABIInfo() {
   return getTargetCIRGenInfo().getABIInfo();
 }
 
+bool CIRGenModule::shouldEmitCUDAGlobalVar(const VarDecl *global) const {
+  assert(langOpts.CUDA && "Should not be called by non-CUDA languages");
+  // We need to emit host-side 'shadows' for all global
+  // 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>();
+}
+
 void CIRGenModule::emitGlobal(GlobalDecl GD) {
   llvm::TimeTraceScope scope("build CIR Global", [&]() -> std::string {
     auto *ND = dyn_cast<NamedDecl>(GD.getDecl());
@@ -554,8 +571,10 @@ void CIRGenModule::emitGlobal(GlobalDecl GD) {
       }
     }
 
-    if (dyn_cast<VarDecl>(Global))
-      llvm_unreachable("NYI");
+    if (const auto *VD = dyn_cast<VarDecl>(Global)) {
+      if (!shouldEmitCUDAGlobalVar(VD))
+        return;
+    }
   }
 
   if (langOpts.OpenMP) {
@@ -599,7 +618,6 @@ void CIRGenModule::emitGlobal(GlobalDecl GD) {
       return;
     }
   } else {
-    assert(!langOpts.CUDA && "NYI");
     const auto *VD = cast<VarDecl>(Global);
     assert(VD->isFileVarDecl() && "Cannot emit local var decl as global.");
     if (VD->isThisDeclarationADefinition() != VarDecl::Definition &&
@@ -1149,8 +1167,11 @@ CIRGenModule::getOrCreateCIRGlobal(StringRef MangledName, mlir::Type Ty,
 
     // External HIP managed variables needed to be recorded for transformation
     // in both device and host compilations.
-    if (getLangOpts().CUDA)
-      assert(0 && "not implemented");
+    // External HIP managed variables needed to be recorded for transformation
+    // in both device and host compilations.
+    if (getLangOpts().CUDA && D && D->hasAttr<HIPManagedAttr>() &&
+        D->hasExternalStorage())
+      llvm_unreachable("NYI");
   }
 
   // TODO(cir): address space cast when needed for DAddrSpace.
@@ -1422,9 +1443,6 @@ void CIRGenModule::emitGlobalVarDefinition(const clang::VarDecl *D,
   // the device. [...]"
   // CUDA B.2.2 "The __constant__ qualifier, optionally used together with
   // __device__, declares a variable that: [...]
-  if (GV && getLangOpts().CUDA) {
-    assert(0 && "not implemented");
-  }
 
   // Set initializer and finalize emission
   CIRGenModule::setInitializer(GV, Init);
@@ -4012,9 +4030,17 @@ LangAS CIRGenModule::getGlobalVarAddressSpace(const VarDecl *D) {
     llvm_unreachable("NYI");
 
   if (langOpts.CUDA && langOpts.CUDAIsDevice) {
-    if (D && D->hasAttr<CUDASharedAttr>())
-      return LangAS::cuda_shared;
-    llvm_unreachable("NYI");
+    if (D) {
+      if (D->hasAttr<CUDAConstantAttr>())
+        return LangAS::cuda_constant;
+      if (D->hasAttr<CUDASharedAttr>())
+        return LangAS::cuda_shared;
+      if (D->hasAttr<CUDADeviceAttr>())
+        return LangAS::cuda_device;
+      if (D->getType().isConstQualified())
+        return LangAS::cuda_constant;
+    }
+    return LangAS::cuda_device;
   }
 
   if (langOpts.OpenMP)

clang/lib/CIR/CodeGen/CIRGenModule.h

@@ -271,6 +271,9 @@ class CIRGenModule : public CIRGenTypeCache {
 
   // Return whether structured convergence intrinsics should be generated for
   // this target.
+
+  bool shouldEmitCUDAGlobalVar(const VarDecl *global) const;
+
   bool shouldEmitConvergenceTokens() const {
     // TODO: this shuld probably become unconditional once the controlled
     // convergence becomes the norm.

clang/lib/CIR/Dialect/IR/CIRAttrs.cpp

@@ -649,6 +649,8 @@ AddressSpaceAttr::getValueFromLangAS(clang::LangAS langAS) {
     // Local means local among the work-group (OpenCL) or block (CUDA).
     // All threads inside the kernel can access local memory.
     return Kind::offload_local;
+  case LangAS::cuda_device:
+    return Kind::offload_global;
   case LangAS::opencl_constant:
     return Kind::offload_constant;
   case LangAS::opencl_private:
@@ -658,7 +660,6 @@ AddressSpaceAttr::getValueFromLangAS(clang::LangAS langAS) {
 
   case LangAS::opencl_global_device:
   case LangAS::opencl_global_host:
-  case LangAS::cuda_device:
   case LangAS::cuda_constant:
   case LangAS::sycl_global:
   case LangAS::sycl_global_device:

clang/test/CIR/CodeGen/CUDA/global-vars.cu

@@ -0,0 +1,11 @@
+#include "../Inputs/cuda.h"
+
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fclangir \
+// RUN:            -fcuda-is-device -emit-cir -target-sdk-version=12.3 \
+// RUN:            %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR-DEVICE --input-file=%t.cir %s
+
+
+__device__ int a;
+
+// CIR-DEVICE: cir.global external addrspace(offload_global) @a = #cir.int<0> : !s32i {alignment = 4 : i64} loc(#loc3)