[llvm-branch-commits] [clang] [CIR] Global w/ Poison Attr lowering and CUDA `__shared__` global lowering (PR #186564)

[David Rivera via llvm-branch-commits]

https://github.com/RiverDave created https://github.com/llvm/llvm-project/pull/186564

None

>From db5ba3a066398a212901703c2c233cd62f21fbd8 Mon Sep 17 00:00:00 2001
From: David Rivera <davidriverg at gmail.com>
Date: Sat, 14 Mar 2026 02:08:52 -0400
Subject: [PATCH] [CIR] Global w/ Poison Attr lowering and CUDA `__shared__`
 global lowering

---
 .../CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp   | 17 ++++++++++++-----
 clang/test/CIR/CodeGenCUDA/address-spaces.cu    | 14 ++++++++++++++
 2 files changed, 26 insertions(+), 5 deletions(-)

diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
index 00b1070576cfb..b380429b5d220 100644
--- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
+++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
@@ -365,7 +365,7 @@ class CIRAttrToValue {
         .Case<cir::IntAttr, cir::FPAttr, cir::ConstComplexAttr,
               cir::ConstArrayAttr, cir::ConstRecordAttr, cir::ConstVectorAttr,
               cir::ConstPtrAttr, cir::GlobalViewAttr, cir::TypeInfoAttr,
-              cir::UndefAttr, cir::VTableAttr, cir::ZeroAttr>(
+              cir::PoisonAttr, cir::UndefAttr, cir::VTableAttr, cir::ZeroAttr>(
             [&](auto attrT) { return visitCirAttr(attrT); })
         .Default([&](auto attrT) { return mlir::Value(); });
   }
@@ -379,6 +379,7 @@ class CIRAttrToValue {
   mlir::Value visitCirAttr(cir::ConstVectorAttr attr);
   mlir::Value visitCirAttr(cir::GlobalViewAttr attr);
   mlir::Value visitCirAttr(cir::TypeInfoAttr attr);
+  mlir::Value visitCirAttr(cir::PoisonAttr attr);
   mlir::Value visitCirAttr(cir::UndefAttr attr);
   mlir::Value visitCirAttr(cir::VTableAttr attr);
   mlir::Value visitCirAttr(cir::ZeroAttr attr);
@@ -725,6 +726,12 @@ mlir::Value CIRAttrToValue::visitCirAttr(cir::TypeInfoAttr typeInfoAttr) {
 }
 
 /// UndefAttr visitor.
+mlir::Value CIRAttrToValue::visitCirAttr(cir::PoisonAttr poisonAttr) {
+  mlir::Location loc = parentOp->getLoc();
+  return mlir::LLVM::PoisonOp::create(
+      rewriter, loc, converter->convertType(poisonAttr.getType()));
+}
+
 mlir::Value CIRAttrToValue::visitCirAttr(cir::UndefAttr undefAttr) {
   mlir::Location loc = parentOp->getLoc();
   return mlir::LLVM::UndefOp::create(
@@ -2516,8 +2523,8 @@ CIRToLLVMGlobalOpLowering::matchAndRewriteRegionInitializedGlobal(
   assert(
       (isa<cir::ConstArrayAttr, cir::ConstRecordAttr, cir::ConstVectorAttr,
            cir::ConstPtrAttr, cir::ConstComplexAttr, cir::GlobalViewAttr,
-           cir::TypeInfoAttr, cir::UndefAttr, cir::VTableAttr, cir::ZeroAttr>(
-          init)));
+           cir::TypeInfoAttr, cir::PoisonAttr, cir::UndefAttr, cir::VTableAttr,
+           cir::ZeroAttr>(init)));
 
   // TODO(cir): once LLVM's dialect has proper equivalent attributes this
   // should be updated. For now, we use a custom op to initialize globals
@@ -2577,8 +2584,8 @@ mlir::LogicalResult CIRToLLVMGlobalOpLowering::matchAndRewrite(
     } else if (mlir::isa<cir::ConstArrayAttr, cir::ConstVectorAttr,
                          cir::ConstRecordAttr, cir::ConstPtrAttr,
                          cir::ConstComplexAttr, cir::GlobalViewAttr,
-                         cir::TypeInfoAttr, cir::UndefAttr, cir::VTableAttr,
-                         cir::ZeroAttr>(init.value())) {
+                         cir::TypeInfoAttr, cir::PoisonAttr, cir::UndefAttr,
+                         cir::VTableAttr, cir::ZeroAttr>(init.value())) {
       // TODO(cir): once LLVM's dialect has proper equivalent attributes this
       // should be updated. For now, we use a custom op to initialize globals
       // to the appropriate value.
diff --git a/clang/test/CIR/CodeGenCUDA/address-spaces.cu b/clang/test/CIR/CodeGenCUDA/address-spaces.cu
index 54e85ab75bd76..49dacf1200a2d 100644
--- a/clang/test/CIR/CodeGenCUDA/address-spaces.cu
+++ b/clang/test/CIR/CodeGenCUDA/address-spaces.cu
@@ -7,6 +7,10 @@
 // RUN:   -fcuda-is-device -fclangir -emit-cir %s -o %t.cir
 // RUN: FileCheck --check-prefix=CIR-POST --input-file=%t.cir %s
 
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -x cuda \
+// RUN:   -fcuda-is-device -fclangir -emit-llvm %s -o %t-cir.ll
+// RUN: FileCheck --check-prefix=CIR-LLVM --input-file=%t-cir.ll %s
+
 // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -x cuda \
 // RUN:   -fcuda-is-device -emit-llvm %s -o %t.ll
 // RUN: FileCheck --check-prefix=OGCG --input-file=%t.ll %s
@@ -17,34 +21,44 @@
 
 // CIR-PRE: cir.global external  lang_address_space(offload_global) @i = #cir.int<0> : !s32i
 // CIR-POST: cir.global external  target_address_space(1) @i = #cir.int<0> : !s32i
+// CIR-LLVM-DAG: @i = addrspace(1) global i32 0, align 4
 // OGCG-DAG: @i = addrspace(1) externally_initialized global i32 0, align 4
 __device__ int i;
 
 // CIR-PRE: cir.global constant external  lang_address_space(offload_constant) @j = #cir.int<0> : !s32i
 // CIR-POST: cir.global constant external  target_address_space(4) @j = #cir.int<0> : !s32i
+// CIR-LLVM-DAG: @j = addrspace(4) constant i32 0, align 4
 // OGCG-DAG: @j = addrspace(4) externally_initialized constant i32 0, align 4
 __constant__ int j;
 
 // CIR-PRE: cir.global external  lang_address_space(offload_local) @k = #cir.poison : !s32i
 // CIR-POST: cir.global external  target_address_space(3) @k = #cir.poison : !s32i
+// CIR-LLVM-DAG: @k = addrspace(3) global i32 poison, align 4
 // OGCG-DAG: @k = addrspace(3) global i32 undef, align 4
 __shared__ int k;
 
 // CIR-PRE: cir.global external  lang_address_space(offload_local) @b = #cir.poison : !cir.float
 // CIR-POST: cir.global external  target_address_space(3) @b = #cir.poison : !cir.float
+// CIR-LLVM-DAG: @b = addrspace(3) global float poison, align 4
 // OGCG-DAG: @b = addrspace(3) global float undef, align 4
 __shared__ float b;
 
 __device__ void foo() {
   // CIR-PRE: cir.get_global @i : !cir.ptr<!s32i, lang_address_space(offload_global)>
   // CIR-POST: cir.get_global @i : !cir.ptr<!s32i, target_address_space(1)>
+  // CIR-LLVM: load i32, ptr addrspace(1) @i
+  // OGCG: load i32, ptr addrspacecast (ptr addrspace(1) @i to ptr)
   i++;
 
   // CIR-PRE: cir.get_global @j : !cir.ptr<!s32i, lang_address_space(offload_constant)>
   // CIR-POST: cir.get_global @j : !cir.ptr<!s32i, target_address_space(4)>
+  // CIR-LLVM: load i32, ptr addrspace(4) @j
+  // OGCG: load i32, ptr addrspacecast (ptr addrspace(4) @j to ptr)
   j++;
 
   // CIR-PRE: cir.get_global @k : !cir.ptr<!s32i, lang_address_space(offload_local)>
   // CIR-POST: cir.get_global @k : !cir.ptr<!s32i, target_address_space(3)>
+  // CIR-LLVM: load i32, ptr addrspace(3) @k
+  // OGCG: load i32, ptr addrspacecast (ptr addrspace(3) @k to ptr)
   k++;
 }