Skip to content

Commit a7976c0

Browse files
ronliebronlieb
authored
merge main into amd-staging (llvm#3558)
2 parents ca6337f + 54df6cf commit a7976c0

File tree

243 files changed

+5510
-2128
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

243 files changed

+5510
-2128
lines changed

‎clang/include/clang/Basic/BuiltinsAMDGPU.def

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -508,6 +508,9 @@ TARGET_BUILTIN(__builtin_amdgcn_s_barrier_signal, "vIi", "n", "gfx12-insts")
508508
TARGET_BUILTIN(__builtin_amdgcn_s_barrier_signal_var, "vv*i", "n", "gfx12-insts")
509509
TARGET_BUILTIN(__builtin_amdgcn_s_barrier_wait, "vIs", "n", "gfx12-insts")
510510
TARGET_BUILTIN(__builtin_amdgcn_s_barrier_signal_isfirst, "bIi", "n", "gfx12-insts")
511+
TARGET_BUILTIN(__builtin_amdgcn_s_barrier_init, "vv*i", "n", "gfx12-insts")
512+
TARGET_BUILTIN(__builtin_amdgcn_s_barrier_join, "vv*", "n", "gfx12-insts")
513+
TARGET_BUILTIN(__builtin_amdgcn_s_barrier_leave, "vIs", "n", "gfx12-insts")
511514
TARGET_BUILTIN(__builtin_amdgcn_s_get_barrier_state, "Uii", "n", "gfx12-insts")
512515
TARGET_BUILTIN(__builtin_amdgcn_s_get_named_barrier_state, "Uiv*", "n", "gfx12-insts")
513516
TARGET_BUILTIN(__builtin_amdgcn_s_prefetch_data, "vvC*Ui", "nc", "gfx12-insts")

‎clang/include/clang/CIR/Dialect/IR/CIRDataLayout.h

Lines changed: 15 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -34,6 +34,21 @@ class CIRDataLayout {
3434
void reset(mlir::DataLayoutSpecInterface spec);
3535

3636
bool isBigEndian() const { return bigEndian; }
37+
38+
/// Returns the maximum number of bytes that may be overwritten by
39+
/// storing the specified type.
40+
///
41+
/// If Ty is a scalable vector type, the scalable property will be set and
42+
/// the runtime size will be a positive integer multiple of the base size.
43+
///
44+
/// For example, returns 5 for i36 and 10 for x86_fp80.
45+
llvm::TypeSize getTypeStoreSize(mlir::Type ty) const {
46+
llvm::TypeSize baseSize = getTypeSizeInBits(ty);
47+
return {llvm::divideCeil(baseSize.getKnownMinValue(), 8),
48+
baseSize.isScalable()};
49+
}
50+
51+
llvm::TypeSize getTypeSizeInBits(mlir::Type ty) const;
3752
};
3853

3954
} // namespace cir

‎clang/include/clang/CIR/MissingFeatures.h

Lines changed: 10 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -87,7 +87,6 @@ struct MissingFeatures {
8787
static bool setFunctionAttributes() { return false; }
8888

8989
// CallOp handling
90-
static bool opCallPseudoDtor() { return false; }
9190
static bool opCallAggregateArgs() { return false; }
9291
static bool opCallPaddingArgs() { return false; }
9392
static bool opCallABIExtendArg() { return false; }
@@ -162,6 +161,13 @@ struct MissingFeatures {
162161
static bool addressIsKnownNonNull() { return false; }
163162
static bool addressPointerAuthInfo() { return false; }
164163

164+
// Atomic
165+
static bool atomicExpr() { return false; }
166+
static bool atomicInfo() { return false; }
167+
static bool atomicInfoGetAtomicPointer() { return false; }
168+
static bool atomicInfoGetAtomicAddress() { return false; }
169+
static bool atomicUseLibCall() { return false; }
170+
165171
// Misc
166172
static bool abiArgInfo() { return false; }
167173
static bool addHeapAllocSiteMetadata() { return false; }
@@ -197,7 +203,9 @@ struct MissingFeatures {
197203
static bool ctorMemcpyizer() { return false; }
198204
static bool cudaSupport() { return false; }
199205
static bool cxxRecordStaticMembers() { return false; }
206+
static bool dataLayoutTypeIsSized() { return false; }
200207
static bool dataLayoutTypeAllocSize() { return false; }
208+
static bool dataLayoutTypeStoreSize() { return false; }
201209
static bool deferredCXXGlobalInit() { return false; }
202210
static bool ehCleanupFlags() { return false; }
203211
static bool ehCleanupScope() { return false; }
@@ -238,6 +246,7 @@ struct MissingFeatures {
238246
static bool objCBlocks() { return false; }
239247
static bool objCGC() { return false; }
240248
static bool objCLifetime() { return false; }
249+
static bool openCL() { return false; }
241250
static bool openMP() { return false; }
242251
static bool opTBAA() { return false; }
243252
static bool peepholeProtection() { return false; }

‎clang/lib/CIR/CodeGen/CIRGenAtomic.cpp

Lines changed: 230 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,230 @@
1+
//===--- CIRGenAtomic.cpp - Emit CIR for atomic operations ----------------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
//
9+
// This file contains the code for emitting atomic operations.
10+
//
11+
//===----------------------------------------------------------------------===//
12+
13+
#include "CIRGenFunction.h"
14+
#include "clang/CIR/MissingFeatures.h"
15+
16+
using namespace clang;
17+
using namespace clang::CIRGen;
18+
using namespace cir;
19+
20+
namespace {
21+
class AtomicInfo {
22+
CIRGenFunction &cgf;
23+
QualType atomicTy;
24+
QualType valueTy;
25+
uint64_t atomicSizeInBits = 0;
26+
uint64_t valueSizeInBits = 0;
27+
CharUnits atomicAlign;
28+
CharUnits valueAlign;
29+
TypeEvaluationKind evaluationKind = cir::TEK_Scalar;
30+
LValue lvalue;
31+
mlir::Location loc;
32+
33+
public:
34+
AtomicInfo(CIRGenFunction &cgf, LValue &lvalue, mlir::Location loc)
35+
: cgf(cgf), loc(loc) {
36+
assert(!lvalue.isGlobalReg());
37+
ASTContext &ctx = cgf.getContext();
38+
if (lvalue.isSimple()) {
39+
atomicTy = lvalue.getType();
40+
if (auto *ty = atomicTy->getAs<AtomicType>())
41+
valueTy = ty->getValueType();
42+
else
43+
valueTy = atomicTy;
44+
evaluationKind = cgf.getEvaluationKind(valueTy);
45+
46+
TypeInfo valueTypeInfo = ctx.getTypeInfo(valueTy);
47+
TypeInfo atomicTypeInfo = ctx.getTypeInfo(atomicTy);
48+
uint64_t valueAlignInBits = valueTypeInfo.Align;
49+
uint64_t atomicAlignInBits = atomicTypeInfo.Align;
50+
valueSizeInBits = valueTypeInfo.Width;
51+
atomicSizeInBits = atomicTypeInfo.Width;
52+
assert(valueSizeInBits <= atomicSizeInBits);
53+
assert(valueAlignInBits <= atomicAlignInBits);
54+
55+
atomicAlign = ctx.toCharUnitsFromBits(atomicAlignInBits);
56+
valueAlign = ctx.toCharUnitsFromBits(valueAlignInBits);
57+
if (lvalue.getAlignment().isZero())
58+
lvalue.setAlignment(atomicAlign);
59+
60+
this->lvalue = lvalue;
61+
} else {
62+
assert(!cir::MissingFeatures::atomicInfo());
63+
cgf.cgm.errorNYI(loc, "AtomicInfo: non-simple lvalue");
64+
}
65+
66+
assert(!cir::MissingFeatures::atomicUseLibCall());
67+
}
68+
69+
QualType getValueType() const { return valueTy; }
70+
CharUnits getAtomicAlignment() const { return atomicAlign; }
71+
TypeEvaluationKind getEvaluationKind() const { return evaluationKind; }
72+
mlir::Value getAtomicPointer() const {
73+
if (lvalue.isSimple())
74+
return lvalue.getPointer();
75+
assert(!cir::MissingFeatures::atomicInfoGetAtomicPointer());
76+
return nullptr;
77+
}
78+
Address getAtomicAddress() const {
79+
mlir::Type elemTy;
80+
if (lvalue.isSimple()) {
81+
elemTy = lvalue.getAddress().getElementType();
82+
} else {
83+
assert(!cir::MissingFeatures::atomicInfoGetAtomicAddress());
84+
cgf.cgm.errorNYI(loc, "AtomicInfo::getAtomicAddress: non-simple lvalue");
85+
}
86+
return Address(getAtomicPointer(), elemTy, getAtomicAlignment());
87+
}
88+
89+
/// Is the atomic size larger than the underlying value type?
90+
///
91+
/// Note that the absence of padding does not mean that atomic
92+
/// objects are completely interchangeable with non-atomic
93+
/// objects: we might have promoted the alignment of a type
94+
/// without making it bigger.
95+
bool hasPadding() const { return (valueSizeInBits != atomicSizeInBits); }
96+
97+
bool emitMemSetZeroIfNecessary() const;
98+
99+
/// Copy an atomic r-value into atomic-layout memory.
100+
void emitCopyIntoMemory(RValue rvalue) const;
101+
102+
/// Project an l-value down to the value field.
103+
LValue projectValue() const {
104+
assert(lvalue.isSimple());
105+
Address addr = getAtomicAddress();
106+
if (hasPadding()) {
107+
cgf.cgm.errorNYI(loc, "AtomicInfo::projectValue: padding");
108+
}
109+
110+
assert(!cir::MissingFeatures::opTBAA());
111+
return LValue::makeAddr(addr, getValueType(), lvalue.getBaseInfo());
112+
}
113+
114+
private:
115+
bool requiresMemSetZero(mlir::Type ty) const;
116+
};
117+
} // namespace
118+
119+
/// Does a store of the given IR type modify the full expected width?
120+
static bool isFullSizeType(CIRGenModule &cgm, mlir::Type ty,
121+
uint64_t expectedSize) {
122+
return cgm.getDataLayout().getTypeStoreSize(ty) * 8 == expectedSize;
123+
}
124+
125+
/// Does the atomic type require memsetting to zero before initialization?
126+
///
127+
/// The IR type is provided as a way of making certain queries faster.
128+
bool AtomicInfo::requiresMemSetZero(mlir::Type ty) const {
129+
// If the atomic type has size padding, we definitely need a memset.
130+
if (hasPadding())
131+
return true;
132+
133+
// Otherwise, do some simple heuristics to try to avoid it:
134+
switch (getEvaluationKind()) {
135+
// For scalars and complexes, check whether the store size of the
136+
// type uses the full size.
137+
case cir::TEK_Scalar:
138+
return !isFullSizeType(cgf.cgm, ty, atomicSizeInBits);
139+
case cir::TEK_Complex:
140+
cgf.cgm.errorNYI(loc, "AtomicInfo::requiresMemSetZero: complex type");
141+
return false;
142+
143+
// Padding in structs has an undefined bit pattern. User beware.
144+
case cir::TEK_Aggregate:
145+
return false;
146+
}
147+
llvm_unreachable("bad evaluation kind");
148+
}
149+
150+
bool AtomicInfo::emitMemSetZeroIfNecessary() const {
151+
assert(lvalue.isSimple());
152+
Address addr = lvalue.getAddress();
153+
if (!requiresMemSetZero(addr.getElementType()))
154+
return false;
155+
156+
cgf.cgm.errorNYI(loc,
157+
"AtomicInfo::emitMemSetZeroIfNecessary: emit memset zero");
158+
return false;
159+
}
160+
161+
/// Copy an r-value into memory as part of storing to an atomic type.
162+
/// This needs to create a bit-pattern suitable for atomic operations.
163+
void AtomicInfo::emitCopyIntoMemory(RValue rvalue) const {
164+
assert(lvalue.isSimple());
165+
166+
// If we have an r-value, the rvalue should be of the atomic type,
167+
// which means that the caller is responsible for having zeroed
168+
// any padding. Just do an aggregate copy of that type.
169+
if (rvalue.isAggregate()) {
170+
cgf.cgm.errorNYI("copying aggregate into atomic lvalue");
171+
return;
172+
}
173+
174+
// Okay, otherwise we're copying stuff.
175+
176+
// Zero out the buffer if necessary.
177+
emitMemSetZeroIfNecessary();
178+
179+
// Drill past the padding if present.
180+
LValue tempLValue = projectValue();
181+
182+
// Okay, store the rvalue in.
183+
if (rvalue.isScalar()) {
184+
cgf.emitStoreOfScalar(rvalue.getValue(), tempLValue, /*isInit=*/true);
185+
} else {
186+
cgf.cgm.errorNYI("copying complex into atomic lvalue");
187+
}
188+
}
189+
190+
RValue CIRGenFunction::emitAtomicExpr(AtomicExpr *e) {
191+
QualType atomicTy = e->getPtr()->getType()->getPointeeType();
192+
QualType memTy = atomicTy;
193+
if (const auto *ty = atomicTy->getAs<AtomicType>())
194+
memTy = ty->getValueType();
195+
196+
Address ptr = emitPointerWithAlignment(e->getPtr());
197+
198+
assert(!cir::MissingFeatures::openCL());
199+
if (e->getOp() == AtomicExpr::AO__c11_atomic_init) {
200+
LValue lvalue = makeAddrLValue(ptr, atomicTy);
201+
emitAtomicInit(e->getVal1(), lvalue);
202+
return RValue::get(nullptr);
203+
}
204+
205+
assert(!cir::MissingFeatures::atomicExpr());
206+
cgm.errorNYI(e->getSourceRange(), "atomic expr is NYI");
207+
return RValue::get(nullptr);
208+
}
209+
210+
void CIRGenFunction::emitAtomicInit(Expr *init, LValue dest) {
211+
AtomicInfo atomics(*this, dest, getLoc(init->getSourceRange()));
212+
213+
switch (atomics.getEvaluationKind()) {
214+
case cir::TEK_Scalar: {
215+
mlir::Value value = emitScalarExpr(init);
216+
atomics.emitCopyIntoMemory(RValue::get(value));
217+
return;
218+
}
219+
220+
case cir::TEK_Complex:
221+
cgm.errorNYI(init->getSourceRange(), "emitAtomicInit: complex type");
222+
return;
223+
224+
case cir::TEK_Aggregate:
225+
cgm.errorNYI(init->getSourceRange(), "emitAtomicInit: aggregate type");
226+
return;
227+
}
228+
229+
llvm_unreachable("bad evaluation kind");
230+
}

‎clang/lib/CIR/CodeGen/CIRGenCall.h

Lines changed: 21 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -46,6 +46,7 @@ class CIRGenCallee {
4646
enum class SpecialKind : uintptr_t {
4747
Invalid,
4848
Builtin,
49+
PseudoDestructor,
4950

5051
Last = Builtin,
5152
};
@@ -54,12 +55,16 @@ class CIRGenCallee {
5455
const clang::FunctionDecl *decl;
5556
unsigned id;
5657
};
58+
struct PseudoDestructorInfoStorage {
59+
const clang::CXXPseudoDestructorExpr *expr;
60+
};
5761

5862
SpecialKind kindOrFunctionPtr;
5963

6064
union {
6165
CIRGenCalleeInfo abstractInfo;
6266
BuiltinInfoStorage builtinInfo;
67+
PseudoDestructorInfoStorage pseudoDestructorInfo;
6368
};
6469

6570
explicit CIRGenCallee(SpecialKind kind) : kindOrFunctionPtr(kind) {}
@@ -98,6 +103,22 @@ class CIRGenCallee {
98103
return result;
99104
}
100105

106+
static CIRGenCallee
107+
forPseudoDestructor(const clang::CXXPseudoDestructorExpr *expr) {
108+
CIRGenCallee result(SpecialKind::PseudoDestructor);
109+
result.pseudoDestructorInfo.expr = expr;
110+
return result;
111+
}
112+
113+
bool isPseudoDestructor() const {
114+
return kindOrFunctionPtr == SpecialKind::PseudoDestructor;
115+
}
116+
117+
const CXXPseudoDestructorExpr *getPseudoDestructorExpr() const {
118+
assert(isPseudoDestructor());
119+
return pseudoDestructorInfo.expr;
120+
}
121+
101122
bool isOrdinary() const {
102123
return uintptr_t(kindOrFunctionPtr) > uintptr_t(SpecialKind::Last);
103124
}

‎clang/lib/CIR/CodeGen/CIRGenExpr.cpp

Lines changed: 9 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -184,8 +184,11 @@ Address CIRGenFunction::emitPointerWithAlignment(const Expr *expr,
184184
if (const UnaryOperator *uo = dyn_cast<UnaryOperator>(expr)) {
185185
// TODO(cir): maybe we should use cir.unary for pointers here instead.
186186
if (uo->getOpcode() == UO_AddrOf) {
187-
cgm.errorNYI(expr->getSourceRange(), "emitPointerWithAlignment: unary &");
188-
return Address::invalid();
187+
LValue lv = emitLValue(uo->getSubExpr());
188+
if (baseInfo)
189+
*baseInfo = lv.getBaseInfo();
190+
assert(!cir::MissingFeatures::opTBAA());
191+
return lv.getAddress();
189192
}
190193
}
191194

@@ -1544,10 +1547,10 @@ CIRGenCallee CIRGenFunction::emitCallee(const clang::Expr *e) {
15441547
cgm.errorNYI(e->getSourceRange(),
15451548
"emitCallee: call to member function is NYI");
15461549
return {};
1550+
} else if (auto *pde = dyn_cast<CXXPseudoDestructorExpr>(e)) {
1551+
return CIRGenCallee::forPseudoDestructor(pde);
15471552
}
15481553

1549-
assert(!cir::MissingFeatures::opCallPseudoDtor());
1550-
15511554
// Otherwise, we have an indirect reference.
15521555
mlir::Value calleePtr;
15531556
QualType functionType;
@@ -1599,10 +1602,8 @@ RValue CIRGenFunction::emitCallExpr(const clang::CallExpr *e,
15991602
return emitBuiltinExpr(callee.getBuiltinDecl(), callee.getBuiltinID(), e,
16001603
returnValue);
16011604

1602-
if (isa<CXXPseudoDestructorExpr>(e->getCallee())) {
1603-
cgm.errorNYI(e->getSourceRange(), "call to pseudo destructor");
1604-
}
1605-
assert(!cir::MissingFeatures::opCallPseudoDtor());
1605+
if (callee.isPseudoDestructor())
1606+
return emitCXXPseudoDestructorExpr(callee.getPseudoDestructorExpr());
16061607

16071608
return emitCall(e->getCallee()->getType(), callee, e, returnValue);
16081609
}

0 commit comments

Comments
 (0)