constvar + address space stuff (llvm#2)

* emit const attr for cuda

getValueFromLangAS: implement cuda constant

fix shouldEmitCUDAGlobalVar
- copies og

emitGlobalVarDefinition: set constant for cudaconstant

* emitGlobal: move shouldEmitCUDAGlobalVar to first check
- matches og

* tests: add constant to global-vars.cu
- also fixed __device__ check

* implement const

* fix address space for offload_constant

* Fix address space values for NVPTX

* add testcase

* test: enable const test in addrspace-lowering.cu

* remove duplicated case

Author: anominos

clang/lib/CIR/CodeGen/CIRGenModule.cpp

@@ -569,12 +569,12 @@ bool CIRGenModule::shouldEmitCUDAGlobalVar(const VarDecl *global) const {
   // size and host-side address in order to provide access to
   // their device-side incarnations.
 
-  if (global->hasAttr<CUDAConstantAttr>() ||
-      global->getType()->isCUDADeviceBuiltinTextureType()) {
+  if (global->getType()->isCUDADeviceBuiltinTextureType()) {
     llvm_unreachable("NYI");
   }
 
   return !langOpts.CUDAIsDevice || global->hasAttr<CUDADeviceAttr>() ||
+         global->hasAttr<CUDAConstantAttr>() ||
          global->hasAttr<CUDASharedAttr>() ||
          global->getType()->isCUDADeviceBuiltinSurfaceType();
 }
@@ -1492,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->hasAttr<CUDAConstantAttr>() ||
          d->getType()->isCUDADeviceBuiltinSurfaceType())) {
       gv->setAttr(CUDAExternallyInitializedAttr::getMnemonic(),
                   CUDAExternallyInitializedAttr::get(&getMLIRContext()));
@@ -1505,8 +1505,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>())

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

@@ -652,13 +652,12 @@ 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;
   case LangAS::opencl_generic:
     return Kind::offload_generic;
-  case LangAS::cuda_constant:
-    return Kind::offload_constant;
   case LangAS::opencl_global_device:
   case LangAS::opencl_global_host:
   case LangAS::sycl_global:

clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/NVPTX.cpp

@@ -52,9 +52,9 @@ class NVPTXTargetLoweringInfo : public TargetLoweringInfo {
     case Kind::offload_global:
       return 1;
     case Kind::offload_constant:
-      return 2;
-    case Kind::offload_generic:
       return 4;
+    case Kind::offload_generic:
+      return 0;
     default:
       cir_cconv_unreachable("Unknown CIR address space for this target");
     }

clang/test/CIR/CodeGen/CUDA/addrspace-lowering.cu

@@ -0,0 +1,19 @@
+#include "../Inputs/cuda.h"
+
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fclangir \
+// RUN:            -fcuda-is-device -emit-llvm -target-sdk-version=12.3 \
+// RUN:            %s -o %t.ll
+// RUN: FileCheck --check-prefix=LLVM-DEVICE --input-file=%t.ll %s
+
+
+__shared__ int a;
+
+// LLVM-DEVICE: @a = addrspace(3) {{.*}}
+
+__device__ int b;
+
+// LLVM-DEVICE: @b = addrspace(1) {{.*}}
+
+__constant__ int c;
+
+// LLVM-DEVICE: @c = addrspace(4) {{.*}}

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

@@ -17,3 +17,7 @@ __device__ int a;
 __shared__ int shared;
 // CIR-DEVICE: cir.global external addrspace(offload_local) @shared = #cir.undef
 // LLVM-DEVICE: @shared = addrspace(3) global i32 undef, align 4
+
+__constant__ int b;
+// CIR-DEVICE: cir.global constant external addrspace(offload_constant) @b = #cir.int<0> : !s32i {alignment = 4 : i64, cu.externally_initialized = #cir.cu.externally_initialized}
+// LLVM-DEVICE: @b = addrspace(4) externally_initialized constant i32 0, align 4