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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
