[clang] [CIR] Lower undef and poison initializers of any type in CXXABILowering (PR #230901)

[via cfe-commits]

llvmorg-github-actions[bot] wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-clang

Author: Arseniy Obolenskiy (aobolensk)

<details>
<summary>Changes</summary>

CUDA host shadows of device arrays, pointers and member pointers crashed the pass

---
Full diff: https://github.com/llvm/llvm-project/pull/230901.diff


3 Files Affected:

- (modified) clang/lib/CIR/Dialect/Transforms/CXXABILowering.cpp (+6-8) 
- (added) clang/test/CIR/CodeGenCUDA/shadow-var-undef.cu (+27) 
- (modified) clang/test/CIR/Transforms/cxx-abi-lowering-attrs.cir (+8) 


``````````diff
diff --git a/clang/lib/CIR/Dialect/Transforms/CXXABILowering.cpp b/clang/lib/CIR/Dialect/Transforms/CXXABILowering.cpp
index fd4e1863de5e5..86c74ce14842b 100644
--- a/clang/lib/CIR/Dialect/Transforms/CXXABILowering.cpp
+++ b/clang/lib/CIR/Dialect/Transforms/CXXABILowering.cpp
@@ -409,6 +409,12 @@ static mlir::TypedAttr lowerInitialValue(const LowerModule *lowerModule,
                                          const mlir::TypeConverter &tc,
                                          mlir::Type ty,
                                          mlir::Attribute initVal) {
+  if (mlir::isa_and_present<cir::UndefAttr>(initVal))
+    return cir::UndefAttr::get(tc.convertType(ty));
+
+  if (mlir::isa_and_present<cir::PoisonAttr>(initVal))
+    return cir::PoisonAttr::get(tc.convertType(ty));
+
   if (mlir::isa<cir::DataMemberType>(ty)) {
     // Members without a CIR field index (e.g. no_unique_address empty fields)
     // are represented by an explicit byte offset instead of a field path.
@@ -469,14 +475,6 @@ static mlir::TypedAttr lowerInitialValue(const LowerModule *lowerModule,
     if (auto recVal = mlir::dyn_cast_if_present<cir::ZeroAttr>(initVal))
       return cir::ZeroAttr::get(convertedTy);
 
-    if (auto undefVal = mlir::dyn_cast_if_present<cir::UndefAttr>(initVal))
-      return cir::UndefAttr::get(convertedTy);
-
-    // This might not be possible from Clang directly, but we can get here with
-    // hand-written IR.
-    if (auto poisonVal = mlir::dyn_cast_if_present<cir::PoisonAttr>(initVal))
-      return cir::PoisonAttr::get(convertedTy);
-
     if (auto recVal =
             mlir::dyn_cast_if_present<cir::ConstRecordAttr>(initVal)) {
       auto recordMembers = mlir::cast<ArrayAttr>(recVal.getMembers());
diff --git a/clang/test/CIR/CodeGenCUDA/shadow-var-undef.cu b/clang/test/CIR/CodeGenCUDA/shadow-var-undef.cu
new file mode 100644
index 0000000000000..a6da66f66cb6b
--- /dev/null
+++ b/clang/test/CIR/CodeGenCUDA/shadow-var-undef.cu
@@ -0,0 +1,27 @@
+#include "Inputs/cuda.h"
+
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir \
+// RUN:            -x cuda -emit-llvm -target-sdk-version=12.3 \
+// RUN:            %s -o %t-cir.ll
+// RUN: FileCheck --check-prefix=HOST --input-file=%t-cir.ll %s
+
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu \
+// RUN:            -x cuda -emit-llvm -target-sdk-version=12.3 \
+// RUN:            %s -o %t.ll
+// RUN: FileCheck --check-prefix=HOST --input-file=%t.ll %s
+
+// Host shadows of device variables are undef for every type.
+
+struct S { float x, y; };
+
+// HOST: @arr = internal global [4 x %struct.S] undef
+__constant__ S arr[4];
+
+// HOST: @ptr = internal global ptr undef
+__device__ S *ptr;
+
+// HOST: @dm = internal global i64 undef
+__device__ int S::*dm;
+
+// HOST: @mf = internal global { i64, i64 } undef
+__device__ float (S::*mf)();
diff --git a/clang/test/CIR/Transforms/cxx-abi-lowering-attrs.cir b/clang/test/CIR/Transforms/cxx-abi-lowering-attrs.cir
index 06259eda595a1..d0576ac004160 100644
--- a/clang/test/CIR/Transforms/cxx-abi-lowering-attrs.cir
+++ b/clang/test/CIR/Transforms/cxx-abi-lowering-attrs.cir
@@ -25,10 +25,18 @@ module attributes {cir.triple = "x86_64-unknown-linux-gnu"} {
     // cir::PoisonAttr
     %poison1 = cir.const #cir.poison : !rec_Struct
     // CHECK: cir.const #cir.poison : !rec_Struct
+    %poison2 = cir.const #cir.poison : !cir.array<!dmp x 2>
+    // CHECK: cir.const #cir.poison : !cir.array<!s64i x 2>
 
     // cir::UndefAttr
     %undef1 = cir.const #cir.undef : !rec_Struct
     // CHECK: cir.const #cir.undef : !rec_Struct
+    %undef2 = cir.const #cir.undef : !dmp
+    // CHECK: cir.const #cir.undef : !s64i
+    %undef3 = cir.const #cir.undef : !dmp_ptr
+    // CHECK: cir.const #cir.undef : !cir.ptr<!s64i>
+    %undef4 = cir.const #cir.undef : !cir.array<!dmp x 2>
+    // CHECK: cir.const #cir.undef : !cir.array<!s64i x 2>
 
     // cir::ConstPtrAttr
     %constptr1 = cir.const #cir.ptr<null> : !cir.ptr<!rec_Struct>

``````````

</details>


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