https://github.com/aobolensk created https://github.com/llvm/llvm-project/pull/230901
CUDA host shadows of device arrays, pointers and member pointers crashed the pass >From 5cb98cae275fb4b0dafbaffb3fc3524513f90cce Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Sun, 11 Oct 2026 07:53:48 +0200 Subject: [PATCH] [CIR] Lower undef and poison initializers of any type in CXXABILowering CUDA host shadows of device arrays, pointers and member pointers crashed the pass --- .../CIR/Dialect/Transforms/CXXABILowering.cpp | 14 +++++----- .../test/CIR/CodeGenCUDA/shadow-var-undef.cu | 27 +++++++++++++++++++ .../CIR/Transforms/cxx-abi-lowering-attrs.cir | 8 ++++++ 3 files changed, 41 insertions(+), 8 deletions(-) create mode 100644 clang/test/CIR/CodeGenCUDA/shadow-var-undef.cu 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> _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
