Skip to content

Commit 1094c85

Browse files
author
iclsrc
committed
Merge from 'sycl' to 'sycl-web' (#9)
2 parents a34c9f0 + ee37fce commit 1094c85

File tree

28 files changed

+422
-144
lines changed

28 files changed

+422
-144
lines changed

‎clang/lib/Basic/Targets/NVPTX.h

+3-2
Original file line numberDiff line numberDiff line change
@@ -46,8 +46,8 @@ static const int NVPTXDWARFAddrSpaceMap[] = {
4646
-1, // Default, opencl_private or opencl_generic - not defined
4747
5, // opencl_global
4848
-1,
49-
8, // opencl_local or cuda_shared
50-
4, // opencl_constant or cuda_constant
49+
8, // opencl_local or cuda_shared
50+
4, // opencl_constant or cuda_constant
5151
};
5252

5353
class LLVM_LIBRARY_VISIBILITY NVPTXTargetInfo : public TargetInfo {
@@ -144,6 +144,7 @@ class LLVM_LIBRARY_VISIBILITY NVPTXTargetInfo : public TargetInfo {
144144
Opts.support("cl_khr_int64_base_atomics");
145145
Opts.support("cl_khr_int64_extended_atomics");
146146
Opts.support("cl_khr_fp16");
147+
Opts.support("cl_khr_3d_image_writes");
147148
}
148149

149150
/// \returns If a target requires an address within a target specific address

‎clang/lib/Driver/ToolChains/Clang.cpp

+1-1
Original file line numberDiff line numberDiff line change
@@ -7901,5 +7901,5 @@ void PartialLink::ConstructJob(Compilation &C, const JobAction &JA,
79017901
LinkArgs.push_back(TCArgs.MakeArgString(HTC->GetFilePath("crtn.o")));
79027902
const char *Exec = TCArgs.MakeArgString(getToolChain().GetLinkerPath());
79037903
C.addCommand(std::make_unique<Command>(
7904-
JA, *this, ResponseFileSupport::None(), Exec, LinkArgs, Inputs));
7904+
JA, *this, ResponseFileSupport::AtFileCurCP(), Exec, LinkArgs, Inputs));
79057905
}

‎clang/lib/Driver/ToolChains/Cuda.h

+1-1
Original file line numberDiff line numberDiff line change
@@ -203,7 +203,7 @@ class LLVM_LIBRARY_VISIBILITY CudaToolChain : public ToolChain {
203203

204204
unsigned GetDefaultDwarfVersion() const override { return 2; }
205205

206-
Tool *SelectTool(const JobAction &JA) const;
206+
Tool *SelectTool(const JobAction &JA) const override;
207207

208208
const ToolChain &HostTC;
209209
CudaInstallationDetector CudaInstallation;

‎clang/lib/Driver/ToolChains/SYCL.cpp

+3-3
Original file line numberDiff line numberDiff line change
@@ -51,7 +51,7 @@ const char *SYCL::Linker::constructLLVMSpirvCommand(Compilation &C,
5151
llvm::sys::path::append(LLVMSpirvPath, "llvm-spirv");
5252
const char *LLVMSpirv = C.getArgs().MakeArgString(LLVMSpirvPath);
5353
C.addCommand(std::make_unique<Command>(
54-
JA, *this, ResponseFileSupport::None(), LLVMSpirv, CmdArgs, None));
54+
JA, *this, ResponseFileSupport::AtFileUTF8(), LLVMSpirv, CmdArgs, None));
5555
return OutputFileName;
5656
}
5757

@@ -128,7 +128,7 @@ const char *SYCL::Linker::constructLLVMLinkCommand(Compilation &C,
128128
llvm::sys::path::append(ExecPath, "llvm-link");
129129
const char *Exec = C.getArgs().MakeArgString(ExecPath);
130130
C.addCommand(std::make_unique<Command>(
131-
JA, *this, ResponseFileSupport::None(), Exec, CmdArgs, None));
131+
JA, *this, ResponseFileSupport::AtFileUTF8(), Exec, CmdArgs, None));
132132
return OutputFileName;
133133
}
134134

@@ -142,7 +142,7 @@ void SYCL::Linker::constructLlcCommand(Compilation &C, const JobAction &JA,
142142
llvm::sys::path::append(LlcPath, "llc");
143143
const char *Llc = C.getArgs().MakeArgString(LlcPath);
144144
C.addCommand(std::make_unique<Command>(
145-
JA, *this, ResponseFileSupport::None(), Llc, LlcArgs, None));
145+
JA, *this, ResponseFileSupport::AtFileUTF8(), Llc, LlcArgs, None));
146146
}
147147

148148
// For SYCL the inputs of the linker job are SPIR-V binaries and output is

‎clang/lib/Sema/SemaSYCL.cpp

+57-50
Original file line numberDiff line numberDiff line change
@@ -935,7 +935,7 @@ void KernelObjVisitor::VisitRecord(CXXRecordDecl *Owner, ParentTy &Parent,
935935

936936
// A base type that the SYCL OpenCL Kernel construction task uses to implement
937937
// individual tasks.
938-
template <typename Derived> class SyclKernelFieldHandler {
938+
class SyclKernelFieldHandler {
939939
protected:
940940
Sema &SemaRef;
941941
SyclKernelFieldHandler(Sema &S) : SemaRef(S) {}
@@ -1001,11 +1001,12 @@ template <typename Derived> class SyclKernelFieldHandler {
10011001
virtual bool enterArray() { return true; }
10021002
virtual bool nextElement(QualType) { return true; }
10031003
virtual bool leaveArray(FieldDecl *, QualType, int64_t) { return true; }
1004+
1005+
virtual ~SyclKernelFieldHandler() = default;
10041006
};
10051007

10061008
// A type to check the validity of all of the argument types.
1007-
class SyclKernelFieldChecker
1008-
: public SyclKernelFieldHandler<SyclKernelFieldChecker> {
1009+
class SyclKernelFieldChecker : public SyclKernelFieldHandler {
10091010
bool IsInvalid = false;
10101011
DiagnosticsEngine &Diag;
10111012

@@ -1101,8 +1102,7 @@ class SyclKernelFieldChecker
11011102
};
11021103

11031104
// A type to Create and own the FunctionDecl for the kernel.
1104-
class SyclKernelDeclCreator
1105-
: public SyclKernelFieldHandler<SyclKernelDeclCreator> {
1105+
class SyclKernelDeclCreator : public SyclKernelFieldHandler {
11061106
FunctionDecl *KernelDecl;
11071107
llvm::SmallVector<ParmVarDecl *, 8> Params;
11081108
SyclKernelFieldChecker &ArgChecker;
@@ -1286,8 +1286,7 @@ class SyclKernelDeclCreator
12861286
using SyclKernelFieldHandler::handleSyclSamplerType;
12871287
};
12881288

1289-
class SyclKernelBodyCreator
1290-
: public SyclKernelFieldHandler<SyclKernelBodyCreator> {
1289+
class SyclKernelBodyCreator : public SyclKernelFieldHandler {
12911290
SyclKernelDeclCreator &DeclCreator;
12921291
llvm::SmallVector<Stmt *, 16> BodyStmts;
12931292
llvm::SmallVector<Stmt *, 16> FinalizeStmts;
@@ -1696,12 +1695,8 @@ class SyclKernelBodyCreator
16961695
using SyclKernelFieldHandler::leaveStruct;
16971696
};
16981697

1699-
class SyclKernelIntHeaderCreator
1700-
: public SyclKernelFieldHandler<SyclKernelIntHeaderCreator> {
1698+
class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
17011699
SYCLIntegrationHeader &Header;
1702-
const CXXRecordDecl *KernelObj;
1703-
// Necessary to figure out the offset of the base class.
1704-
const CXXRecordDecl *CurStruct = nullptr;
17051700
int64_t CurOffset = 0;
17061701

17071702
void addParam(const FieldDecl *FD, QualType ArgTy,
@@ -1720,7 +1715,7 @@ class SyclKernelIntHeaderCreator
17201715
SyclKernelIntHeaderCreator(Sema &S, SYCLIntegrationHeader &H,
17211716
const CXXRecordDecl *KernelObj, QualType NameType,
17221717
StringRef Name, StringRef StableName)
1723-
: SyclKernelFieldHandler(S), Header(H), KernelObj(KernelObj) {
1718+
: SyclKernelFieldHandler(S), Header(H) {
17241719
Header.startKernel(Name, NameType, StableName, KernelObj->getLocation());
17251720
}
17261721

@@ -2120,13 +2115,19 @@ static const char *paramKind2Str(KernelParamKind K) {
21202115
#undef CASE
21212116
}
21222117

2123-
// Removes all "(anonymous namespace)::" substrings from given string
2124-
static std::string eraseAnonNamespace(std::string S) {
2118+
// Removes all "(anonymous namespace)::" substrings from given string, and emits
2119+
// it.
2120+
static void emitWithoutAnonNamespaces(llvm::raw_ostream &OS, StringRef Source) {
21252121
const char S1[] = "(anonymous namespace)::";
21262122

2127-
for (auto Pos = S.find(S1); Pos != StringRef::npos; Pos = S.find(S1, Pos))
2128-
S.erase(Pos, sizeof(S1) - 1);
2129-
return S;
2123+
size_t Pos;
2124+
2125+
while ((Pos = Source.find(S1)) != StringRef::npos) {
2126+
OS << Source.take_front(Pos);
2127+
Source = Source.drop_front(Pos + sizeof(S1) - 1);
2128+
}
2129+
2130+
OS << Source;
21302131
}
21312132

21322133
static bool checkEnumTemplateParameter(const EnumDecl *ED,
@@ -2375,19 +2376,19 @@ void SYCLIntegrationHeader::emitForwardClassDecls(
23752376
}
23762377
}
23772378

2378-
static std::string getCPPTypeString(QualType Ty) {
2379+
static void emitCPPTypeString(raw_ostream &OS, QualType Ty) {
23792380
LangOptions LO;
23802381
PrintingPolicy P(LO);
23812382
P.SuppressTypedefs = true;
2382-
return eraseAnonNamespace(Ty.getAsString(P));
2383+
emitWithoutAnonNamespaces(OS, Ty.getAsString(P));
23832384
}
23842385

23852386
static void printArguments(ASTContext &Ctx, raw_ostream &ArgOS,
23862387
ArrayRef<TemplateArgument> Args,
23872388
const PrintingPolicy &P);
23882389

2389-
static std::string getKernelNameTypeString(QualType T, ASTContext &Ctx,
2390-
const PrintingPolicy &TypePolicy);
2390+
static void emitKernelNameType(QualType T, ASTContext &Ctx, raw_ostream &OS,
2391+
const PrintingPolicy &TypePolicy);
23912392

23922393
static void printArgument(ASTContext &Ctx, raw_ostream &ArgOS,
23932394
TemplateArgument Arg, const PrintingPolicy &P) {
@@ -2418,7 +2419,8 @@ static void printArgument(ASTContext &Ctx, raw_ostream &ArgOS,
24182419
TypePolicy.SuppressTypedefs = true;
24192420
TypePolicy.SuppressTagKeyword = true;
24202421
QualType T = Arg.getAsType();
2421-
ArgOS << getKernelNameTypeString(T, Ctx, TypePolicy);
2422+
2423+
emitKernelNameType(T, Ctx, ArgOS, TypePolicy);
24222424
break;
24232425
}
24242426
case TemplateArgument::ArgKind::Template: {
@@ -2456,42 +2458,46 @@ static void printTemplateArguments(ASTContext &Ctx, raw_ostream &ArgOS,
24562458
ArgOS << ">";
24572459
}
24582460

2459-
static std::string printRecordType(QualType T, const CXXRecordDecl *RD,
2460-
const PrintingPolicy &TypePolicy) {
2461+
static void emitRecordType(raw_ostream &OS, QualType T, const CXXRecordDecl *RD,
2462+
const PrintingPolicy &TypePolicy) {
24612463
SmallString<64> Buf;
2462-
llvm::raw_svector_ostream OS(Buf);
2463-
T.getCanonicalType().getQualifiers().print(OS, TypePolicy,
2464+
llvm::raw_svector_ostream RecOS(Buf);
2465+
T.getCanonicalType().getQualifiers().print(RecOS, TypePolicy,
24642466
/*appendSpaceIfNotEmpty*/ true);
24652467
if (const auto *TSD = dyn_cast<ClassTemplateSpecializationDecl>(RD)) {
24662468

24672469
// Print template class name
2468-
TSD->printQualifiedName(OS, TypePolicy, /*WithGlobalNsPrefix*/ true);
2470+
TSD->printQualifiedName(RecOS, TypePolicy, /*WithGlobalNsPrefix*/ true);
24692471

24702472
// Print template arguments substituting enumerators
24712473
ASTContext &Ctx = RD->getASTContext();
24722474
const TemplateArgumentList &Args = TSD->getTemplateArgs();
2473-
printTemplateArguments(Ctx, OS, Args.asArray(), TypePolicy);
2475+
printTemplateArguments(Ctx, RecOS, Args.asArray(), TypePolicy);
24742476

2475-
return eraseAnonNamespace(OS.str().str());
2477+
emitWithoutAnonNamespaces(OS, RecOS.str());
2478+
return;
24762479
}
2477-
if (RD->getDeclContext()->isFunctionOrMethod())
2478-
return eraseAnonNamespace(T.getCanonicalType().getAsString(TypePolicy));
2480+
if (RD->getDeclContext()->isFunctionOrMethod()) {
2481+
emitWithoutAnonNamespaces(OS, T.getCanonicalType().getAsString(TypePolicy));
2482+
return;
2483+
}
2484+
24792485
const NamespaceDecl *NS = dyn_cast<NamespaceDecl>(RD->getDeclContext());
2480-
RD->printQualifiedName(OS, TypePolicy, !(NS && NS->isAnonymousNamespace()));
2481-
return eraseAnonNamespace(OS.str().str());
2486+
RD->printQualifiedName(RecOS, TypePolicy,
2487+
!(NS && NS->isAnonymousNamespace()));
2488+
emitWithoutAnonNamespaces(OS, RecOS.str());
24822489
}
24832490

2484-
static std::string getKernelNameTypeString(QualType T, ASTContext &Ctx,
2485-
const PrintingPolicy &TypePolicy) {
2486-
if (T->isRecordType())
2487-
return printRecordType(T, T->getAsCXXRecordDecl(), TypePolicy);
2488-
if (T->isEnumeralType()) {
2489-
SmallString<64> Buf;
2490-
llvm::raw_svector_ostream OS(Buf);
2491-
OS << "::" << T.getCanonicalType().getAsString(TypePolicy);
2492-
return eraseAnonNamespace(OS.str().str());
2493-
}
2494-
return eraseAnonNamespace(T.getCanonicalType().getAsString(TypePolicy));
2491+
static void emitKernelNameType(QualType T, ASTContext &Ctx, raw_ostream &OS,
2492+
const PrintingPolicy &TypePolicy) {
2493+
if (T->isRecordType()) {
2494+
emitRecordType(OS, T, T->getAsCXXRecordDecl(), TypePolicy);
2495+
return;
2496+
}
2497+
2498+
if (T->isEnumeralType())
2499+
OS << "::";
2500+
emitWithoutAnonNamespaces(OS, T.getCanonicalType().getAsString(TypePolicy));
24952501
}
24962502

24972503
void SYCLIntegrationHeader::emit(raw_ostream &O) {
@@ -2519,9 +2525,9 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
25192525
});
25202526
O << "// Specialization constants IDs:\n";
25212527
for (const auto &P : llvm::make_range(SpecConsts.begin(), End)) {
2522-
std::string CPPName = getCPPTypeString(P.first);
2523-
O << "template <> struct sycl::detail::SpecConstantInfo<" << CPPName
2524-
<< "> {\n";
2528+
O << "template <> struct sycl::detail::SpecConstantInfo<";
2529+
emitCPPTypeString(O, P.first);
2530+
O << "> {\n";
25252531
O << " static constexpr const char* getName() {\n";
25262532
O << " return \"" << P.second << "\";\n";
25272533
O << " }\n";
@@ -2613,8 +2619,9 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
26132619
LangOptions LO;
26142620
PrintingPolicy P(LO);
26152621
P.SuppressTypedefs = true;
2616-
O << "template <> struct KernelInfo<"
2617-
<< getKernelNameTypeString(K.NameType, S.getASTContext(), P) << "> {\n";
2622+
O << "template <> struct KernelInfo<";
2623+
emitKernelNameType(K.NameType, S.getASTContext(), O, P);
2624+
O << "> {\n";
26182625
}
26192626
O << " DLL_LOCAL\n";
26202627
O << " static constexpr const char* getName() { return \"" << K.Name

0 commit comments

Comments
 (0)