diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp index 16a2b443563c..4a6275e07fcc 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/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() || - global->getType()->isCUDADeviceBuiltinTextureType()) { + if (global->getType()->isCUDADeviceBuiltinTextureType()) { llvm_unreachable("NYI"); } return !langOpts.CUDAIsDevice || global->hasAttr() || + global->hasAttr() || global->hasAttr() || 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() || + (d->hasAttr() || d->hasAttr() || 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() && 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()) diff --git a/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp b/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp index ef31582f65be..32ceb0096d71 100644 --- a/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp +++ b/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: diff --git a/clang/test/CIR/CodeGen/CUDA/addrspace-lowering.cu b/clang/test/CIR/CodeGen/CUDA/addrspace-lowering.cu index d6c8e3590968..91f26fa29597 100644 --- a/clang/test/CIR/CodeGen/CUDA/addrspace-lowering.cu +++ b/clang/test/CIR/CodeGen/CUDA/addrspace-lowering.cu @@ -14,6 +14,6 @@ __device__ int b; // LLVM-DEVICE: @b = addrspace(1) {{.*}} -// __constant__ int c; +__constant__ int c; -// XFAIL-LLVM-DEVICE: @c = addrspace(4) {{.*}} +// LLVM-DEVICE: @c = addrspace(4) {{.*}} diff --git a/clang/test/CIR/CodeGen/CUDA/global-vars.cu b/clang/test/CIR/CodeGen/CUDA/global-vars.cu index f6e630a7e797..25de5f28dd91 100644 --- a/clang/test/CIR/CodeGen/CUDA/global-vars.cu +++ b/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