Skip to content

Commit

Permalink
Merge from 'sycl' to 'sycl-web'
Browse files Browse the repository at this point in the history
  • Loading branch information
iclsrc committed Oct 22, 2020
2 parents cbe0886 + 51d3c20 commit 5ab867b
Show file tree
Hide file tree
Showing 85 changed files with 2,240 additions and 1,236 deletions.
227 changes: 102 additions & 125 deletions clang/lib/Sema/SemaSYCL.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -3322,21 +3322,6 @@ static const char *paramKind2Str(KernelParamKind K) {
#undef CASE
}

// Removes all "(anonymous namespace)::" substrings from given string, and emits
// it.
static void emitWithoutAnonNamespaces(llvm::raw_ostream &OS, StringRef Source) {
const char S1[] = "(anonymous namespace)::";

size_t Pos;

while ((Pos = Source.find(S1)) != StringRef::npos) {
OS << Source.take_front(Pos);
Source = Source.drop_front(Pos + sizeof(S1) - 1);
}

OS << Source;
}

// Emits a forward declaration
void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D,
SourceLocation KernelLocation) {
Expand Down Expand Up @@ -3555,139 +3540,133 @@ void SYCLIntegrationHeader::emitForwardClassDecls(
}
}

static void emitCPPTypeString(raw_ostream &OS, QualType Ty) {
LangOptions LO;
PrintingPolicy P(LO);
P.SuppressTypedefs = true;
emitWithoutAnonNamespaces(OS, Ty.getAsString(P));
}

static void printArguments(ASTContext &Ctx, raw_ostream &ArgOS,
ArrayRef<TemplateArgument> Args,
const PrintingPolicy &P);
class SYCLKernelNameTypePrinter
: public TypeVisitor<SYCLKernelNameTypePrinter>,
public ConstTemplateArgumentVisitor<SYCLKernelNameTypePrinter> {
using InnerTypeVisitor = TypeVisitor<SYCLKernelNameTypePrinter>;
using InnerTemplArgVisitor =
ConstTemplateArgumentVisitor<SYCLKernelNameTypePrinter>;
raw_ostream &OS;
PrintingPolicy &Policy;

void printTemplateArgs(ArrayRef<TemplateArgument> Args) {
for (size_t I = 0, E = Args.size(); I < E; ++I) {
const TemplateArgument &Arg = Args[I];
// If argument is an empty pack argument, skip printing comma and
// argument.
if (Arg.getKind() == TemplateArgument::ArgKind::Pack && !Arg.pack_size())
continue;

static void emitKernelNameType(QualType T, ASTContext &Ctx, raw_ostream &OS,
const PrintingPolicy &TypePolicy);
if (I)
OS << ", ";

static void printArgument(ASTContext &Ctx, raw_ostream &ArgOS,
TemplateArgument Arg, const PrintingPolicy &P) {
switch (Arg.getKind()) {
case TemplateArgument::ArgKind::Pack: {
printArguments(Ctx, ArgOS, Arg.getPackAsArray(), P);
break;
}
case TemplateArgument::ArgKind::Integral: {
QualType T = Arg.getIntegralType();
const EnumType *ET = T->getAs<EnumType>();

if (ET) {
const llvm::APSInt &Val = Arg.getAsIntegral();
ArgOS << "static_cast<"
<< ET->getDecl()->getQualifiedNameAsString(
/*WithGlobalNsPrefix*/ true)
<< ">"
<< "(" << Val << ")";
} else {
Arg.print(P, ArgOS);
Visit(Arg);
}
break;
}
case TemplateArgument::ArgKind::Type: {
LangOptions LO;
PrintingPolicy TypePolicy(LO);
TypePolicy.SuppressTypedefs = true;
TypePolicy.SuppressTagKeyword = true;
QualType T = Arg.getAsType();

emitKernelNameType(T, Ctx, ArgOS, TypePolicy);
break;
}
case TemplateArgument::ArgKind::Template: {
TemplateDecl *TD = Arg.getAsTemplate().getAsTemplateDecl();
ArgOS << TD->getQualifiedNameAsString();
break;
}
default:
Arg.print(P, ArgOS);
void VisitQualifiers(Qualifiers Quals) {
Quals.print(OS, Policy, /*appendSpaceIfNotEmpty*/ true);
}
}

static void printArguments(ASTContext &Ctx, raw_ostream &ArgOS,
ArrayRef<TemplateArgument> Args,
const PrintingPolicy &P) {
for (unsigned I = 0; I < Args.size(); I++) {
const TemplateArgument &Arg = Args[I];
public:
SYCLKernelNameTypePrinter(raw_ostream &OS, PrintingPolicy &Policy)
: OS(OS), Policy(Policy) {}

// If argument is an empty pack argument, skip printing comma and argument.
if (Arg.getKind() == TemplateArgument::ArgKind::Pack && !Arg.pack_size())
continue;
void Visit(QualType T) {
if (T.isNull())
return;

if (I != 0)
ArgOS << ", ";
QualType CT = T.getCanonicalType();
VisitQualifiers(CT.getQualifiers());

printArgument(Ctx, ArgOS, Arg, P);
InnerTypeVisitor::Visit(CT.getTypePtr());
}
}

static void printTemplateArguments(ASTContext &Ctx, raw_ostream &ArgOS,
ArrayRef<TemplateArgument> Args,
const PrintingPolicy &P) {
ArgOS << "<";
printArguments(Ctx, ArgOS, Args, P);
ArgOS << ">";
}
void VisitType(const Type *T) {
OS << QualType::getAsString(T, Qualifiers(), Policy);
}

static void emitRecordType(raw_ostream &OS, QualType T, const CXXRecordDecl *RD,
const PrintingPolicy &TypePolicy) {
SmallString<64> Buf;
llvm::raw_svector_ostream RecOS(Buf);
T.getCanonicalType().getQualifiers().print(RecOS, TypePolicy,
/*appendSpaceIfNotEmpty*/ true);
if (const auto *TSD = dyn_cast<ClassTemplateSpecializationDecl>(RD)) {
void Visit(const TemplateArgument &TA) {
if (TA.isNull())
return;
InnerTemplArgVisitor::Visit(TA);
}

void VisitTagType(const TagType *T) {
TagDecl *RD = T->getDecl();
if (const auto *TSD = dyn_cast<ClassTemplateSpecializationDecl>(RD)) {

// Print template class name
TSD->printQualifiedName(RecOS, TypePolicy, /*WithGlobalNsPrefix*/ true);
// Print template class name
TSD->printQualifiedName(OS, Policy, /*WithGlobalNsPrefix*/ true);

// Print template arguments substituting enumerators
ASTContext &Ctx = RD->getASTContext();
const TemplateArgumentList &Args = TSD->getTemplateArgs();
printTemplateArguments(Ctx, RecOS, Args.asArray(), TypePolicy);
ArrayRef<TemplateArgument> Args = TSD->getTemplateArgs().asArray();
OS << "<";
printTemplateArgs(Args);
OS << ">";

emitWithoutAnonNamespaces(OS, RecOS.str());
return;
return;
}
// TODO: Next part of code results in printing of "class" keyword before
// class name in case if kernel name doesn't belong to some namespace. It
// seems if we don't print it, the integration header still represents valid
// c++ code. Probably we don't need to print it at all.
if (RD->getDeclContext()->isFunctionOrMethod()) {
OS << QualType::getAsString(T, Qualifiers(), Policy);
return;
}

const NamespaceDecl *NS = dyn_cast<NamespaceDecl>(RD->getDeclContext());
RD->printQualifiedName(OS, Policy, !(NS && NS->isAnonymousNamespace()));
}
if (RD->getDeclContext()->isFunctionOrMethod()) {
emitWithoutAnonNamespaces(OS, T.getCanonicalType().getAsString(TypePolicy));
return;

void VisitTemplateArgument(const TemplateArgument &TA) {
TA.print(Policy, OS);
}

const NamespaceDecl *NS = dyn_cast<NamespaceDecl>(RD->getDeclContext());
RD->printQualifiedName(RecOS, TypePolicy,
!(NS && NS->isAnonymousNamespace()));
emitWithoutAnonNamespaces(OS, RecOS.str());
}
void VisitTypeTemplateArgument(const TemplateArgument &TA) {
Policy.SuppressTagKeyword = true;
QualType T = TA.getAsType();
Visit(T);
Policy.SuppressTagKeyword = false;
}

static void emitKernelNameType(QualType T, ASTContext &Ctx, raw_ostream &OS,
const PrintingPolicy &TypePolicy) {
if (T->isRecordType()) {
emitRecordType(OS, T, T->getAsCXXRecordDecl(), TypePolicy);
return;
void VisitIntegralTemplateArgument(const TemplateArgument &TA) {
QualType T = TA.getIntegralType();
if (const EnumType *ET = T->getAs<EnumType>()) {
const llvm::APSInt &Val = TA.getAsIntegral();
OS << "static_cast<";
ET->getDecl()->printQualifiedName(OS, Policy,
/*WithGlobalNsPrefix*/ true);
OS << ">(" << Val << ")";
} else {
TA.print(Policy, OS);
}
}

if (T->isEnumeralType())
OS << "::";
emitWithoutAnonNamespaces(OS, T.getCanonicalType().getAsString(TypePolicy));
}
void VisitTemplateTemplateArgument(const TemplateArgument &TA) {
TemplateDecl *TD = TA.getAsTemplate().getAsTemplateDecl();
TD->printQualifiedName(OS, Policy);
}

void VisitPackTemplateArgument(const TemplateArgument &TA) {
printTemplateArgs(TA.getPackAsArray());
}
};

void SYCLIntegrationHeader::emit(raw_ostream &O) {
O << "// This is auto-generated SYCL integration header.\n";
O << "\n";

O << "#include <CL/sycl/detail/defines.hpp>\n";
O << "#include <CL/sycl/detail/defines_elementary.hpp>\n";
O << "#include <CL/sycl/detail/kernel_desc.hpp>\n";

O << "\n";

LangOptions LO;
PrintingPolicy Policy(LO);
Policy.SuppressTypedefs = true;
Policy.SuppressUnwrittenScope = true;

if (SpecConsts.size() > 0) {
// Remove duplicates.
std::sort(SpecConsts.begin(), SpecConsts.end(),
Expand All @@ -3705,7 +3684,7 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
O << "// Specialization constants IDs:\n";
for (const auto &P : llvm::make_range(SpecConsts.begin(), End)) {
O << "template <> struct sycl::detail::SpecConstantInfo<";
emitCPPTypeString(O, P.first);
O << P.first.getAsString(Policy);
O << "> {\n";
O << " static constexpr const char* getName() {\n";
O << " return \"" << P.second << "\";\n";
Expand Down Expand Up @@ -3773,19 +3752,17 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
O << "', '" << c;
O << "'> {\n";
} else {
LangOptions LO;
PrintingPolicy P(LO);
P.SuppressTypedefs = true;
O << "template <> struct KernelInfo<";
emitKernelNameType(K.NameType, S.getASTContext(), O, P);
SYCLKernelNameTypePrinter Printer(O, Policy);
Printer.Visit(K.NameType);
O << "> {\n";
}
O << " DLL_LOCAL\n";
O << " __SYCL_DLL_LOCAL\n";
O << " static constexpr const char* getName() { return \"" << K.Name
<< "\"; }\n";
O << " DLL_LOCAL\n";
O << " __SYCL_DLL_LOCAL\n";
O << " static constexpr unsigned getNumParams() { return " << N << "; }\n";
O << " DLL_LOCAL\n";
O << " __SYCL_DLL_LOCAL\n";
O << " static constexpr const kernel_param_desc_t& ";
O << "getParamDesc(unsigned i) {\n";
O << " return kernel_signatures[i+" << CurStart << "];\n";
Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/stdtypes_kernel_type.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -2,7 +2,7 @@
// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s
// RUN: FileCheck -input-file=%t.h %s
//
// CHECK: #include <CL/sycl/detail/defines.hpp>
// CHECK: #include <CL/sycl/detail/defines_elementary.hpp>
// CHECK-NEXT: #include <CL/sycl/detail/kernel_desc.hpp>
//
// CHECK: static constexpr
Expand Down
Loading

0 comments on commit 5ab867b

Please sign in to comment.