llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clangir Author: Steffen Larsen (steffenlarsen) <details> <summary>Changes</summary> Port the rest of CodeGenFunction::tryEmitAsConstant from OGCG. Previously only EnumConstantDecl was folded, so const-qualified variables were always loaded from memory. Now const, non-volatile variables of scalar, complex and trivially copyable record types are emitted as constants when their value can be evaluated as an r-value. Additionally, references that evaluate to an l-value are emitted as a constant address. This implements ConstantEmission::getReferenceLValue and the reference case of emitScalarConstant. Assisted-by: Claude Opus 5.5 --- Patch is 21.72 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/231872.diff 8 Files Affected: - (modified) clang/lib/CIR/CodeGen/CIRGenExpr.cpp (+109-14) - (modified) clang/lib/CIR/CodeGen/CIRGenFunction.h (+4-3) - (modified) clang/test/CIR/CodeGen/assume-attr.cpp (+6-11) - (modified) clang/test/CIR/CodeGen/basic.cpp (+2-2) - (added) clang/test/CIR/CodeGen/decl-ref-constant-emission.cpp (+112) - (modified) clang/test/CIR/CodeGen/global-init-fenv.cpp (+6-8) - (added) clang/test/CIR/CodeGenCUDA/lambda-reference-var.cu (+54) - (modified) clang/test/CIR/Lowering/basic.cpp (+1-2) ``````````diff diff --git a/clang/lib/CIR/CodeGen/CIRGenExpr.cpp b/clang/lib/CIR/CodeGen/CIRGenExpr.cpp index 9c71d987c37e1..b30896f025f78 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,6 +3056,54 @@ 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 @@ -3068,23 +3117,65 @@ 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 (getLangOpts().CUDAIsDevice && result.Val.isLValue() && + refExpr->refersToEnclosingVariableOrCapture()) { + const auto *md = dyn_cast_or_null<CXXMethodDecl>(curCodeDecl); + if (isLambdaMethod(md) && md->getOverloadedOperator() == OO_Call) { + const APValue::LValueBase &base = result.Val.getLValueBase(); + if (const auto *d = base.dyn_cast<const ValueDecl *>()) { + if (const auto *vd = dyn_cast<VarDecl>(d)) { + if (!vd->hasAttr<CUDADeviceAttr>()) + return ConstantEmission(); + } + } + } + } // Emit as a constant. // FIXME(cir): have emitAbstract build a TypedAttr instead (this requires @@ -3096,6 +3187,10 @@ CIRGenFunction::tryEmitAsConstant(const DeclRefExpr *refExpr) { 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 +3204,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/CIRGenFunction.h b/clang/lib/CIR/CodeGen/CIRGenFunction.h index 8937e237a6414..8010be0b9d923 100644 --- a/clang/lib/CIR/CodeGen/CIRGenFunction.h +++ b/clang/lib/CIR/CodeGen/CIRGenFunction.h @@ -802,9 +802,10 @@ class CIRGenFunction : public CIRGenTypeCache { bool isReference() const { return valueAndIsReference.getInt(); } LValue getReferenceLValue(CIRGenFunction &cgf, Expr *refExpr) const { assert(isReference()); - cgf.cgm.errorNYI(refExpr->getSourceRange(), - "ConstantEmission::getReferenceLValue"); - return {}; + auto attr = mlir::cast<mlir::TypedAttr>(valueAndIsReference.getPointer()); + mlir::Value val = cgf.getBuilder().getConstant( + cgf.getLoc(refExpr->getSourceRange()), attr); + return cgf.makeNaturalAlignAddrLValue(val, refExpr->getType()); } mlir::TypedAttr getValue() const { diff --git a/clang/test/CIR/CodeGen/assume-attr.cpp b/clang/test/CIR/CodeGen/assume-attr.cpp index a3ea3bb6faaf6..ee4bdec795a7f 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 6a31e07a38285..cbf99c0d0eaba 100644 --- a/clang/test/CIR/CodeGen/basic.cpp +++ b/clang/test/CIR/CodeGen/basic.cpp @@ -49,8 +49,8 @@ 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: cir.store{{.*}} %[[I]], %[[RV]] : !s32i, !cir.ptr<!s32i> +// CHECK: %[[TWO_FOLDED:.*]] = cir.const #cir.int<2> : !s32i +// CHECK: cir.store{{.*}} %[[TWO_FOLDED]], %[[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/decl-ref-constant-emission.cpp b/clang/test/CIR/CodeGen/decl-ref-constant-emission.cpp new file mode 100644 index 0000000000000..68ca2e6b6ebd9 --- /dev/null +++ b/clang/test/CIR/CodeGen/decl-ref-constant-emission.cpp @@ -0,0 +1,112 @@ +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir -emit-cir %s -o - \ +// RUN: | FileCheck --check-prefix=CIR %s +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir -emit-llvm %s -o - \ +// RUN: | FileCheck --check-prefix=LLVM %s +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm %s -o - \ +// RUN: | FileCheck --check-prefix=OGCG %s + +int g; +int &rg = g; +int ref_only() { return rg; } + +// CIR-LABEL: cir.func {{.*}} @_Z8ref_onlyv() +// CIR: %[[RETVAL:.*]] = cir.alloca "__retval" +// CIR-NEXT: %[[G:.*]] = cir.const #cir.global_view<@g> : !cir.ptr<!s32i> +// CIR-NEXT: %[[VAL:.*]] = cir.load align(4) %[[G]] : !cir.ptr<!s32i>, !s32i +// CIR-NEXT: cir.store %[[VAL]], %[[RETVAL]] + +// LLVM-LABEL: define {{.*}} i32 @_Z8ref_onlyv() +// LLVM: %[[RETVAL:.*]] = alloca i32 +// LLVM-NEXT: %[[VAL:.*]] = load i32, ptr @g, align 4 +// LLVM-NEXT: store i32 %[[VAL]], ptr %[[RETVAL]] + +// OGCG-LABEL: define {{.*}} i32 @_Z8ref_onlyv() +// OGCG: %[[VAL:.*]] = load i32, ptr @g, align 4 +// OGCG-NEXT: ret i32 %[[VAL]] + +int g2; +const int &crg = g2; +int value_or_ref_as_ref() { return crg; } + +// CIR-LABEL: cir.func {{.*}} @_Z19value_or_ref_as_refv() +// CIR: %[[RETVAL:.*]] = cir.alloca "__retval" +// CIR-NEXT: %[[G2:.*]] = cir.const #cir.global_view<@g2> : !cir.ptr<!s32i> +// CIR-NEXT: %[[VAL:.*]] = cir.load align(4) %[[G2]] : !cir.ptr<!s32i>, !s32i +// CIR-NEXT: cir.store %[[VAL]], %[[RETVAL]] + +// LLVM-LABEL: define {{.*}} i32 @_Z19value_or_ref_as_refv() +// LLVM: %[[RETVAL:.*]] = alloca i32 +// LLVM-NEXT: %[[VAL:.*]] = load i32, ptr @g2, align 4 +// LLVM-NEXT: store i32 %[[VAL]], ptr %[[RETVAL]] + +// OGCG-LABEL: define {{.*}} i32 @_Z19value_or_ref_as_refv() +// OGCG: %[[VAL:.*]] = load i32, ptr @g2, align 4 +// OGCG-NEXT: ret i32 %[[VAL]] + +const int k = 7; +const int &crk = k; +int value_or_ref_as_value() { return crk; } + +// CIR-LABEL: cir.func {{.*}} @_Z21value_or_ref_as_valuev() +// CIR: %[[RETVAL:.*]] = cir.alloca "__retval" +// CIR-NEXT: %[[SEVEN:.*]] = cir.const #cir.int<7> : !s32i +// CIR-NEXT: cir.store %[[SEVEN]], %[[RETVAL]] + +// LLVM-LABEL: define {{.*}} i32 @_Z21value_or_ref_as_valuev() +// LLVM: %[[RETVAL:.*]] = alloca i32 +// LLVM-NEXT: store i32 7, ptr %[[RETVAL]] + +// OGCG-LABEL: define {{.*}} i32 @_Z21value_or_ref_as_valuev() +// OGCG: ret i32 7 + +struct S { static const int sm = 9; }; +int static_member() { return S::sm; } + +// CIR-LABEL: cir.func {{.*}} @_Z13static_memberv() +// CIR: %[[RETVAL:.*]] = cir.alloca "__retval" +// CIR-NEXT: %[[NINE:.*]] = cir.const #cir.int<9> : !s32i +// CIR-NEXT: cir.store %[[NINE]], %[[RETVAL]] + +// LLVM-LABEL: define {{.*}} i32 @_Z13static_memberv() +// LLVM: %[[RETVAL:.*]] = alloca i32 +// LLVM-NEXT: store i32 9, ptr %[[RETVAL]] + +// OGCG-LABEL: define {{.*}} i32 @_Z13static_memberv() +// OGCG: ret i32 9 + +const _Complex float cc = 1.0f; +_Complex float complex_value() { return cc; } + +// CIR-LABEL: cir.func {{.*}} @_Z13complex_valuev() +// CIR: %[[RETVAL:.*]] = cir.alloca "__retval" +// CIR-NEXT: %[[CPLX:.*]] = cir.const #cir.const_complex<#cir.fp<1.000000e+00> : !cir.float, #cir.fp<0.000000e+00> : !cir.float> : !cir.complex<!cir.float> +// CIR-NEXT: cir.store align(4) %[[CPLX]], %[[RETVAL]] + +// LLVM-LABEL: define {{.*}} <2 x float> @_Z13complex_valuev() +// LLVM: %[[RETVAL:.*]] = alloca { float, float }, align 4 +// LLVM-NEXT: store { float, float } { float 1.000000e+00, float 0.000000e+00 }, ptr %[[RETVAL]] + +// OGCG-LABEL: define {{.*}} <2 x float> @_Z13complex_valuev() +// OGCG: %[[REALP:.*]] = getelementptr inbounds nuw { float, float }, ptr %{{.*}}, i32 0, i32 0 +// OGCG-NEXT: %[[IMAGP:.*]] = getelementptr inbounds nuw { float, float }, ptr %{{.*}}, i32 0, i32 1 +// OGCG-NEXT: store float 1.000000e+00, ptr %[[REALP]] +// OGCG-NEXT: store float 0.000000e+00, ptr %[[IMAGP]] + +_Complex float cg; +_Complex float &rcg = cg; +_Complex float complex_ref() { return rcg; } + +// CIR-LABEL: cir.func {{.*}} @_Z11complex_refv() +// CIR: %[[RETVAL:.*]] = cir.alloca "__retval" +// CIR-NEXT: %[[CG:.*]] = cir.const #cir.global_view<@cg> : !cir.ptr<!cir.complex<!cir.float>> +// CIR-NEXT: %[[VAL:.*]] = cir.load align(4) %[[CG]] : !cir.ptr<!cir.complex<!cir.float>>, !cir.complex<!cir.float> +// CIR-NEXT: cir.store align(4) %[[VAL]], %[[RETVAL]] + +// LLVM-LABEL: define {{.*}} <2 x float> @_Z11complex_refv() +// LLVM: %[[RETVAL:.*]] = alloca { float, float }, align 4 +// LLVM-NEXT: %[[VAL:.*]] = load { float, float }, ptr @cg, align 4 +// LLVM-NEXT: store { float, float } %[[VAL]], ptr %[[RETVAL]] + +// OGCG-LABEL: define {{.*}} <2 x float> @_Z11complex_refv() +// OGCG: %{{.*}} = load float, ptr @cg, align 4 +// OGCG-NEXT: %{{.*}} = load float, ptr getelementptr inbounds nuw (i8, ptr @cg, i64 4), align 4 diff --git a/clang/test/CIR/CodeGen/global-init-fenv.cpp b/clang/test/CIR/CodeGen/global-init-fenv.cpp index 80f09793cb50e..2cd580fa2c0f5 100644 --- a/clang/test/CIR/CodeGen/global-init-fenv.cpp +++ b/clang/test/CIR/CodeGen/global-init-fenv.cpp @@ -69,11 +69,11 @@ 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-STRICT: %[[X:.*]] = cir.const #cir.fp<1.000000e+00> : !cir.float +// CIR-STRICT: %[[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: %[[X:.*]] = cir.const #cir.fp<1.000000e+00> : !cir.float +// CIR-MAYTRAP: %[[Y:.*]] = cir.const #cir.fp<1.000000e+01> : !cir.float // 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 +96,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/CodeGenCUDA/lambda-reference-var.cu b/clang/test/CIR/CodeGenCUDA/lambda-reference-var.cu new file mode 100644 index 0000000000000..d68d9bca0af34 --- /dev/null +++ b/clang/test/CIR/CodeGenCUDA/lambda-reference-var.cu @@ -0,0 +1,54 @@ +#include "Inputs/cuda.h" + +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -x cuda -fclangir -fcuda-is-device -emit-cir %s -o - \ +// RUN: | FileCheck --check-prefix=CIR %s +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -x cuda -fclangir -fcuda-is-device -emit-llvm %s -o - \ +// RUN: | FileCheck --check-prefix=LLVM %s +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -x cuda -fcuda-is-device -emit-llvm %s -o - \ +// RUN: | FileCheck --check-prefix=LLVM %s + +// RUN: %clang_cc1 -triple amdgpu9-amd-amdhsa -x hip -fclangir -fcuda-is-device -emit-cir %s -o - \ +// RUN: | FileCheck --check-prefix=CIR %s +// RUN: %clang_cc1 -triple amdgpu9-amd-amdhsa -x hip -fclangir -fcuda-is-device -emit-llvm %s -o - \ +// RUN: | FileCheck --check-prefix=LLVM %s +// RUN: %clang_cc1 -triple amdgpu9-amd-amdhsa -x hip -fcuda-is-device -emit-llvm %s -o - \ +// RUN: | FileCheck --check-prefix=LLVM %s + +int global_host_var; +__device__ int global_device_var; + +template <class F> +__global__ void kern(F f) { f(); } + +__device__ void dev_capture_dev_ref_by_copy(int *out) { + int &ref = global_device_var; + [=]() { *out = ref; }(); +} + +// CIR-LABEL: cir.func {{.*}} @_ZZ27dev_capture_dev_ref_by_copyPiENKUlvE_clEv( +// CIR: %[[GV:.*]] = cir.const #cir.global_view<@global_device_var> : !cir.ptr<!s32i> +// CIR-NEXT: %[[VAL:.*]] = cir.load align(4) %[[GV]] : !cir.ptr<!s32i>, !s32i +// CIR: cir.store align(4) %[[VAL]], %{{.*}} : !s32i, !cir.ptr<!s32i> + +// LLVM-LABEL: define internal void @_ZZ27dev_capture_dev_ref_by_copyPiENKUlvE_clEv( +// LLVM: %[[VAL:.*]] = load i32, ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr), align 4 +// LLVM: store i32 %[[VAL]], ptr %{{... [truncated] `````````` </details> https://github.com/llvm/llvm-project/pull/231872 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
