Skip to content

Commit 63ff3ac

Browse files
committed
Merge from 'sycl' to 'sycl-web' (28 commits)
CONFLICT (content): Merge conflict in clang/lib/CodeGen/CGExprCXX.cpp
2 parents bd06f83 + 4a08a6c commit 63ff3ac

File tree

167 files changed

+3265
-3554
lines changed

Some content is hidden

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

167 files changed

+3265
-3554
lines changed

.github/workflows/sycl-update-gpu-driver.yml

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -30,5 +30,6 @@ jobs:
3030
git checkout -B $BRANCH
3131
git add -u
3232
git commit -m "[GHA] Uplift Linux GPU RT version to $NEW_DRIVER_VERSION" || exit 0 # exit if commit is empty
33+
git show
3334
git push https://[email protected]/${{ github.repository }} ${BRANCH}
3435
gh pr create --head $BRANCH --title "[GHA] Uplift Linux GPU RT version to $NEW_DRIVER_VERSION" --body "Scheduled drivers uplift"

clang/include/clang/Basic/DiagnosticSemaKinds.td

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -191,7 +191,7 @@ def err_intel_sycl_alloca_wrong_arg
191191
"'sycl::kernel_handler &'. Got %0">;
192192
def err_intel_sycl_alloca_wrong_type
193193
: Error<"__builtin_intel_sycl_alloca can only return 'sycl::private_ptr' "
194-
"to a cv-unqualified object type. Got %0">;
194+
"to a cv-unqualified trivial type. Got %0">;
195195
def err_intel_sycl_alloca_wrong_size
196196
: Error<"__builtin_intel_sycl_alloca must be passed a specialization "
197197
"constant of integral value type as a template argument. Got %1 (%0)">;

clang/include/clang/Basic/LangOptions.def

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -303,6 +303,7 @@ LANGOPT(
303303
ENUM_LANGOPT(SYCLRangeRounding, SYCLRangeRoundingPreference, 2,
304304
SYCLRangeRoundingPreference::On,
305305
"Preference for SYCL parallel_for range rounding")
306+
LANGOPT(SYCLExperimentalRangeRounding, 1, 0, "Use experimental parallel for range rounding")
306307
LANGOPT(SYCLEnableIntHeaderDiags, 1, 0, "Enable diagnostics that require the "
307308
"SYCL integration header")
308309
LANGOPT(SYCLAllowVirtualFunctions, 1, 0,

clang/include/clang/Driver/Options.td

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4014,6 +4014,10 @@ def fsycl_disable_range_rounding : Flag<["-"], "fsycl-disable-range-rounding">,
40144014
Alias<fsycl_range_rounding_EQ>, AliasArgs<["disable"]>,
40154015
HelpText<"Deprecated: please use -fsycl-range-rounding=disable instead.">,
40164016
Flags<[Deprecated]>;
4017+
def fsycl_exp_range_rounding : Flag<["-"], "fsycl-exp-range-rounding">,
4018+
Visibility<[ClangOption, CLOption, DXCOption, CC1Option]>,
4019+
HelpText<"Use experimental range rounding.">,
4020+
MarshallingInfoFlag<LangOpts<"SYCLExperimentalRangeRounding">>;
40174021
def fno_sycl_use_footer : Flag<["-"], "fno-sycl-use-footer">, Visibility<[ClangOption, CLOption, DXCOption]>,
40184022
HelpText<"Disable usage of the integration footer during SYCL enabled "
40194023
"compilations.">;

clang/lib/CodeGen/ABIInfoImpl.cpp

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -187,7 +187,7 @@ CodeGen::emitVoidPtrDirectVAArg(CodeGenFunction &CGF, Address VAListAddr,
187187
CharUnits FullDirectSize = DirectSize.alignTo(SlotSize);
188188
Address NextPtr =
189189
CGF.Builder.CreateConstInBoundsByteGEP(Addr, FullDirectSize, "argp.next");
190-
CGF.Builder.CreateStore(NextPtr.emitRawPointer(CGF), VAListAddr);
190+
CGF.Builder.CreateStore(NextPtr.getPointer(), VAListAddr);
191191

192192
// If the argument is smaller than a slot, and this is a big-endian
193193
// target, the argument will be right-adjusted in its slot.
@@ -239,8 +239,8 @@ Address CodeGen::emitMergePHI(CodeGenFunction &CGF, Address Addr1,
239239
const llvm::Twine &Name) {
240240
assert(Addr1.getType() == Addr2.getType());
241241
llvm::PHINode *PHI = CGF.Builder.CreatePHI(Addr1.getType(), 2, Name);
242-
PHI->addIncoming(Addr1.emitRawPointer(CGF), Block1);
243-
PHI->addIncoming(Addr2.emitRawPointer(CGF), Block2);
242+
PHI->addIncoming(Addr1.getPointer(), Block1);
243+
PHI->addIncoming(Addr2.getPointer(), Block2);
244244
CharUnits Align = std::min(Addr1.getAlignment(), Addr2.getAlignment());
245245
return Address(PHI, Addr1.getElementType(), Align);
246246
}
@@ -400,7 +400,7 @@ Address CodeGen::EmitVAArgInstr(CodeGenFunction &CGF, Address VAListAddr,
400400
llvm::Type *ElementTy = CGF.ConvertTypeForMem(Ty);
401401
llvm::Type *BaseTy = llvm::PointerType::getUnqual(ElementTy);
402402
llvm::Value *Addr =
403-
CGF.Builder.CreateVAArg(VAListAddr.emitRawPointer(CGF), BaseTy);
403+
CGF.Builder.CreateVAArg(VAListAddr.getPointer(), BaseTy);
404404
return Address(Addr, ElementTy, TyAlignForABI);
405405
} else {
406406
assert((AI.isDirect() || AI.isExtend()) &&
@@ -416,7 +416,7 @@ Address CodeGen::EmitVAArgInstr(CodeGenFunction &CGF, Address VAListAddr,
416416
"Unexpected CoerceToType seen in arginfo in generic VAArg emitter!");
417417

418418
Address Temp = CGF.CreateMemTemp(Ty, "varet");
419-
Val = CGF.Builder.CreateVAArg(VAListAddr.emitRawPointer(CGF),
419+
Val = CGF.Builder.CreateVAArg(VAListAddr.getPointer(),
420420
CGF.ConvertTypeForMem(Ty));
421421
CGF.Builder.CreateStore(Val, Temp);
422422
return Temp;

clang/lib/CodeGen/Address.h

Lines changed: 28 additions & 167 deletions
Original file line numberDiff line numberDiff line change
@@ -15,49 +15,35 @@
1515
#define LLVM_CLANG_LIB_CODEGEN_ADDRESS_H
1616

1717
#include "clang/AST/CharUnits.h"
18-
#include "clang/AST/Type.h"
1918
#include "llvm/ADT/PointerIntPair.h"
2019
#include "llvm/IR/Constants.h"
2120
#include "llvm/Support/MathExtras.h"
2221

2322
namespace clang {
2423
namespace CodeGen {
2524

26-
class Address;
27-
class CGBuilderTy;
28-
class CodeGenFunction;
29-
class CodeGenModule;
30-
3125
// Indicates whether a pointer is known not to be null.
3226
enum KnownNonNull_t { NotKnownNonNull, KnownNonNull };
3327

34-
/// An abstract representation of an aligned address. This is designed to be an
35-
/// IR-level abstraction, carrying just the information necessary to perform IR
36-
/// operations on an address like loads and stores. In particular, it doesn't
37-
/// carry C type information or allow the representation of things like
38-
/// bit-fields; clients working at that level should generally be using
39-
/// `LValue`.
40-
/// The pointer contained in this class is known to be unsigned.
41-
class RawAddress {
28+
/// An aligned address.
29+
class Address {
4230
llvm::PointerIntPair<llvm::Value *, 1, bool> PointerAndKnownNonNull;
4331
llvm::Type *ElementType;
4432
CharUnits Alignment;
4533

4634
protected:
47-
RawAddress(std::nullptr_t) : ElementType(nullptr) {}
35+
Address(std::nullptr_t) : ElementType(nullptr) {}
4836

4937
public:
50-
RawAddress(llvm::Value *Pointer, llvm::Type *ElementType, CharUnits Alignment,
51-
KnownNonNull_t IsKnownNonNull = NotKnownNonNull)
38+
Address(llvm::Value *Pointer, llvm::Type *ElementType, CharUnits Alignment,
39+
KnownNonNull_t IsKnownNonNull = NotKnownNonNull)
5240
: PointerAndKnownNonNull(Pointer, IsKnownNonNull),
5341
ElementType(ElementType), Alignment(Alignment) {
5442
assert(Pointer != nullptr && "Pointer cannot be null");
5543
assert(ElementType != nullptr && "Element type cannot be null");
5644
}
5745

58-
inline RawAddress(Address Addr);
59-
60-
static RawAddress invalid() { return RawAddress(nullptr); }
46+
static Address invalid() { return Address(nullptr); }
6147
bool isValid() const {
6248
return PointerAndKnownNonNull.getPointer() != nullptr;
6349
}
@@ -94,133 +80,6 @@ class RawAddress {
9480
return Alignment;
9581
}
9682

97-
/// Return address with different element type, but same pointer and
98-
/// alignment.
99-
RawAddress withElementType(llvm::Type *ElemTy) const {
100-
return RawAddress(getPointer(), ElemTy, getAlignment(), isKnownNonNull());
101-
}
102-
103-
KnownNonNull_t isKnownNonNull() const {
104-
assert(isValid());
105-
return (KnownNonNull_t)PointerAndKnownNonNull.getInt();
106-
}
107-
};
108-
109-
/// Like RawAddress, an abstract representation of an aligned address, but the
110-
/// pointer contained in this class is possibly signed.
111-
class Address {
112-
friend class CGBuilderTy;
113-
114-
// The boolean flag indicates whether the pointer is known to be non-null.
115-
llvm::PointerIntPair<llvm::Value *, 1, bool> Pointer;
116-
117-
/// The expected IR type of the pointer. Carrying accurate element type
118-
/// information in Address makes it more convenient to work with Address
119-
/// values and allows frontend assertions to catch simple mistakes.
120-
llvm::Type *ElementType = nullptr;
121-
122-
CharUnits Alignment;
123-
124-
/// Offset from the base pointer.
125-
llvm::Value *Offset = nullptr;
126-
127-
llvm::Value *emitRawPointerSlow(CodeGenFunction &CGF) const;
128-
129-
protected:
130-
Address(std::nullptr_t) : ElementType(nullptr) {}
131-
132-
public:
133-
Address(llvm::Value *pointer, llvm::Type *elementType, CharUnits alignment,
134-
KnownNonNull_t IsKnownNonNull = NotKnownNonNull)
135-
: Pointer(pointer, IsKnownNonNull), ElementType(elementType),
136-
Alignment(alignment) {
137-
assert(pointer != nullptr && "Pointer cannot be null");
138-
assert(elementType != nullptr && "Element type cannot be null");
139-
assert(!alignment.isZero() && "Alignment cannot be zero");
140-
}
141-
142-
Address(llvm::Value *BasePtr, llvm::Type *ElementType, CharUnits Alignment,
143-
llvm::Value *Offset, KnownNonNull_t IsKnownNonNull = NotKnownNonNull)
144-
: Pointer(BasePtr, IsKnownNonNull), ElementType(ElementType),
145-
Alignment(Alignment), Offset(Offset) {}
146-
147-
Address(RawAddress RawAddr)
148-
: Pointer(RawAddr.isValid() ? RawAddr.getPointer() : nullptr),
149-
ElementType(RawAddr.isValid() ? RawAddr.getElementType() : nullptr),
150-
Alignment(RawAddr.isValid() ? RawAddr.getAlignment()
151-
: CharUnits::Zero()) {}
152-
153-
static Address invalid() { return Address(nullptr); }
154-
bool isValid() const { return Pointer.getPointer() != nullptr; }
155-
156-
/// This function is used in situations where the caller is doing some sort of
157-
/// opaque "laundering" of the pointer.
158-
void replaceBasePointer(llvm::Value *P) {
159-
assert(isValid() && "pointer isn't valid");
160-
assert(P->getType() == Pointer.getPointer()->getType() &&
161-
"Pointer's type changed");
162-
Pointer.setPointer(P);
163-
assert(isValid() && "pointer is invalid after replacement");
164-
}
165-
166-
CharUnits getAlignment() const { return Alignment; }
167-
168-
void setAlignment(CharUnits Value) { Alignment = Value; }
169-
170-
llvm::Value *getBasePointer() const {
171-
assert(isValid() && "pointer isn't valid");
172-
return Pointer.getPointer();
173-
}
174-
175-
/// Return the type of the pointer value.
176-
llvm::PointerType *getType() const {
177-
return llvm::PointerType::get(
178-
ElementType,
179-
llvm::cast<llvm::PointerType>(Pointer.getPointer()->getType())
180-
->getAddressSpace());
181-
}
182-
183-
/// Return the type of the values stored in this address.
184-
llvm::Type *getElementType() const {
185-
assert(isValid());
186-
return ElementType;
187-
}
188-
189-
/// Return the address space that this address resides in.
190-
unsigned getAddressSpace() const { return getType()->getAddressSpace(); }
191-
192-
/// Return the IR name of the pointer value.
193-
llvm::StringRef getName() const { return Pointer.getPointer()->getName(); }
194-
195-
// This function is called only in CGBuilderBaseTy::CreateElementBitCast.
196-
void setElementType(llvm::Type *Ty) {
197-
assert(hasOffset() &&
198-
"this funcion shouldn't be called when there is no offset");
199-
ElementType = Ty;
200-
}
201-
202-
/// Whether the pointer is known not to be null.
203-
KnownNonNull_t isKnownNonNull() const {
204-
assert(isValid());
205-
return (KnownNonNull_t)Pointer.getInt();
206-
}
207-
208-
Address setKnownNonNull() {
209-
assert(isValid());
210-
Pointer.setInt(KnownNonNull);
211-
return *this;
212-
}
213-
214-
bool hasOffset() const { return Offset; }
215-
216-
llvm::Value *getOffset() const { return Offset; }
217-
218-
/// Return the pointer contained in this class after authenticating it and
219-
/// adding offset to it if necessary.
220-
llvm::Value *emitRawPointer(CodeGenFunction &CGF) const {
221-
return getBasePointer();
222-
}
223-
22483
/// Return address with different pointer, but same element type and
22584
/// alignment.
22685
Address withPointer(llvm::Value *NewPointer,
@@ -232,59 +91,61 @@ class Address {
23291
/// Return address with different alignment, but same pointer and element
23392
/// type.
23493
Address withAlignment(CharUnits NewAlignment) const {
235-
return Address(Pointer.getPointer(), getElementType(), NewAlignment,
94+
return Address(getPointer(), getElementType(), NewAlignment,
23695
isKnownNonNull());
23796
}
23897

23998
/// Return address with different element type, but same pointer and
24099
/// alignment.
241100
Address withElementType(llvm::Type *ElemTy) const {
242-
if (!hasOffset())
243-
return Address(getBasePointer(), ElemTy, getAlignment(), nullptr,
244-
isKnownNonNull());
245-
Address A(*this);
246-
A.ElementType = ElemTy;
247-
return A;
101+
return Address(getPointer(), ElemTy, getAlignment(), isKnownNonNull());
248102
}
249-
};
250103

251-
inline RawAddress::RawAddress(Address Addr)
252-
: PointerAndKnownNonNull(Addr.isValid() ? Addr.getBasePointer() : nullptr,
253-
Addr.isValid() ? Addr.isKnownNonNull()
254-
: NotKnownNonNull),
255-
ElementType(Addr.isValid() ? Addr.getElementType() : nullptr),
256-
Alignment(Addr.isValid() ? Addr.getAlignment() : CharUnits::Zero()) {}
104+
/// Whether the pointer is known not to be null.
105+
KnownNonNull_t isKnownNonNull() const {
106+
assert(isValid());
107+
return (KnownNonNull_t)PointerAndKnownNonNull.getInt();
108+
}
109+
110+
/// Set the non-null bit.
111+
Address setKnownNonNull() {
112+
assert(isValid());
113+
PointerAndKnownNonNull.setInt(true);
114+
return *this;
115+
}
116+
};
257117

258118
/// A specialization of Address that requires the address to be an
259119
/// LLVM Constant.
260-
class ConstantAddress : public RawAddress {
261-
ConstantAddress(std::nullptr_t) : RawAddress(nullptr) {}
120+
class ConstantAddress : public Address {
121+
ConstantAddress(std::nullptr_t) : Address(nullptr) {}
262122

263123
public:
264124
ConstantAddress(llvm::Constant *pointer, llvm::Type *elementType,
265125
CharUnits alignment)
266-
: RawAddress(pointer, elementType, alignment) {}
126+
: Address(pointer, elementType, alignment) {}
267127

268128
static ConstantAddress invalid() {
269129
return ConstantAddress(nullptr);
270130
}
271131

272132
llvm::Constant *getPointer() const {
273-
return llvm::cast<llvm::Constant>(RawAddress::getPointer());
133+
return llvm::cast<llvm::Constant>(Address::getPointer());
274134
}
275135

276136
ConstantAddress withElementType(llvm::Type *ElemTy) const {
277137
return ConstantAddress(getPointer(), ElemTy, getAlignment());
278138
}
279139

280-
static bool isaImpl(RawAddress addr) {
140+
static bool isaImpl(Address addr) {
281141
return llvm::isa<llvm::Constant>(addr.getPointer());
282142
}
283-
static ConstantAddress castImpl(RawAddress addr) {
143+
static ConstantAddress castImpl(Address addr) {
284144
return ConstantAddress(llvm::cast<llvm::Constant>(addr.getPointer()),
285145
addr.getElementType(), addr.getAlignment());
286146
}
287147
};
148+
288149
}
289150

290151
// Present a minimal LLVM-like casting interface.

clang/lib/CodeGen/BackendUtil.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1138,10 +1138,10 @@ void EmitAssemblyHelper::RunOptimizationPipeline(
11381138

11391139
// Add SPIRITTAnnotations pass to the pass manager if
11401140
// -fsycl-instrument-device-code option was passed. This option can be
1141-
// used only with spir triple.
1141+
// used only with spir or spirv triple.
11421142
if (CodeGenOpts.SPIRITTAnnotations) {
11431143
assert(
1144-
TargetTriple.isSPIR() &&
1144+
TargetTriple.isSPIROrSPIRV() &&
11451145
"ITT annotations can only be added to a module with spir target");
11461146
MPM.addPass(SPIRITTAnnotationsPass());
11471147
}

0 commit comments

Comments
 (0)