https://github.com/koparasy updated https://github.com/llvm/llvm-project/pull/231689
>From d211acdc555e12453f1d85b0ff92d54b27a41164 Mon Sep 17 00:00:00 2001 From: Konstantinos Parasyris <[email protected]> Date: Tue, 13 Oct 2026 14:03:46 -0700 Subject: [PATCH 1/3] [CIR] Emit references to constant-initialized const variables as constants Co-Authored-By: Claude Opus 5.5 (1M context) <[email protected]> --- clang/include/clang/CIR/Dialect/IR/CIROps.td | 2 + clang/include/clang/CIR/MissingFeatures.h | 1 - clang/lib/CIR/CodeGen/CIRGenExpr.cpp | 139 ++++++-- clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp | 36 +- clang/lib/CIR/CodeGen/CIRGenFunction.h | 9 +- .../CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 28 +- clang/test/CIR/CodeGen/assume-attr.cpp | 17 +- clang/test/CIR/CodeGen/basic.cpp | 2 +- .../CodeGen/const-var-emit-as-constant.cpp | 319 ++++++++++++++++++ clang/test/CIR/CodeGen/global-init-fenv.cpp | 16 +- .../CIR/CodeGenHIP/lambda-reference-var.hip | 161 +++++++++ .../CIR/CodeGenSYCL/lambda-capture-const.cpp | 33 ++ clang/test/CIR/Lowering/basic.cpp | 3 +- 13 files changed, 686 insertions(+), 80 deletions(-) create mode 100644 clang/test/CIR/CodeGen/const-var-emit-as-constant.cpp create mode 100644 clang/test/CIR/CodeGenHIP/lambda-reference-var.hip create mode 100644 clang/test/CIR/CodeGenSYCL/lambda-capture-const.cpp diff --git a/clang/include/clang/CIR/Dialect/IR/CIROps.td b/clang/include/clang/CIR/Dialect/IR/CIROps.td index c631d23b6e7f9a..1f2dad7d9fdb42 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIROps.td +++ b/clang/include/clang/CIR/Dialect/IR/CIROps.td @@ -680,6 +680,8 @@ def CIR_ConstantOp : CIR_Op<"const", [ let hasCXXABILowering = true; let isLLVMLoweringRecursive = true; + let customLLVMLoweringConstructorDecl = + LoweringBuilders<(ins "LLVMBlockAddressInfo &":$blockInfoAddr)>; } //===----------------------------------------------------------------------===// diff --git a/clang/include/clang/CIR/MissingFeatures.h b/clang/include/clang/CIR/MissingFeatures.h index cd1e43b17ec3f0..deb469f8fe66ea 100644 --- a/clang/include/clang/CIR/MissingFeatures.h +++ b/clang/include/clang/CIR/MissingFeatures.h @@ -304,7 +304,6 @@ struct MissingFeatures { static bool targetCIRGenInfoOS() { return false; } static bool targetCodeGenInfoGetNullPointer() { return false; } static bool thunks() { return false; } - static bool tryEmitAsConstant() { return false; } static bool typeAwareAllocation() { return false; } static bool typeChecks() { return false; } static bool useEHCleanupForArray() { return false; } diff --git a/clang/lib/CIR/CodeGen/CIRGenExpr.cpp b/clang/lib/CIR/CodeGen/CIRGenExpr.cpp index 9c71d987c37e16..b15f80bc461e6e 100644 --- a/clang/lib/CIR/CodeGen/CIRGenExpr.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenExpr.cpp @@ -19,6 +19,7 @@ #include "mlir/Dialect/Ptr/IR/MemorySpaceInterfaces.h" #include "mlir/IR/BuiltinAttributes.h" #include "mlir/IR/Value.h" +#include "clang/AST/ASTLambda.h" #include "clang/AST/Attr.h" #include "clang/AST/CharUnits.h" #include "clang/AST/Decl.h" @@ -3055,36 +3056,122 @@ Address CIRGenFunction::createDefaultAlignTempAlloca(mlir::Type ty, return createTempAlloca(ty, align, loc, name); } +/// Given an object of the given canonical type, can we safely copy a +/// value out of it based on its initializer? +static bool isConstantEmittableObjectType(QualType type) { + assert(type.isCanonical()); + assert(!type->isReferenceType()); + + // Must be const-qualified but non-volatile. + Qualifiers qs = type.getLocalQualifiers(); + if (!qs.hasConst() || qs.hasVolatile()) + return false; + + // Otherwise, all object types satisfy this except C++ classes with + // mutable subobjects or non-trivial copy/destroy behavior. + if (const auto *rt = dyn_cast<clang::RecordType>(type)) + if (const auto *rd = dyn_cast<CXXRecordDecl>(rt->getDecl())) { + rd = rd->getDefinitionOrSelf(); + if (rd->hasMutableFields() || !rd->isTrivial()) + return false; + } + + return true; +} + +/// Can we constant-emit a load of a reference to a variable of the +/// given type? This is different from predicates like +/// Decl::mightBeUsableInConstantExpressions because we do want it to apply +/// in situations that don't necessarily satisfy the language's rules +/// for this (e.g. C++'s ODR-use rules). For example, we want to able +/// to do this with const float variables even if those variables +/// aren't marked 'constexpr'. +enum ConstantEmissionKind { + CEK_None, + CEK_AsReferenceOnly, + CEK_AsValueOrReference, + CEK_AsValueOnly +}; +static ConstantEmissionKind checkVarTypeForConstantEmission(QualType type) { + type = type.getCanonicalType(); + if (const auto *ref = dyn_cast<ReferenceType>(type)) { + if (isConstantEmittableObjectType(ref->getPointeeType())) + return CEK_AsValueOrReference; + return CEK_AsReferenceOnly; + } + if (isConstantEmittableObjectType(type)) + return CEK_AsValueOnly; + return CEK_None; +} + /// Try to emit a reference to the given value without producing it as -/// an l-value. For many cases, this is just an optimization, but it avoids -/// us needing to emit global copies of variables if they're named without -/// triggering a formal use in a context where we can't emit a direct -/// reference to them, for instance if a block or lambda or a member of a -/// local class uses a const int variable or constexpr variable from an -/// enclosing function. -/// -/// For named members of enums, this is the only way they are emitted. +/// an l-value. This is just an optimization, but it avoids us needing +/// to emit global copies of variables if they're named without triggering +/// a formal use in a context where we can't emit a direct reference to them, +/// for instance if a block or lambda or a member of a local class uses a +/// const int variable or constexpr variable from an enclosing function. CIRGenFunction::ConstantEmission CIRGenFunction::tryEmitAsConstant(const DeclRefExpr *refExpr) { const ValueDecl *value = refExpr->getDecl(); - // There is a lot more to do here, but for now only EnumConstantDecl is - // supported. - assert(!cir::MissingFeatures::tryEmitAsConstant()); - // The value needs to be an enum constant or a constant variable. - if (!isa<EnumConstantDecl>(value)) + ConstantEmissionKind cek; + if (isa<ParmVarDecl>(value)) + cek = CEK_None; + else if (const auto *var = dyn_cast<VarDecl>(value)) + cek = checkVarTypeForConstantEmission(var->getType()); + else if (isa<EnumConstantDecl>(value)) + cek = CEK_AsValueOnly; + else + cek = CEK_None; + if (cek == CEK_None) return ConstantEmission(); Expr::EvalResult result; - if (!refExpr->EvaluateAsRValue(result, getContext())) + bool resultIsReference; + QualType resultType; + + // It's best to evaluate all the way as an r-value if that's permitted. + if (cek != CEK_AsReferenceOnly && + refExpr->EvaluateAsRValue(result, getContext())) { + resultIsReference = false; + resultType = refExpr->getType().getUnqualifiedType(); + + // Otherwise, try to evaluate as an l-value. + } else if (cek != CEK_AsValueOnly && + refExpr->EvaluateAsLValue(result, getContext())) { + resultIsReference = true; + resultType = value->getType(); + + // Failure. + } else { return ConstantEmission(); + } - QualType resultType = refExpr->getType(); + // In any case, if the initializer has side-effects, abandon ship. + if (result.HasSideEffects) + return ConstantEmission(); - // As long as we're only handling EnumConstantDecl, there should be no - // side-effects. - assert(!result.HasSideEffects); + // In CUDA/HIP device compilation, a lambda may capture a reference variable + // referencing a global host variable by copy. In this case the lambda should + // make a copy of the value of the global host variable. The DRE of the + // captured reference variable cannot be emitted as load from the host + // global variable as compile time constant, since the host variable is not + // accessible on device. The DRE of the captured reference variable has to be + // loaded from captures. + if (cgm.getLangOpts().CUDAIsDevice && result.Val.isLValue() && + refExpr->refersToEnclosingVariableOrCapture()) { + auto *md = dyn_cast_or_null<CXXMethodDecl>(curCodeDecl); + if (isLambdaMethod(md) && md->getOverloadedOperator() == OO_Call) { + const APValue::LValueBase &base = result.Val.getLValueBase(); + if (const ValueDecl *d = base.dyn_cast<const ValueDecl *>()) { + if (const auto *vd = dyn_cast<const VarDecl>(d)) { + if (!vd->hasAttr<CUDADeviceAttr>()) + return ConstantEmission(); + } + } + } + } // Emit as a constant. // FIXME(cir): have emitAbstract build a TypedAttr instead (this requires @@ -3092,10 +3179,16 @@ CIRGenFunction::tryEmitAsConstant(const DeclRefExpr *refExpr) { mlir::Attribute c = ConstantEmitter(*this).emitAbstract( refExpr->getLocation(), result.Val, resultType); mlir::TypedAttr cstToEmit = mlir::dyn_cast_if_present<mlir::TypedAttr>(c); - assert(cstToEmit && "expected a typed attribute"); + if (!cstToEmit) + return ConstantEmission(); + // Make sure we emit a debug reference to the global variable. assert(!cir::MissingFeatures::generateDebugInfo()); + // If we emitted a reference constant, we need to dereference that. + if (resultIsReference) + return ConstantEmission::forReference(cstToEmit); + return ConstantEmission::forValue(cstToEmit); } @@ -3109,10 +3202,10 @@ CIRGenFunction::tryEmitAsConstant(const MemberExpr *me) { mlir::Value CIRGenFunction::emitScalarConstant( const CIRGenFunction::ConstantEmission &constant, Expr *e) { assert(constant && "not a constant"); - if (constant.isReference()) { - cgm.errorNYI(e->getSourceRange(), "emitScalarConstant: reference"); - return {}; - } + if (constant.isReference()) + return emitLoadOfLValue(constant.getReferenceLValue(*this, e), + e->getExprLoc()) + .getValue(); return builder.getConstant(getLoc(e->getSourceRange()), constant.getValue()); } diff --git a/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp b/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp index 8e8eea7c21ab55..7e95274fc649fe 100644 --- a/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp @@ -2949,25 +2949,27 @@ mlir::Value ScalarExprEmitter::VisitCallExpr(const CallExpr *e) { } mlir::Value ScalarExprEmitter::VisitMemberExpr(MemberExpr *e) { - // TODO(cir): The classic codegen calls tryEmitAsConstant() here. Folding - // constants sound like work for MLIR optimizers, but we'll keep an assertion - // for now. - assert(!cir::MissingFeatures::tryEmitAsConstant()); - Expr::EvalResult result; - if (e->EvaluateAsInt(result, cgf.getContext(), Expr::SE_AllowSideEffects)) { - llvm::APSInt value = result.Val.getInt(); + if (CIRGenFunction::ConstantEmission constant = cgf.tryEmitAsConstant(e)) { cgf.emitIgnoredExpr(e->getBase()); - mlir::Location loc = cgf.getLoc(e->getExprLoc()); - // The constant is folded from an APSInt with the source-type's bit width - // (1 for bool), but the AST's expression type is what later consumers of - // this value see. For a bool member we have to emit a !cir.bool constant - // -- otherwise downstream ops (cir.call into a bool parameter, cir.if / - // cir.ternary on the value, ...) would all reject the !cir.int<u, 1> the - // raw APSInt would produce. - if (e->getType()->isBooleanType()) - return builder.getBool(value.getBoolValue(), loc); - return builder.getConstInt(loc, value); + return cgf.emitScalarConstant(constant, e); + } else { + Expr::EvalResult result; + if (e->EvaluateAsInt(result, cgf.getContext(), Expr::SE_AllowSideEffects)) { + llvm::APSInt value = result.Val.getInt(); + cgf.emitIgnoredExpr(e->getBase()); + mlir::Location loc = cgf.getLoc(e->getExprLoc()); + // The constant is folded from an APSInt with the source-type's bit width + // (1 for bool), but the AST's expression type is what later consumers of + // this value see. For a bool member we have to emit a !cir.bool constant + // -- otherwise downstream ops (cir.call into a bool parameter, cir.if / + // cir.ternary on the value, ...) would all reject the !cir.int<u, 1> the + // raw APSInt would produce. + if (e->getType()->isBooleanType()) + return builder.getBool(value.getBoolValue(), loc); + return builder.getConstInt(loc, value); + } } + return emitLoadOfLValue(e); } diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.h b/clang/lib/CIR/CodeGen/CIRGenFunction.h index 8937e237a64141..5e0bd5c57c59be 100644 --- a/clang/lib/CIR/CodeGen/CIRGenFunction.h +++ b/clang/lib/CIR/CodeGen/CIRGenFunction.h @@ -800,11 +800,12 @@ class CIRGenFunction : public CIRGenTypeCache { } bool isReference() const { return valueAndIsReference.getInt(); } - LValue getReferenceLValue(CIRGenFunction &cgf, Expr *refExpr) const { + LValue getReferenceLValue(CIRGenFunction &cgf, const Expr *refExpr) const { assert(isReference()); - cgf.cgm.errorNYI(refExpr->getSourceRange(), - "ConstantEmission::getReferenceLValue"); - return {}; + mlir::Value addr = cgf.getBuilder().getConstant( + cgf.getLoc(refExpr->getSourceRange()), + mlir::cast<mlir::TypedAttr>(valueAndIsReference.getPointer())); + return cgf.makeNaturalAlignAddrLValue(addr, refExpr->getType()); } mlir::TypedAttr getValue() const { diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp index b724236ed16a34..33dea5c6f7655c 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp @@ -2662,10 +2662,12 @@ mlir::LogicalResult CIRToLLVMConstantOpLowering::matchAndRewrite( attr = rewriter.getIntegerAttr(typeConverter->convertType(op.getType()), value); } else if (mlir::isa<cir::IntType>(op.getType())) { - // Lower GlobalViewAttr to llvm.mlir.addressof + llvm.mlir.ptrtoint - if (auto gv = mlir::dyn_cast<cir::GlobalViewAttr>(op.getValue())) { - auto newOp = lowerCirAttrAsValue(op, gv, rewriter, symbolTables, - getTypeConverter()); + // Lower GlobalViewAttr to llvm.mlir.addressof + llvm.mlir.ptrtoint and + // BlockAddrDiffAttr to the difference of two llvm.blockaddress. + if (mlir::isa<cir::GlobalViewAttr, cir::BlockAddrDiffAttr>(op.getValue())) { + auto newOp = + lowerCirAttrAsValue(op, op.getValue(), rewriter, symbolTables, + getTypeConverter(), &blockInfoAddr); rewriter.replaceOp(op, newOp); return mlir::success(); } @@ -2686,10 +2688,14 @@ mlir::LogicalResult CIRToLLVMConstantOpLowering::matchAndRewrite( return mlir::success(); } } - // Lower GlobalViewAttr and GlobalOffsetAttr to llvm.mlir.addressof - if (mlir::isa<cir::GlobalViewAttr, cir::GlobalOffsetAttr>(op.getValue())) { - auto newOp = lowerCirAttrAsValue(op, op.getValue(), rewriter, - symbolTables, getTypeConverter()); + // Lower GlobalViewAttr and GlobalOffsetAttr to llvm.mlir.addressof, + // non-null ConstPtrAttr to llvm.inttoptr and BlockAddrInfoAttr to + // llvm.blockaddress. + if (mlir::isa<cir::GlobalViewAttr, cir::GlobalOffsetAttr, cir::ConstPtrAttr, + cir::BlockAddrInfoAttr>(op.getValue())) { + auto newOp = + lowerCirAttrAsValue(op, op.getValue(), rewriter, symbolTables, + getTypeConverter(), &blockInfoAddr); rewriter.replaceOp(op, newOp); return mlir::success(); } @@ -4636,9 +4642,9 @@ void ConvertCIRToLLVMPass::runOnOperation() { /// repeated O(M) module-wide symbol scans for every call site. mlir::SymbolTableCollection symbolTables; mlir::RewritePatternSet patterns(&getContext()); - patterns.add<CIRToLLVMBlockAddressOpLowering, CIRToLLVMGlobalOpLowering, - CIRToLLVMLabelOpLowering>(converter, patterns.getContext(), dl, - symbolTables, blockInfoAddr); + patterns.add<CIRToLLVMBlockAddressOpLowering, CIRToLLVMConstantOpLowering, + CIRToLLVMGlobalOpLowering, CIRToLLVMLabelOpLowering>( + converter, patterns.getContext(), dl, symbolTables, blockInfoAddr); patterns.add< #define GET_LLVM_LOWERING_PATTERNS_LIST diff --git a/clang/test/CIR/CodeGen/assume-attr.cpp b/clang/test/CIR/CodeGen/assume-attr.cpp index a3ea3bb6faaf66..ee4bdec795a7fa 100644 --- a/clang/test/CIR/CodeGen/assume-attr.cpp +++ b/clang/test/CIR/CodeGen/assume-attr.cpp @@ -27,15 +27,12 @@ void usage(int i, int j) { [[assume(globalVar == i)]] ; - // CIR: %[[GET_GLOB_VAR:.*]] = cir.get_global @_ZL9globalVar - // CIR: %[[LOAD_GLOB_VAR:.*]] = cir.load {{.*}}%[[GET_GLOB_VAR]] + // CIR: %[[GLOB_VAR:.*]] = cir.const #cir.int<5> : !s32i // CIR: %[[LOAD_I:.*]] = cir.load {{.*}}%[[I]] - // CIR: %[[CMP:.*]] = cir.cmp eq %[[LOAD_GLOB_VAR]], %[[LOAD_I]] + // CIR: %[[CMP:.*]] = cir.cmp eq %[[GLOB_VAR]], %[[LOAD_I]] // CIR: cir.assume %[[CMP]] : !cir.bool - // LLVM: %[[LOAD_GLOB_VAR:.*]] = load i32, ptr @_ZL9globalVar // BOTH: %[[LOAD_I:.*]] = load i32, ptr %[[I]] - // LLVM: %[[CMP:.*]] = icmp eq i32 %[[LOAD_GLOB_VAR]], %[[LOAD_I]] - // OGCG: %[[CMP:.*]] = icmp eq i32 5, %[[LOAD_I]] + // BOTH: %[[CMP:.*]] = icmp eq i32 5, %[[LOAD_I]] // BOTH: call void @llvm.assume(i1 %[[CMP]]) assume1(); // CIR: cir.call @_Z7assume1v() @@ -49,8 +46,8 @@ void usage(int i, int j) { // CIR: %[[J_EQ_I:.*]] = cir.cmp eq %[[LOAD_J]], %[[LOAD_I]] // CIR: %[[TERN:.*]] = cir.ternary(%[[J_EQ_I]], true { // CIR-NEXT: %[[LOAD_J:.*]] = cir.load {{.*}}%[[J]] - // CIR-NEXT: %[[LOAD_LOCAL:.*]] = cir.load {{.*}}%[[LOCAL]] - // CIR-NEXT: %[[J_GT_LOCAL:.*]] = cir.cmp gt %[[LOAD_J]], %[[LOAD_LOCAL]] + // CIR-NEXT: %[[LOCAL_VAL:.*]] = cir.const #cir.int<12> : !s32i + // CIR-NEXT: %[[J_GT_LOCAL:.*]] = cir.cmp gt %[[LOAD_J]], %[[LOCAL_VAL]] // CIR-NEXT: cir.yield %[[J_GT_LOCAL]] // CIR-NEXT: }, false { // CIR-NEXT: %[[FALSE:.*]] = cir.const #false @@ -64,9 +61,7 @@ void usage(int i, int j) { // BOTH: br i1 %[[J_EQ_I]], label %[[TRUE:.*]], label %[[FALSE:.*]] // BOTH: [[TRUE]]: // BOTH: %[[LOAD_J:.*]] = load i32, ptr %[[J]] - // LLVM: %[[LOAD_LOCAL:.*]] = load i32, ptr %[[LOCAL]] - // LLVM: %[[J_GT_LOCAL:.*]] = icmp sgt i32 %[[LOAD_J]], %[[LOAD_LOCAL]] - // OGCG: %[[J_GT_LOCAL:.*]] = icmp sgt i32 %[[LOAD_J]], 12 + // BOTH: %[[J_GT_LOCAL:.*]] = icmp sgt i32 %[[LOAD_J]], 12 // LLVM: br label %[[DONE:.*]] // OGCG: br label %[[FALSE]] // LLVM: [[FALSE]]: diff --git a/clang/test/CIR/CodeGen/basic.cpp b/clang/test/CIR/CodeGen/basic.cpp index 6a31e07a382859..d91cb4ee8081c5 100644 --- a/clang/test/CIR/CodeGen/basic.cpp +++ b/clang/test/CIR/CodeGen/basic.cpp @@ -49,7 +49,7 @@ int f2() { // CHECK: %[[I_PTR:.*]] = cir.alloca "i" align(4) init const : !cir.ptr<!s32i> // CHECK: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i // CHECK: cir.store{{.*}} %[[TWO]], %[[I_PTR]] : !s32i, !cir.ptr<!s32i> -// CHECK: %[[I:.*]] = cir.load{{.*}} %[[I_PTR]] : !cir.ptr<!s32i>, !s32i +// CHECK: %[[I:.*]] = cir.const #cir.int<2> : !s32i // CHECK: cir.store{{.*}} %[[I]], %[[RV]] : !s32i, !cir.ptr<!s32i> // CHECK: %[[R:.*]] = cir.load{{.*}} %[[RV]] : !cir.ptr<!s32i>, !s32i // CHECK: cir.return %[[R]] : !s32i diff --git a/clang/test/CIR/CodeGen/const-var-emit-as-constant.cpp b/clang/test/CIR/CodeGen/const-var-emit-as-constant.cpp new file mode 100644 index 00000000000000..a02c4b3bec4ce0 --- /dev/null +++ b/clang/test/CIR/CodeGen/const-var-emit-as-constant.cpp @@ -0,0 +1,319 @@ +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir -emit-llvm %s -o %t-cir.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t-cir.ll %s +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=OGCG --input-file=%t.ll %s + +// References to const, non-volatile variables whose initializer can be +// constant evaluated are emitted as the constant value, even when the variable +// is not usable in constant expressions (e.g. 'const float') and even when the +// reference is to an enclosing variable captured by a lambda. + +// CIR-NOT: cir.global {{.*}}@__const. +// LLVM-NOT: @__const. +// OGCG-NOT: @__const. + +float use(float); +int side(); +template <class F> auto call(F f) { return f(); } + +float captured() { + const float k = 0.9f; + const int n = 4; + return call([=] { return use(k) + n; }); +} + +int non_captured() { + const int n = 4; + return n; +} + +// CIR-LABEL: cir.func {{.*}} @_Z12non_capturedv() +// CIR: %[[RV:.*]] = cir.alloca "__retval" +// CIR: %[[N:.*]] = cir.alloca "n" +// CIR: %[[FOUR:.*]] = cir.const #cir.int<4> : !s32i +// CIR: cir.store {{.*}}%[[FOUR]], %[[N]] +// CIR: %[[FOUR2:.*]] = cir.const #cir.int<4> : !s32i +// CIR: cir.store {{.*}}%[[FOUR2]], %[[RV]] + +// LLVM-LABEL: define {{.*}}i32 @_Z12non_capturedv() +// LLVM: store i32 4, ptr %{{.*}} +// LLVM: store i32 4, ptr %{{.*}} + +// OGCG-LABEL: define {{.*}}i32 @_Z12non_capturedv() +// OGCG: store i32 4, ptr %{{.*}} +// OGCG: ret i32 4 + +int constexpr_var() { + constexpr int c = 7; + return c * 2; +} + +// CIR-LABEL: cir.func {{.*}} @_Z13constexpr_varv() +// CIR: %[[C:.*]] = cir.alloca "c" +// CIR: cir.store {{.*}}, %[[C]] +// CIR: %[[SEVEN:.*]] = cir.const #cir.int<7> : !s32i +// CIR: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i +// CIR: cir.mul nsw %[[SEVEN]], %[[TWO]] : !s32i + +// LLVM-LABEL: define {{.*}}i32 @_Z13constexpr_varv() +// LLVM: store i32 14, ptr %{{.*}} + +// OGCG-LABEL: define {{.*}}i32 @_Z13constexpr_varv() +// OGCG: ret i32 14 + +_Complex float complex_var() { + const _Complex float z = 1.0f; + return z; +} + +// CIR-LABEL: cir.func {{.*}} @_Z11complex_varv() +// CIR: %[[RV:.*]] = cir.alloca "__retval" +// CIR: %[[Z_ADDR:.*]] = cir.alloca "z" +// CIR: cir.store {{.*}}, %[[Z_ADDR]] +// CIR: %[[Z:.*]] = cir.const #cir.const_complex<#cir.fp<1.000000e+00> : !cir.float, #cir.fp<0.000000e+00> : !cir.float> : !cir.complex<!cir.float> +// CIR: cir.store {{.*}}%[[Z]], %[[RV]] + +// LLVM-LABEL: define {{.*}} @_Z11complex_varv() +// LLVM: store { float, float } { float 1.000000e+00, float 0.000000e+00 }, ptr %{{.*}} + +// OGCG-LABEL: define {{.*}} @_Z11complex_varv() +// OGCG: store float 1.000000e+00, ptr %{{.*}} +// OGCG: store float 0.000000e+00, ptr %{{.*}} + +int global_var = 5; +int captured_ref_to_global() { + int &r = global_var; + return call([=] { return r; }); +} + +// References whose initializer is a constant l-value are emitted as a load of +// the referenced object. + +int const_ref_to_nonconst_global() { + const int &r = global_var; + return r; +} + +// CIR-LABEL: cir.func {{.*}} @_Z28const_ref_to_nonconst_globalv() +// CIR: %[[R:.*]] = cir.alloca "r" +// CIR: cir.store {{.*}}, %[[R]] +// CIR: %[[G:.*]] = cir.const #cir.global_view<@global_var> : !cir.ptr<!s32i> +// CIR: cir.load {{.*}}%[[G]] : !cir.ptr<!s32i>, !s32i + +// LLVM-LABEL: define {{.*}}i32 @_Z28const_ref_to_nonconst_globalv() +// LLVM: store ptr @global_var, ptr %{{.*}} +// LLVM: load i32, ptr @global_var + +// OGCG-LABEL: define {{.*}}i32 @_Z28const_ref_to_nonconst_globalv() +// OGCG: store ptr @global_var, ptr %{{.*}} +// OGCG: load i32, ptr @global_var + +int g_arr[4] = {1, 2, 3, 4}; +int ref_to_array_element() { + int &r = g_arr[1]; + return r; +} + +// CIR-LABEL: cir.func {{.*}} @_Z20ref_to_array_elementv() +// CIR: %[[R:.*]] = cir.alloca "r" +// CIR: cir.store {{.*}}, %[[R]] +// CIR: %[[ELT:.*]] = cir.const #cir.global_view<@g_arr, [1 : i32]> : !cir.ptr<!s32i> +// CIR: cir.load {{.*}}%[[ELT]] : !cir.ptr<!s32i>, !s32i + +// LLVM-LABEL: define {{.*}}i32 @_Z20ref_to_array_elementv() +// LLVM: load i32, ptr getelementptr inbounds nuw (i8, ptr @g_arr, i64 4) + +// OGCG-LABEL: define {{.*}}i32 @_Z20ref_to_array_elementv() +// OGCG: load i32, ptr getelementptr (i8, ptr @g_arr, i64 4) + +_Complex float gcf; +float complex_ref() { + _Complex float &r = gcf; + return __imag__ (r + r); +} + +// CIR-LABEL: cir.func {{.*}} @_Z11complex_refv() +// CIR: %[[R:.*]] = cir.alloca "r" +// CIR: cir.store {{.*}}, %[[R]] +// CIR: %[[G0:.*]] = cir.const #cir.global_view<@gcf> : !cir.ptr<!cir.complex<!cir.float>> +// CIR: %[[V0:.*]] = cir.load {{.*}}%[[G0]] +// CIR: %[[G1:.*]] = cir.const #cir.global_view<@gcf> : !cir.ptr<!cir.complex<!cir.float>> +// CIR: %[[V1:.*]] = cir.load {{.*}}%[[G1]] +// CIR: cir.complex.fadd %[[V0]], %[[V1]] + +// LLVM-LABEL: define {{.*}}float @_Z11complex_refv() +// LLVM: load { float, float }, ptr @gcf +// LLVM: load { float, float }, ptr @gcf + +// OGCG-LABEL: define {{.*}}float @_Z11complex_refv() +// OGCG: load float, ptr @gcf +// OGCG: load float, ptr getelementptr inbounds nuw (i8, ptr @gcf, i64 4) + +// A const pointer whose value is the address of a global. + +int const_ptr_to_global() { + int *const p = &global_var; + return *p; +} + +// CIR-LABEL: cir.func {{.*}} @_Z19const_ptr_to_globalv() +// CIR: %[[P:.*]] = cir.alloca "p" +// CIR: cir.store {{.*}}, %[[P]] +// CIR: %[[G:.*]] = cir.const #cir.global_view<@global_var> : !cir.ptr<!s32i> +// CIR: cir.load {{.*}}%[[G]] : !cir.ptr<!s32i>, !s32i + +// LLVM-LABEL: define {{.*}}i32 @_Z19const_ptr_to_globalv() +// LLVM: load i32, ptr @global_var + +// OGCG-LABEL: define {{.*}}i32 @_Z19const_ptr_to_globalv() +// OGCG: load i32, ptr @global_var + +// Constant pointers and integers that need dedicated LLVM constant expressions. + +int const_absolute_ptr(void *p) { + void *const failed = (void *)-1; + return p == failed; +} + +// CIR-LABEL: cir.func {{.*}} @_Z18const_absolute_ptrPv( +// CIR: %[[FAILED:.*]] = cir.const #cir.ptr<-1 : i64> : !cir.ptr<!void> +// CIR: cir.cmp eq %{{.*}}, %[[FAILED]] : !cir.ptr<!void> + +// LLVM-LABEL: define {{.*}}i32 @_Z18const_absolute_ptrPv( +// LLVM: icmp eq ptr %{{.*}}, inttoptr (i64 -1 to ptr) + +// OGCG-LABEL: define {{.*}}i32 @_Z18const_absolute_ptrPv( +// OGCG: icmp eq ptr %{{.*}}, inttoptr (i64 -1 to ptr) + +void *const_label_addr(int i) { + void *const target = &&l1; + const long diff = (long)&&l2 - (long)&&l1; + if (i) + goto *target; + return (void *)diff; +l1: + return 0; +l2: + return (void *)1; +} + +// CIR-LABEL: cir.func {{.*}} @_Z16const_label_addri( +// CIR: %[[TARGET:.*]] = cir.const #cir.block_addr_info<@_Z16const_label_addri, "l1"> : !cir.ptr<!void> +// CIR: cir.indirect_goto %[[TARGET]] : !cir.ptr<!void> +// CIR: %[[DIFF:.*]] = cir.const #cir.block_addr_diff<@_Z16const_label_addri, "l2", "l1"> : !s64i +// CIR: cir.cast integral %[[DIFF]] : !s64i -> !u64i + +// LLVM-LABEL: define {{.*}}ptr @_Z16const_label_addri( +// LLVM: store ptr inttoptr (i64 sub (i64 ptrtoint (ptr blockaddress(@_Z16const_label_addri, %[[L2:.*]]) to i64), i64 ptrtoint (ptr blockaddress(@_Z16const_label_addri, %[[L1:.*]]) to i64)) to ptr), ptr %{{.*}} +// LLVM: phi ptr [ blockaddress(@_Z16const_label_addri, %[[L1]]), %{{.*}} ] + +// OGCG-LABEL: define {{.*}}ptr @_Z16const_label_addri( +// OGCG: store ptr inttoptr (i64 sub (i64 ptrtoint (ptr blockaddress(@_Z16const_label_addri, %l2) to i64), i64 ptrtoint (ptr blockaddress(@_Z16const_label_addri, %l1) to i64)) to ptr), ptr %{{.*}} +// OGCG: phi ptr [ blockaddress(@_Z16const_label_addri, %l1), %{{.*}} ] + +// Static data members named through a member access. + +struct S { + static constexpr float kf = 1.5f; +}; +float static_member(S s) { return s.kf; } + +// CIR-LABEL: cir.func {{.*}} @_Z13static_member1S( +// CIR-NOT: cir.get_global +// CIR: cir.const #cir.fp<1.500000e+00> : !cir.float + +// LLVM-LABEL: define {{.*}}float @_Z13static_member1S( +// LLVM-NOT: load float +// LLVM: store float 1.500000e+00 + +// OGCG-LABEL: define {{.*}}float @_Z13static_member1S( +// OGCG: ret float 1.500000e+00 + +// Negative cases: the value is loaded from the variable. + +int volatile_const() { + volatile const int v = 3; + return v; +} + +// CIR-LABEL: cir.func {{.*}} @_Z14volatile_constv() +// CIR: %[[V:.*]] = cir.alloca "v" +// CIR: cir.load volatile {{.*}}%[[V]] + +// LLVM-LABEL: define {{.*}}i32 @_Z14volatile_constv() +// LLVM: load volatile i32, ptr %{{.*}} + +// OGCG-LABEL: define {{.*}}i32 @_Z14volatile_constv() +// OGCG: load volatile i32, ptr %{{.*}} + +int non_const_init() { + const int x = side(); + return x; +} + +// CIR-LABEL: cir.func {{.*}} @_Z14non_const_initv() +// CIR: %[[X:.*]] = cir.alloca "x" +// CIR: cir.call @_Z4sidev() +// CIR: cir.load {{.*}}%[[X]] + +// LLVM-LABEL: define {{.*}}i32 @_Z14non_const_initv() +// LLVM: %[[X:.*]] = alloca i32 +// LLVM: call {{.*}}i32 @_Z4sidev() +// LLVM: load i32, ptr %[[X]] + +// OGCG-LABEL: define {{.*}}i32 @_Z14non_const_initv() +// OGCG: %[[X:.*]] = alloca i32 +// OGCG: call {{.*}}i32 @_Z4sidev() +// OGCG: load i32, ptr %[[X]] + +int param(const int p) { + return p; +} + +// CIR-LABEL: cir.func {{.*}} @_Z5parami( +// CIR: %[[P:.*]] = cir.alloca "p" +// CIR: cir.load {{.*}}%[[P]] + +// LLVM-LABEL: define {{.*}}i32 @_Z5parami( +// LLVM: %[[P:.*]] = alloca i32 +// LLVM: load i32, ptr %[[P]] + +// OGCG-LABEL: define {{.*}}i32 @_Z5parami( +// OGCG: %[[P:.*]] = alloca i32 +// OGCG: load i32, ptr %[[P]] + +// The captured constants are used directly instead of being loaded from the +// closure. + +// CIR-LABEL: cir.func {{.*}} lambda {{.*}} @_ZZ8capturedvENK3$_0clEv( +// CIR-NOT: cir.get_member +// CIR: %[[K:.*]] = cir.const #cir.fp<0.899999976> : !cir.float +// CIR: %[[USE:.*]] = cir.call @_Z3usef(%[[K]]) +// CIR: %[[N:.*]] = cir.const #cir.int<4> : !s32i +// CIR: %[[NF:.*]] = cir.cast int_to_float %[[N]] : !s32i -> !cir.float +// CIR: cir.fadd %[[USE]], %[[NF]] : !cir.float + +// LLVM-LABEL: define internal {{.*}}float @"_ZZ8capturedvENK3$_0clEv"( +// LLVM-NOT: getelementptr +// LLVM: %[[USE:.*]] = call {{.*}}float @_Z3usef(float noundef f0x3F666666) +// LLVM: fadd float %[[USE]], 4.000000e+00 + +// OGCG-LABEL: define internal {{.*}}float @"_ZZ8capturedvENK3$_0clEv"( +// OGCG-NOT: getelementptr +// OGCG: %[[USE:.*]] = call {{.*}}float @_Z3usef(float noundef f0x3F666666) +// OGCG: fadd float %[[USE]], 4.000000e+00 + +// A reference to a global is emitted as a reference constant: the global is +// loaded directly rather than through the reference. + +// CIR-LABEL: cir.func {{.*}} lambda {{.*}} @_ZZ22captured_ref_to_globalvENK3$_0clEv( +// CIR: %[[G:.*]] = cir.const #cir.global_view<@global_var> : !cir.ptr<!s32i> +// CIR: cir.load {{.*}}%[[G]] : !cir.ptr<!s32i>, !s32i + +// LLVM-LABEL: define internal {{.*}}i32 @"_ZZ22captured_ref_to_globalvENK3$_0clEv"( +// LLVM: load i32, ptr @global_var + +// OGCG-LABEL: define internal {{.*}}i32 @"_ZZ22captured_ref_to_globalvENK3$_0clEv"( +// OGCG: load i32, ptr @global_var diff --git a/clang/test/CIR/CodeGen/global-init-fenv.cpp b/clang/test/CIR/CodeGen/global-init-fenv.cpp index 80f09793cb50ef..019c77969ae794 100644 --- a/clang/test/CIR/CodeGen/global-init-fenv.cpp +++ b/clang/test/CIR/CodeGen/global-init-fenv.cpp @@ -1,7 +1,7 @@ // --- -ffp-exception-behavior=strict --- // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fno-rtti -fclangir \ // RUN: -ffp-exception-behavior=strict -emit-cir %s -o %t-strict.cir -// RUN: FileCheck --check-prefixes=CIR,CIR-STRICT --input-file=%t-strict.cir %s +// RUN: FileCheck --check-prefixes=CIR,CIR-CONSTRAINED,CIR-STRICT --input-file=%t-strict.cir %s // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fno-rtti -fclangir \ // RUN: -ffp-exception-behavior=strict -emit-llvm %s -o %t-strict-cir.ll // RUN: FileCheck --check-prefixes=LLVM,LLVM-CONSTRAINED,LLVM-STRICT --input-file=%t-strict-cir.ll %s @@ -12,7 +12,7 @@ // --- -ffp-exception-behavior=maytrap --- // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fno-rtti -fclangir \ // RUN: -ffp-exception-behavior=maytrap -emit-cir %s -o %t-maytrap.cir -// RUN: FileCheck --check-prefixes=CIR,CIR-MAYTRAP --input-file=%t-maytrap.cir %s +// RUN: FileCheck --check-prefixes=CIR,CIR-CONSTRAINED,CIR-MAYTRAP --input-file=%t-maytrap.cir %s // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fno-rtti -fclangir \ // RUN: -ffp-exception-behavior=maytrap -emit-llvm %s -o %t-maytrap-cir.ll // RUN: FileCheck --check-prefixes=LLVM,LLVM-CONSTRAINED,LLVM-MAYTRAP --input-file=%t-maytrap-cir.ll %s @@ -69,10 +69,8 @@ float g_const = c_x / c_y; // CIR-CONSTRAINED: cir.func {{.*}} @__cxx_global_var_init.1() attributes {strictfp} // CIR-CONSTRAINED: %[[G:.*]] = cir.get_global @g_const : !cir.ptr<!cir.float> -// CIR-CONSTRAINED: %[[XP:.*]] = cir.get_global @_ZL3c_x : !cir.ptr<!cir.float> -// CIR-CONSTRAINED: %[[X:.*]] = cir.load {{.*}} %[[XP]] -// CIR-CONSTRAINED: %[[YP:.*]] = cir.get_global @_ZL3c_y : !cir.ptr<!cir.float> -// CIR-CONSTRAINED: %[[Y:.*]] = cir.load {{.*}} %[[YP]] +// CIR-CONSTRAINED: %[[X:.*]] = cir.const #cir.fp<1.000000e+00> : !cir.float +// CIR-CONSTRAINED: %[[Y:.*]] = cir.const #cir.fp<1.000000e+01> : !cir.float // CIR-STRICT: %[[R:.*]] = cir.fdiv %[[X]], %[[Y]] : !cir.float fenv<dynamic_rounding_mode = tonearest, except_mode = unknown, strict_except = true> // CIR-MAYTRAP: %[[R:.*]] = cir.fdiv %[[X]], %[[Y]] : !cir.float fenv<dynamic_rounding_mode = tonearest, except_mode = unknown, strict_except = false> // CIR-CONSTRAINED: cir.store {{.*}} %[[R]], %[[G]] : !cir.float, !cir.ptr<!cir.float> @@ -96,10 +94,8 @@ float g_mixed = ce_x / ce_y / nc_z; // CIR-CONSTRAINED: cir.func {{.*}} @__cxx_global_var_init.2() attributes {strictfp} // CIR-DEFAULT: cir.func {{.*}} @__cxx_global_var_init.1() // CIR: %[[G:.*]] = cir.get_global @g_mixed : !cir.ptr<!cir.float> -// CIR: %[[AP:.*]] = cir.get_global @_ZL4ce_x : !cir.ptr<!cir.float> -// CIR: %[[A:.*]] = cir.load {{.*}} %[[AP]] -// CIR: %[[BP:.*]] = cir.get_global @_ZL4ce_y : !cir.ptr<!cir.float> -// CIR: %[[B:.*]] = cir.load {{.*}} %[[BP]] +// CIR: %[[A:.*]] = cir.const #cir.fp<1.000000e+00> : !cir.float +// CIR: %[[B:.*]] = cir.const #cir.fp<1.000000e+01> : !cir.float // CIR-STRICT: %[[R0:.*]] = cir.fdiv %[[A]], %[[B]] : !cir.float fenv<dynamic_rounding_mode = tonearest, except_mode = unknown, strict_except = true> // CIR-MAYTRAP: %[[R0:.*]] = cir.fdiv %[[A]], %[[B]] : !cir.float fenv<dynamic_rounding_mode = tonearest, except_mode = unknown, strict_except = false> // CIR-DEFAULT: %[[R0:.*]] = cir.fdiv %[[A]], %[[B]] : !cir.float diff --git a/clang/test/CIR/CodeGenHIP/lambda-reference-var.hip b/clang/test/CIR/CodeGenHIP/lambda-reference-var.hip new file mode 100644 index 00000000000000..5b3ad34df9a2a8 --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/lambda-reference-var.hip @@ -0,0 +1,161 @@ +// Based on clang/test/CodeGenCUDA/lambda-reference-var.cu. + +// REQUIRES: amdgpu-registered-target + +// RUN: %clang_cc1 -x hip -std=c++11 -triple x86_64-linux-gnu \ +// RUN: -fhip-new-launch-api -fclangir -emit-cir -I%S/../CodeGenCUDA/Inputs \ +// RUN: %s -o %t-host.cir +// RUN: FileCheck --check-prefix=CIR-HOST --input-file=%t-host.cir %s +// RUN: %clang_cc1 -x hip -std=c++11 -triple x86_64-linux-gnu \ +// RUN: -fhip-new-launch-api -fclangir -emit-llvm -I%S/../CodeGenCUDA/Inputs \ +// RUN: %s -o %t-host-cir.ll +// RUN: FileCheck --check-prefix=HOST --input-file=%t-host-cir.ll %s +// RUN: %clang_cc1 -x hip -std=c++11 -triple x86_64-linux-gnu \ +// RUN: -fhip-new-launch-api -emit-llvm -I%S/../CodeGenCUDA/Inputs \ +// RUN: %s -o %t-host.ll +// RUN: FileCheck --check-prefix=HOST --input-file=%t-host.ll %s + +// RUN: %clang_cc1 -x hip -std=c++11 -triple amdgcn-amd-amdhsa -fcuda-is-device \ +// RUN: -fclangir -emit-cir -I%S/../CodeGenCUDA/Inputs %s -o %t-dev.cir +// RUN: FileCheck --check-prefix=CIR-DEV --input-file=%t-dev.cir %s +// RUN: %clang_cc1 -x hip -std=c++11 -triple amdgcn-amd-amdhsa -fcuda-is-device \ +// RUN: -fclangir -emit-llvm -I%S/../CodeGenCUDA/Inputs %s -o %t-dev-cir.ll +// RUN: FileCheck --check-prefix=DEV --input-file=%t-dev-cir.ll %s +// RUN: %clang_cc1 -x hip -std=c++11 -triple amdgcn-amd-amdhsa -fcuda-is-device \ +// RUN: -emit-llvm -I%S/../CodeGenCUDA/Inputs %s -o %t-dev.ll +// RUN: FileCheck --check-prefix=DEV --input-file=%t-dev.ll %s + +#include "cuda.h" + +int global_host_var; +__device__ int global_device_var; + +template<class F> +__global__ void kern(F f) { f(); } + +// CIR-DEV-LABEL: cir.func {{.*}} @_ZZ27dev_capture_dev_ref_by_copyPiENKUlvE_clEv( +// CIR-DEV: %[[G:.*]] = cir.const #cir.global_view<@global_device_var> : !cir.ptr<!s32i> +// CIR-DEV: %[[VAL:.*]] = cir.load {{.*}}%[[G]] : !cir.ptr<!s32i>, !s32i +// CIR-DEV: cir.store {{.*}}%[[VAL]] +// DEV-LABEL: @_ZZ27dev_capture_dev_ref_by_copyPiENKUlvE_clEv( +// DEV: %[[VAL:.*]] = load i32, ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr) +// DEV: store i32 %[[VAL]] +__device__ void dev_capture_dev_ref_by_copy(int *out) { + int &ref = global_device_var; + [=](){ *out = ref;}(); +} + +// DEV-LABEL: @_ZZ28dev_capture_dev_rval_by_copyPiENKUlvE_clEv( +// DEV: store i32 3 +__device__ void dev_capture_dev_rval_by_copy(int *out) { + constexpr int a = 1; + constexpr int b = 2; + constexpr int c = a + b; + [=](){ *out = c;}(); +} + +// DEV-LABEL: @_ZZ26dev_capture_dev_ref_by_refPiENKUlvE_clEv( +// DEV: %[[VAL:.*]] = load i32, ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr) +// DEV: %[[VAL2:.*]] = add nsw i32 %[[VAL]], 1 +// DEV: store i32 %[[VAL2]], ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr) +// DEV: %[[VAL:.*]] = load i32, ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr) +// DEV: store i32 %[[VAL]] +__device__ void dev_capture_dev_ref_by_ref(int *out) { + int &ref = global_device_var; + [&](){ ref++; *out = ref;}(); +} + +// DEV-LABEL: define{{.*}} void @_Z7dev_refPi( +// DEV: %[[VAL:.*]] = load i32, ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr) +// DEV: %[[VAL2:.*]] = add nsw i32 %[[VAL]], 1 +// DEV: store i32 %[[VAL2]], ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr) +// DEV: %[[VAL:.*]] = load i32, ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr) +// DEV: store i32 %[[VAL]] +__device__ void dev_ref(int *out) { + int &ref = global_device_var; + ref++; + *out = ref; +} + +// DEV-LABEL: @_ZZ14dev_lambda_refPiENKUlvE_clEv( +// DEV: %[[VAL:.*]] = load i32, ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr) +// DEV: %[[VAL2:.*]] = add nsw i32 %[[VAL]], 1 +// DEV: store i32 %[[VAL2]], ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr) +// DEV: %[[VAL:.*]] = load i32, ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr) +// DEV: store i32 %[[VAL]] +__device__ void dev_lambda_ref(int *out) { + [=](){ + int &ref = global_device_var; + ref++; + *out = ref; + }(); +} + +// CIR-HOST-LABEL: cir.func {{.*}} @_ZZ29host_capture_host_ref_by_copyPiENKUlvE_clEv( +// CIR-HOST: %[[G:.*]] = cir.const #cir.global_view<@global_host_var> : !cir.ptr<!s32i> +// CIR-HOST: %[[VAL:.*]] = cir.load {{.*}}%[[G]] : !cir.ptr<!s32i>, !s32i +// CIR-HOST: cir.store {{.*}}%[[VAL]] +// HOST-LABEL: @_ZZ29host_capture_host_ref_by_copyPiENKUlvE_clEv( +// HOST: %[[VAL:.*]] = load i32, ptr @global_host_var +// HOST: store i32 %[[VAL]] +void host_capture_host_ref_by_copy(int *out) { + int &ref = global_host_var; + [=](){ *out = ref;}(); +} + +// HOST-LABEL: @_ZZ28host_capture_host_ref_by_refPiENKUlvE_clEv( +// HOST: %[[CAP:.*]] = getelementptr inbounds nuw %{{.*}}, ptr %{{.*}}, i32 0, i32 0 +// HOST: %[[REF:.*]] = load ptr, ptr %[[CAP]] +// HOST: %[[VAL:.*]] = load i32, ptr %[[REF]] +// HOST: %[[VAL2:.*]] = add nsw i32 %[[VAL]], 1 +// HOST: store i32 %[[VAL2]], ptr %[[REF]] +// HOST: %[[VAL:.*]] = load i32, ptr @global_host_var +// HOST: store i32 %[[VAL]] +void host_capture_host_ref_by_ref(int *out) { + int &ref = global_host_var; + [&](){ ref++; *out = ref;}(); +} + +// HOST-LABEL: define{{.*}} void @_Z8host_refPi( +// HOST: %[[VAL:.*]] = load i32, ptr @global_host_var +// HOST: %[[VAL2:.*]] = add nsw i32 %[[VAL]], 1 +// HOST: store i32 %[[VAL2]], ptr @global_host_var +// HOST: %[[VAL:.*]] = load i32, ptr @global_host_var +// HOST: store i32 %[[VAL]] +void host_ref(int *out) { + int &ref = global_host_var; + ref++; + *out = ref; +} + +// HOST-LABEL: @_ZZ15host_lambda_refPiENKUlvE_clEv( +// HOST: %[[VAL:.*]] = load i32, ptr @global_host_var +// HOST: %[[VAL2:.*]] = add nsw i32 %[[VAL]], 1 +// HOST: store i32 %[[VAL2]], ptr @global_host_var +// HOST: %[[VAL:.*]] = load i32, ptr @global_host_var +// HOST: store i32 %[[VAL]] +void host_lambda_ref(int *out) { + [=](){ + int &ref = global_host_var; + ref++; + *out = ref; + }(); +} + +// HOST-LABEL: define{{.*}} void @_Z28dev_capture_host_ref_by_copyPi( +// HOST: %[[CAP:.*]] = getelementptr inbounds nuw %{{.*}}, ptr %{{.*}}, i32 0, i32 1 +// HOST: %[[VAL:.*]] = load i32, ptr @global_host_var +// HOST: store i32 %[[VAL]], ptr %[[CAP]] +// CIR-DEV-LABEL: cir.func {{.*}} @_ZZ28dev_capture_host_ref_by_copyPiENKUlvE_clEv( +// CIR-DEV-NOT: cir.global_view<@global_host_var> +// CIR-DEV: %[[CAP:.*]] = cir.get_member %{{.*}}[1] +// CIR-DEV: %[[VAL:.*]] = cir.load {{.*}}%[[CAP]] : !cir.ptr<!s32i>, !s32i +// CIR-DEV: cir.store {{.*}}%[[VAL]] +// DEV-LABEL: define internal void @_ZZ28dev_capture_host_ref_by_copyPiENKUlvE_clEv( +// DEV: %[[CAP:.*]] = getelementptr inbounds nuw %{{.*}}, ptr %{{.*}}, i32 0, i32 1 +// DEV: %[[VAL:.*]] = load i32, ptr %[[CAP]] +// DEV: store i32 %[[VAL]] +void dev_capture_host_ref_by_copy(int *out) { + int &ref = global_host_var; + kern<<<1, 1>>>([=]__device__() { *out = ref;}); +} diff --git a/clang/test/CIR/CodeGenSYCL/lambda-capture-const.cpp b/clang/test/CIR/CodeGenSYCL/lambda-capture-const.cpp new file mode 100644 index 00000000000000..d10cece1e5b604 --- /dev/null +++ b/clang/test/CIR/CodeGenSYCL/lambda-capture-const.cpp @@ -0,0 +1,33 @@ +// RUN: %clang_cc1 -std=c++20 -fsycl-is-device -triple spirv64-unknown-unknown -fclangir -emit-cir %s -o %t.cir +// RUN: FileCheck --input-file=%t.cir %s -check-prefix=CIR +// RUN: %clang_cc1 -std=c++20 -fsycl-is-device -triple spirv64-unknown-unknown -fclangir -emit-llvm %s -o %t-cir.ll +// RUN: FileCheck --input-file=%t-cir.ll %s -check-prefix=LLVM-OGCG +// RUN: %clang_cc1 -std=c++20 -fsycl-is-device -triple spirv64-unknown-unknown -emit-llvm %s -o %t.ll +// RUN: FileCheck --input-file=%t.ll %s -check-prefix=LLVM-OGCG + +// Verify that references to constant-initialized const variables captured by +// a SYCL kernel lambda are emitted as constants in device code, matching +// classic codegen. + +template <typename KernelName, typename... Ts> +void sycl_kernel_launch(const char *, Ts...) {} + +template <typename KernelName, typename KernelType> +[[clang::sycl_kernel_entry_point(KernelName)]] +void kernel_single_task(KernelType kf) { kf(); } + +struct KN1; + +void capture_const_float(float *out) { + const float k = 0.9f; + kernel_single_task<KN1>([=]() { *out = k; }); +} + +// CIR-LABEL: cir.func {{.*}} @_ZZ19capture_const_floatPfENKUlvE_clEv( +// CIR-NOT: cir.get_member %{{.*}}[1] +// CIR: %[[K:.*]] = cir.const #cir.fp<0.899999976> : !cir.float +// CIR: cir.store {{.*}}%[[K]] + +// LLVM-OGCG-LABEL: define internal spir_func void @_ZZ19capture_const_floatPfENKUlvE_clEv( +// LLVM-OGCG-NOT: getelementptr {{.*}}, i32 0, i32 1 +// LLVM-OGCG: store float f0x3F666666 diff --git a/clang/test/CIR/Lowering/basic.cpp b/clang/test/CIR/Lowering/basic.cpp index 606fe1c3e6d8dc..adc0e8f3c3ac2c 100644 --- a/clang/test/CIR/Lowering/basic.cpp +++ b/clang/test/CIR/Lowering/basic.cpp @@ -22,8 +22,7 @@ int f2() { // CHECK: %[[RV:.*]] = alloca i32, align 4 // CHECK: %[[I_PTR:.*]] = alloca i32, align 4 // CHECK: store i32 2, ptr %[[I_PTR]], align 4 -// CHECK: %[[I:.*]] = load i32, ptr %[[I_PTR]], align 4 -// CHECK: store i32 %[[I]], ptr %[[RV]], align 4 +// CHECK: store i32 2, ptr %[[RV]], align 4 // CHECK: %[[R:.*]] = load i32, ptr %[[RV]], align 4 // CHECK: ret i32 %[[R]] >From 1e7327edf0cae5e3e7d486e98c4b4f8aa00ac042 Mon Sep 17 00:00:00 2001 From: Konstantinos Parasyris <[email protected]> Date: Tue, 13 Oct 2026 15:36:40 -0700 Subject: [PATCH 2/3] [CIR] Merge identical LLVM and OGCG checks in const-var-emit-as-constant.cpp Co-Authored-By: Claude Opus 5.5 (1M context) <[email protected]> --- .../CodeGen/const-var-emit-as-constant.cpp | 149 ++++++------------ 1 file changed, 50 insertions(+), 99 deletions(-) diff --git a/clang/test/CIR/CodeGen/const-var-emit-as-constant.cpp b/clang/test/CIR/CodeGen/const-var-emit-as-constant.cpp index a02c4b3bec4ce0..24e168f0f2f832 100644 --- a/clang/test/CIR/CodeGen/const-var-emit-as-constant.cpp +++ b/clang/test/CIR/CodeGen/const-var-emit-as-constant.cpp @@ -1,9 +1,9 @@ // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir -emit-cir %s -o %t.cir // RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir -emit-llvm %s -o %t-cir.ll -// RUN: FileCheck --check-prefix=LLVM --input-file=%t-cir.ll %s +// RUN: FileCheck --check-prefixes=LLVM,LLVM-OGCG --input-file=%t-cir.ll %s // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm %s -o %t.ll -// RUN: FileCheck --check-prefix=OGCG --input-file=%t.ll %s +// RUN: FileCheck --check-prefixes=OGCG,LLVM-OGCG --input-file=%t.ll %s // References to const, non-volatile variables whose initializer can be // constant evaluated are emitted as the constant value, even when the variable @@ -11,8 +11,7 @@ // reference is to an enclosing variable captured by a lambda. // CIR-NOT: cir.global {{.*}}@__const. -// LLVM-NOT: @__const. -// OGCG-NOT: @__const. +// LLVM-OGCG-NOT: @__const. float use(float); int side(); @@ -37,13 +36,10 @@ int non_captured() { // CIR: %[[FOUR2:.*]] = cir.const #cir.int<4> : !s32i // CIR: cir.store {{.*}}%[[FOUR2]], %[[RV]] -// LLVM-LABEL: define {{.*}}i32 @_Z12non_capturedv() -// LLVM: store i32 4, ptr %{{.*}} -// LLVM: store i32 4, ptr %{{.*}} - -// OGCG-LABEL: define {{.*}}i32 @_Z12non_capturedv() -// OGCG: store i32 4, ptr %{{.*}} -// OGCG: ret i32 4 +// LLVM-OGCG-LABEL: define {{.*}}i32 @_Z12non_capturedv() +// LLVM-OGCG: store i32 4, ptr %{{.*}} +// LLVM: store i32 4, ptr %{{.*}} +// OGCG: ret i32 4 int constexpr_var() { constexpr int c = 7; @@ -57,11 +53,9 @@ int constexpr_var() { // CIR: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i // CIR: cir.mul nsw %[[SEVEN]], %[[TWO]] : !s32i -// LLVM-LABEL: define {{.*}}i32 @_Z13constexpr_varv() -// LLVM: store i32 14, ptr %{{.*}} - -// OGCG-LABEL: define {{.*}}i32 @_Z13constexpr_varv() -// OGCG: ret i32 14 +// LLVM-OGCG-LABEL: define {{.*}}i32 @_Z13constexpr_varv() +// LLVM: store i32 14, ptr %{{.*}} +// OGCG: ret i32 14 _Complex float complex_var() { const _Complex float z = 1.0f; @@ -75,12 +69,10 @@ _Complex float complex_var() { // CIR: %[[Z:.*]] = cir.const #cir.const_complex<#cir.fp<1.000000e+00> : !cir.float, #cir.fp<0.000000e+00> : !cir.float> : !cir.complex<!cir.float> // CIR: cir.store {{.*}}%[[Z]], %[[RV]] -// LLVM-LABEL: define {{.*}} @_Z11complex_varv() -// LLVM: store { float, float } { float 1.000000e+00, float 0.000000e+00 }, ptr %{{.*}} - -// OGCG-LABEL: define {{.*}} @_Z11complex_varv() -// OGCG: store float 1.000000e+00, ptr %{{.*}} -// OGCG: store float 0.000000e+00, ptr %{{.*}} +// LLVM-OGCG-LABEL: define {{.*}} @_Z11complex_varv() +// LLVM: store { float, float } { float 1.000000e+00, float 0.000000e+00 }, ptr %{{.*}} +// OGCG: store float 1.000000e+00, ptr %{{.*}} +// OGCG: store float 0.000000e+00, ptr %{{.*}} int global_var = 5; int captured_ref_to_global() { @@ -102,13 +94,9 @@ int const_ref_to_nonconst_global() { // CIR: %[[G:.*]] = cir.const #cir.global_view<@global_var> : !cir.ptr<!s32i> // CIR: cir.load {{.*}}%[[G]] : !cir.ptr<!s32i>, !s32i -// LLVM-LABEL: define {{.*}}i32 @_Z28const_ref_to_nonconst_globalv() -// LLVM: store ptr @global_var, ptr %{{.*}} -// LLVM: load i32, ptr @global_var - -// OGCG-LABEL: define {{.*}}i32 @_Z28const_ref_to_nonconst_globalv() -// OGCG: store ptr @global_var, ptr %{{.*}} -// OGCG: load i32, ptr @global_var +// LLVM-OGCG-LABEL: define {{.*}}i32 @_Z28const_ref_to_nonconst_globalv() +// LLVM-OGCG: store ptr @global_var, ptr %{{.*}} +// LLVM-OGCG: load i32, ptr @global_var int g_arr[4] = {1, 2, 3, 4}; int ref_to_array_element() { @@ -122,11 +110,8 @@ int ref_to_array_element() { // CIR: %[[ELT:.*]] = cir.const #cir.global_view<@g_arr, [1 : i32]> : !cir.ptr<!s32i> // CIR: cir.load {{.*}}%[[ELT]] : !cir.ptr<!s32i>, !s32i -// LLVM-LABEL: define {{.*}}i32 @_Z20ref_to_array_elementv() -// LLVM: load i32, ptr getelementptr inbounds nuw (i8, ptr @g_arr, i64 4) - -// OGCG-LABEL: define {{.*}}i32 @_Z20ref_to_array_elementv() -// OGCG: load i32, ptr getelementptr (i8, ptr @g_arr, i64 4) +// LLVM-OGCG-LABEL: define {{.*}}i32 @_Z20ref_to_array_elementv() +// LLVM-OGCG: load i32, ptr getelementptr {{.*}}(i8, ptr @g_arr, i64 4) _Complex float gcf; float complex_ref() { @@ -143,13 +128,11 @@ float complex_ref() { // CIR: %[[V1:.*]] = cir.load {{.*}}%[[G1]] // CIR: cir.complex.fadd %[[V0]], %[[V1]] -// LLVM-LABEL: define {{.*}}float @_Z11complex_refv() -// LLVM: load { float, float }, ptr @gcf -// LLVM: load { float, float }, ptr @gcf - -// OGCG-LABEL: define {{.*}}float @_Z11complex_refv() -// OGCG: load float, ptr @gcf -// OGCG: load float, ptr getelementptr inbounds nuw (i8, ptr @gcf, i64 4) +// LLVM-OGCG-LABEL: define {{.*}}float @_Z11complex_refv() +// LLVM: load { float, float }, ptr @gcf +// LLVM: load { float, float }, ptr @gcf +// OGCG: load float, ptr @gcf +// OGCG: load float, ptr getelementptr inbounds nuw (i8, ptr @gcf, i64 4) // A const pointer whose value is the address of a global. @@ -164,11 +147,8 @@ int const_ptr_to_global() { // CIR: %[[G:.*]] = cir.const #cir.global_view<@global_var> : !cir.ptr<!s32i> // CIR: cir.load {{.*}}%[[G]] : !cir.ptr<!s32i>, !s32i -// LLVM-LABEL: define {{.*}}i32 @_Z19const_ptr_to_globalv() -// LLVM: load i32, ptr @global_var - -// OGCG-LABEL: define {{.*}}i32 @_Z19const_ptr_to_globalv() -// OGCG: load i32, ptr @global_var +// LLVM-OGCG-LABEL: define {{.*}}i32 @_Z19const_ptr_to_globalv() +// LLVM-OGCG: load i32, ptr @global_var // Constant pointers and integers that need dedicated LLVM constant expressions. @@ -181,11 +161,8 @@ int const_absolute_ptr(void *p) { // CIR: %[[FAILED:.*]] = cir.const #cir.ptr<-1 : i64> : !cir.ptr<!void> // CIR: cir.cmp eq %{{.*}}, %[[FAILED]] : !cir.ptr<!void> -// LLVM-LABEL: define {{.*}}i32 @_Z18const_absolute_ptrPv( -// LLVM: icmp eq ptr %{{.*}}, inttoptr (i64 -1 to ptr) - -// OGCG-LABEL: define {{.*}}i32 @_Z18const_absolute_ptrPv( -// OGCG: icmp eq ptr %{{.*}}, inttoptr (i64 -1 to ptr) +// LLVM-OGCG-LABEL: define {{.*}}i32 @_Z18const_absolute_ptrPv( +// LLVM-OGCG: icmp eq ptr %{{.*}}, inttoptr (i64 -1 to ptr) void *const_label_addr(int i) { void *const target = &&l1; @@ -205,13 +182,9 @@ void *const_label_addr(int i) { // CIR: %[[DIFF:.*]] = cir.const #cir.block_addr_diff<@_Z16const_label_addri, "l2", "l1"> : !s64i // CIR: cir.cast integral %[[DIFF]] : !s64i -> !u64i -// LLVM-LABEL: define {{.*}}ptr @_Z16const_label_addri( -// LLVM: store ptr inttoptr (i64 sub (i64 ptrtoint (ptr blockaddress(@_Z16const_label_addri, %[[L2:.*]]) to i64), i64 ptrtoint (ptr blockaddress(@_Z16const_label_addri, %[[L1:.*]]) to i64)) to ptr), ptr %{{.*}} -// LLVM: phi ptr [ blockaddress(@_Z16const_label_addri, %[[L1]]), %{{.*}} ] - -// OGCG-LABEL: define {{.*}}ptr @_Z16const_label_addri( -// OGCG: store ptr inttoptr (i64 sub (i64 ptrtoint (ptr blockaddress(@_Z16const_label_addri, %l2) to i64), i64 ptrtoint (ptr blockaddress(@_Z16const_label_addri, %l1) to i64)) to ptr), ptr %{{.*}} -// OGCG: phi ptr [ blockaddress(@_Z16const_label_addri, %l1), %{{.*}} ] +// LLVM-OGCG-LABEL: define {{.*}}ptr @_Z16const_label_addri( +// LLVM-OGCG: store ptr inttoptr (i64 sub (i64 ptrtoint (ptr blockaddress(@_Z16const_label_addri, %[[L2:.*]]) to i64), i64 ptrtoint (ptr blockaddress(@_Z16const_label_addri, %[[L1:.*]]) to i64)) to ptr), ptr %{{.*}} +// LLVM-OGCG: phi ptr [ blockaddress(@_Z16const_label_addri, %[[L1]]), %{{.*}} ] // Static data members named through a member access. @@ -224,12 +197,10 @@ float static_member(S s) { return s.kf; } // CIR-NOT: cir.get_global // CIR: cir.const #cir.fp<1.500000e+00> : !cir.float -// LLVM-LABEL: define {{.*}}float @_Z13static_member1S( -// LLVM-NOT: load float -// LLVM: store float 1.500000e+00 - -// OGCG-LABEL: define {{.*}}float @_Z13static_member1S( -// OGCG: ret float 1.500000e+00 +// LLVM-OGCG-LABEL: define {{.*}}float @_Z13static_member1S( +// LLVM-OGCG-NOT: load float +// LLVM: store float 1.500000e+00 +// OGCG: ret float 1.500000e+00 // Negative cases: the value is loaded from the variable. @@ -242,11 +213,8 @@ int volatile_const() { // CIR: %[[V:.*]] = cir.alloca "v" // CIR: cir.load volatile {{.*}}%[[V]] -// LLVM-LABEL: define {{.*}}i32 @_Z14volatile_constv() -// LLVM: load volatile i32, ptr %{{.*}} - -// OGCG-LABEL: define {{.*}}i32 @_Z14volatile_constv() -// OGCG: load volatile i32, ptr %{{.*}} +// LLVM-OGCG-LABEL: define {{.*}}i32 @_Z14volatile_constv() +// LLVM-OGCG: load volatile i32, ptr %{{.*}} int non_const_init() { const int x = side(); @@ -258,15 +226,10 @@ int non_const_init() { // CIR: cir.call @_Z4sidev() // CIR: cir.load {{.*}}%[[X]] -// LLVM-LABEL: define {{.*}}i32 @_Z14non_const_initv() -// LLVM: %[[X:.*]] = alloca i32 -// LLVM: call {{.*}}i32 @_Z4sidev() -// LLVM: load i32, ptr %[[X]] - -// OGCG-LABEL: define {{.*}}i32 @_Z14non_const_initv() -// OGCG: %[[X:.*]] = alloca i32 -// OGCG: call {{.*}}i32 @_Z4sidev() -// OGCG: load i32, ptr %[[X]] +// LLVM-OGCG-LABEL: define {{.*}}i32 @_Z14non_const_initv() +// LLVM-OGCG: %[[X:.*]] = alloca i32 +// LLVM-OGCG: call {{.*}}i32 @_Z4sidev() +// LLVM-OGCG: load i32, ptr %[[X]] int param(const int p) { return p; @@ -276,13 +239,9 @@ int param(const int p) { // CIR: %[[P:.*]] = cir.alloca "p" // CIR: cir.load {{.*}}%[[P]] -// LLVM-LABEL: define {{.*}}i32 @_Z5parami( -// LLVM: %[[P:.*]] = alloca i32 -// LLVM: load i32, ptr %[[P]] - -// OGCG-LABEL: define {{.*}}i32 @_Z5parami( -// OGCG: %[[P:.*]] = alloca i32 -// OGCG: load i32, ptr %[[P]] +// LLVM-OGCG-LABEL: define {{.*}}i32 @_Z5parami( +// LLVM-OGCG: %[[P:.*]] = alloca i32 +// LLVM-OGCG: load i32, ptr %[[P]] // The captured constants are used directly instead of being loaded from the // closure. @@ -295,15 +254,10 @@ int param(const int p) { // CIR: %[[NF:.*]] = cir.cast int_to_float %[[N]] : !s32i -> !cir.float // CIR: cir.fadd %[[USE]], %[[NF]] : !cir.float -// LLVM-LABEL: define internal {{.*}}float @"_ZZ8capturedvENK3$_0clEv"( -// LLVM-NOT: getelementptr -// LLVM: %[[USE:.*]] = call {{.*}}float @_Z3usef(float noundef f0x3F666666) -// LLVM: fadd float %[[USE]], 4.000000e+00 - -// OGCG-LABEL: define internal {{.*}}float @"_ZZ8capturedvENK3$_0clEv"( -// OGCG-NOT: getelementptr -// OGCG: %[[USE:.*]] = call {{.*}}float @_Z3usef(float noundef f0x3F666666) -// OGCG: fadd float %[[USE]], 4.000000e+00 +// LLVM-OGCG-LABEL: define internal {{.*}}float @"_ZZ8capturedvENK3$_0clEv"( +// LLVM-OGCG-NOT: getelementptr +// LLVM-OGCG: %[[USE:.*]] = call {{.*}}float @_Z3usef(float noundef f0x3F666666) +// LLVM-OGCG: fadd float %[[USE]], 4.000000e+00 // A reference to a global is emitted as a reference constant: the global is // loaded directly rather than through the reference. @@ -312,8 +266,5 @@ int param(const int p) { // CIR: %[[G:.*]] = cir.const #cir.global_view<@global_var> : !cir.ptr<!s32i> // CIR: cir.load {{.*}}%[[G]] : !cir.ptr<!s32i>, !s32i -// LLVM-LABEL: define internal {{.*}}i32 @"_ZZ22captured_ref_to_globalvENK3$_0clEv"( -// LLVM: load i32, ptr @global_var - -// OGCG-LABEL: define internal {{.*}}i32 @"_ZZ22captured_ref_to_globalvENK3$_0clEv"( -// OGCG: load i32, ptr @global_var +// LLVM-OGCG-LABEL: define internal {{.*}}i32 @"_ZZ22captured_ref_to_globalvENK3$_0clEv"( +// LLVM-OGCG: load i32, ptr @global_var >From f3296ce1244e37bc1bddcb0257c2bec66da8181c Mon Sep 17 00:00:00 2001 From: Konstantinos Parasyris <[email protected]> Date: Thu, 15 Oct 2026 13:53:06 -0700 Subject: [PATCH 3/3] [CIR] Share constant-emission type checks with classic codegen Move isConstantEmittableObjectType and the ConstantEmissionKind classification out of CGExpr.cpp into QualType so that classic codegen and CIRGen use the same implementation instead of a copy. --- clang/include/clang/AST/TypeBase.h | 21 +++++++++ clang/lib/AST/Type.cpp | 33 ++++++++++++++ clang/lib/CIR/CodeGen/CIRGenExpr.cpp | 68 ++++------------------------ clang/lib/CodeGen/CGExpr.cpp | 64 ++++---------------------- 4 files changed, 71 insertions(+), 115 deletions(-) diff --git a/clang/include/clang/AST/TypeBase.h b/clang/include/clang/AST/TypeBase.h index 6373794bb69f54..bc8f3f961a8896 100644 --- a/clang/include/clang/AST/TypeBase.h +++ b/clang/include/clang/AST/TypeBase.h @@ -1135,6 +1135,27 @@ class QualType { /// Return true if this is a trivially copyable type bool isTriviallyCopyConstructibleType(const ASTContext &Context) const; + /// Given an object of this canonical, non-reference type, can we safely + /// copy a value out of it based on its initializer? + bool isConstantEmittableObjectType() const; + + /// How a load of a variable of this type can be emitted as a constant. + enum ConstantEmissionKind { + CEK_None, + CEK_AsReferenceOnly, + CEK_AsValueOrReference, + CEK_AsValueOnly + }; + + /// Can we constant-emit a load of a reference to a variable of this type? + /// This is different from predicates like + /// Decl::mightBeUsableInConstantExpressions because we do want it to apply + /// in situations that don't necessarily satisfy the language's rules + /// for this (e.g. C++'s ODR-use rules). For example, we want to able + /// to do this with const float variables even if those variables + /// aren't marked 'constexpr'. + ConstantEmissionKind getConstantEmissionKind() const; + /// Returns true if the type uses postfix declarator syntax, i.e. the /// declarator component appears after the name (arrays, functions). /// Looks through pointer-like types to the pointee. diff --git a/clang/lib/AST/Type.cpp b/clang/lib/AST/Type.cpp index 4c5f2f60a90005..ee8ab663450e77 100644 --- a/clang/lib/AST/Type.cpp +++ b/clang/lib/AST/Type.cpp @@ -3148,6 +3148,39 @@ bool QualType::isTriviallyCopyConstructibleType( /*IsCopyConstructible=*/true); } +bool QualType::isConstantEmittableObjectType() const { + assert(isCanonical()); + assert(!getTypePtr()->isReferenceType()); + + // Must be const-qualified but non-volatile. + Qualifiers Qs = getLocalQualifiers(); + if (!Qs.hasConst() || Qs.hasVolatile()) + return false; + + // Otherwise, all object types satisfy this except C++ classes with + // mutable subobjects or non-trivial copy/destroy behavior. + if (const auto *RT = dyn_cast<RecordType>(getTypePtr())) + if (const auto *RD = dyn_cast<CXXRecordDecl>(RT->getDecl())) { + RD = RD->getDefinitionOrSelf(); + if (RD->hasMutableFields() || !RD->isTrivial()) + return false; + } + + return true; +} + +QualType::ConstantEmissionKind QualType::getConstantEmissionKind() const { + QualType Canon = getCanonicalType(); + if (const auto *Ref = dyn_cast<ReferenceType>(Canon)) { + if (Ref->getPointeeType().isConstantEmittableObjectType()) + return CEK_AsValueOrReference; + return CEK_AsReferenceOnly; + } + if (Canon.isConstantEmittableObjectType()) + return CEK_AsValueOnly; + return CEK_None; +} + bool QualType::isNonWeakInMRRWithObjCWeak(const ASTContext &Context) const { return !Context.getLangOpts().ObjCAutoRefCount && Context.getLangOpts().ObjCWeak && diff --git a/clang/lib/CIR/CodeGen/CIRGenExpr.cpp b/clang/lib/CIR/CodeGen/CIRGenExpr.cpp index b15f80bc461e6e..9b157080e180e2 100644 --- a/clang/lib/CIR/CodeGen/CIRGenExpr.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenExpr.cpp @@ -3056,54 +3056,6 @@ Address CIRGenFunction::createDefaultAlignTempAlloca(mlir::Type ty, return createTempAlloca(ty, align, loc, name); } -/// Given an object of the given canonical type, can we safely copy a -/// value out of it based on its initializer? -static bool isConstantEmittableObjectType(QualType type) { - assert(type.isCanonical()); - assert(!type->isReferenceType()); - - // Must be const-qualified but non-volatile. - Qualifiers qs = type.getLocalQualifiers(); - if (!qs.hasConst() || qs.hasVolatile()) - return false; - - // Otherwise, all object types satisfy this except C++ classes with - // mutable subobjects or non-trivial copy/destroy behavior. - if (const auto *rt = dyn_cast<clang::RecordType>(type)) - if (const auto *rd = dyn_cast<CXXRecordDecl>(rt->getDecl())) { - rd = rd->getDefinitionOrSelf(); - if (rd->hasMutableFields() || !rd->isTrivial()) - return false; - } - - return true; -} - -/// Can we constant-emit a load of a reference to a variable of the -/// given type? This is different from predicates like -/// Decl::mightBeUsableInConstantExpressions because we do want it to apply -/// in situations that don't necessarily satisfy the language's rules -/// for this (e.g. C++'s ODR-use rules). For example, we want to able -/// to do this with const float variables even if those variables -/// aren't marked 'constexpr'. -enum ConstantEmissionKind { - CEK_None, - CEK_AsReferenceOnly, - CEK_AsValueOrReference, - CEK_AsValueOnly -}; -static ConstantEmissionKind checkVarTypeForConstantEmission(QualType type) { - type = type.getCanonicalType(); - if (const auto *ref = dyn_cast<ReferenceType>(type)) { - if (isConstantEmittableObjectType(ref->getPointeeType())) - return CEK_AsValueOrReference; - return CEK_AsReferenceOnly; - } - if (isConstantEmittableObjectType(type)) - return CEK_AsValueOnly; - return CEK_None; -} - /// Try to emit a reference to the given value without producing it as /// an l-value. This is just an optimization, but it avoids us needing /// to emit global copies of variables if they're named without triggering @@ -3115,16 +3067,16 @@ CIRGenFunction::tryEmitAsConstant(const DeclRefExpr *refExpr) { const ValueDecl *value = refExpr->getDecl(); // The value needs to be an enum constant or a constant variable. - ConstantEmissionKind cek; + QualType::ConstantEmissionKind cek; if (isa<ParmVarDecl>(value)) - cek = CEK_None; + cek = QualType::CEK_None; else if (const auto *var = dyn_cast<VarDecl>(value)) - cek = checkVarTypeForConstantEmission(var->getType()); + cek = var->getType().getConstantEmissionKind(); else if (isa<EnumConstantDecl>(value)) - cek = CEK_AsValueOnly; + cek = QualType::CEK_AsValueOnly; else - cek = CEK_None; - if (cek == CEK_None) + cek = QualType::CEK_None; + if (cek == QualType::CEK_None) return ConstantEmission(); Expr::EvalResult result; @@ -3132,18 +3084,14 @@ CIRGenFunction::tryEmitAsConstant(const DeclRefExpr *refExpr) { QualType resultType; // It's best to evaluate all the way as an r-value if that's permitted. - if (cek != CEK_AsReferenceOnly && + if (cek != QualType::CEK_AsReferenceOnly && refExpr->EvaluateAsRValue(result, getContext())) { resultIsReference = false; resultType = refExpr->getType().getUnqualifiedType(); - - // Otherwise, try to evaluate as an l-value. - } else if (cek != CEK_AsValueOnly && + } else if (cek != QualType::CEK_AsValueOnly && refExpr->EvaluateAsLValue(result, getContext())) { resultIsReference = true; resultType = value->getType(); - - // Failure. } else { return ConstantEmission(); } diff --git a/clang/lib/CodeGen/CGExpr.cpp b/clang/lib/CodeGen/CGExpr.cpp index dcc2443994d21b..5051d1027a9737 100644 --- a/clang/lib/CodeGen/CGExpr.cpp +++ b/clang/lib/CodeGen/CGExpr.cpp @@ -1949,53 +1949,6 @@ LValue CodeGenFunction::EmitLValueHelper(const Expr *E, } } -/// Given an object of the given canonical type, can we safely copy a -/// value out of it based on its initializer? -static bool isConstantEmittableObjectType(QualType type) { - assert(type.isCanonical()); - assert(!type->isReferenceType()); - - // Must be const-qualified but non-volatile. - Qualifiers qs = type.getLocalQualifiers(); - if (!qs.hasConst() || qs.hasVolatile()) return false; - - // Otherwise, all object types satisfy this except C++ classes with - // mutable subobjects or non-trivial copy/destroy behavior. - if (const auto *RT = dyn_cast<RecordType>(type)) - if (const auto *RD = dyn_cast<CXXRecordDecl>(RT->getDecl())) { - RD = RD->getDefinitionOrSelf(); - if (RD->hasMutableFields() || !RD->isTrivial()) - return false; - } - - return true; -} - -/// Can we constant-emit a load of a reference to a variable of the -/// given type? This is different from predicates like -/// Decl::mightBeUsableInConstantExpressions because we do want it to apply -/// in situations that don't necessarily satisfy the language's rules -/// for this (e.g. C++'s ODR-use rules). For example, we want to able -/// to do this with const float variables even if those variables -/// aren't marked 'constexpr'. -enum ConstantEmissionKind { - CEK_None, - CEK_AsReferenceOnly, - CEK_AsValueOrReference, - CEK_AsValueOnly -}; -static ConstantEmissionKind checkVarTypeForConstantEmission(QualType type) { - type = type.getCanonicalType(); - if (const auto *ref = dyn_cast<ReferenceType>(type)) { - if (isConstantEmittableObjectType(ref->getPointeeType())) - return CEK_AsValueOrReference; - return CEK_AsReferenceOnly; - } - if (isConstantEmittableObjectType(type)) - return CEK_AsValueOnly; - return CEK_None; -} - /// Try to emit a reference to the given value without producing it as /// an l-value. This is just an optimization, but it avoids us needing /// to emit global copies of variables if they're named without triggering @@ -2007,30 +1960,31 @@ CodeGenFunction::tryEmitAsConstant(const DeclRefExpr *RefExpr) { const ValueDecl *Value = RefExpr->getDecl(); // The value needs to be an enum constant or a constant variable. - ConstantEmissionKind CEK; + QualType::ConstantEmissionKind CEK; if (isa<ParmVarDecl>(Value)) { - CEK = CEK_None; + CEK = QualType::CEK_None; } else if (const auto *var = dyn_cast<VarDecl>(Value)) { - CEK = checkVarTypeForConstantEmission(var->getType()); + CEK = var->getType().getConstantEmissionKind(); } else if (isa<EnumConstantDecl>(Value)) { - CEK = CEK_AsValueOnly; + CEK = QualType::CEK_AsValueOnly; } else { - CEK = CEK_None; + CEK = QualType::CEK_None; } - if (CEK == CEK_None) return ConstantEmission(); + if (CEK == QualType::CEK_None) + return ConstantEmission(); Expr::EvalResult result; bool resultIsReference; QualType resultType; // It's best to evaluate all the way as an r-value if that's permitted. - if (CEK != CEK_AsReferenceOnly && + if (CEK != QualType::CEK_AsReferenceOnly && RefExpr->EvaluateAsRValue(result, getContext())) { resultIsReference = false; resultType = RefExpr->getType().getUnqualifiedType(); // Otherwise, try to evaluate as an l-value. - } else if (CEK != CEK_AsValueOnly && + } else if (CEK != QualType::CEK_AsValueOnly && RefExpr->EvaluateAsLValue(result, getContext())) { resultIsReference = true; resultType = Value->getType(); _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
