Skip to content

[SYCL-PTX] Add a JSON backend for the ProgModel builtin #1781

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Closed
wants to merge 9 commits into from
Closed

[SYCL-PTX] Add a JSON backend for the ProgModel builtin #1781

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
9 commits
Select commit Hold shift + click to select a range
02b7975
[SYCL-PTX] Add a JSON backend for the ProgModel builtin
Naghasan May 18, 2020
8a9f2de
[SYCL-PTX] Review comment: handle type description logic using a flag…
Naghasan Jun 1, 2020
720e62d
[SYCL-PTX] Coding style fixes
Naghasan Jun 1, 2020
14d3dbd
[SYCL-PTX] Review comment: moving code around
Naghasan Jun 1, 2020
695b984
[SYCL-PTX] clang-format
Naghasan Jun 1, 2020
414d9f2
[SYCL-PTX] Apply review comment: pedentic changes regarding auto
Naghasan Jun 2, 2020
90ba58a
[SYCL-PTX] add assert for type
Naghasan Jun 2, 2020
f16c4f6
[SYCL-PTX] Apply review comment: some more pedantic changes
Naghasan Jun 3, 2020
75dfbac
[SYCL-PTX] Apply review comment: make the unscoped enum UnderlyingTyp…
Naghasan Jun 3, 2020
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
5 changes: 5 additions & 0 deletions clang/lib/Sema/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -23,6 +23,11 @@ clang_tablegen(SPIRVBuiltins.inc -gen-clang-spirv-builtins
TARGET ClangSPIRVBuiltinsImpl
)

clang_tablegen(SPIRVBuiltins.json --gen-clang-progmodel-builtins-as-json
SOURCE SPIRVBuiltins.td
TARGET ClangSPIRVBuiltinsJSON
)

add_clang_library(clangSema
AnalysisBasedWarnings.cpp
CodeCompleteConsumer.cpp
Expand Down
303 changes: 302 additions & 1 deletion clang/utils/TableGen/ClangProgModelBuiltinEmitter.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -63,18 +63,26 @@
#include "llvm/ADT/StringRef.h"
#include "llvm/ADT/StringSet.h"
#include "llvm/ADT/StringSwitch.h"
#include "llvm/Support/CommandLine.h"
#include "llvm/Support/ErrorHandling.h"
#include "llvm/Support/raw_ostream.h"
#include "llvm/TableGen/Error.h"
#include "llvm/TableGen/Record.h"
#include "llvm/TableGen/StringMatcher.h"
#include "llvm/TableGen/TableGenBackend.h"
#include <numeric>
#include <set>

using namespace llvm;

namespace {

cl::opt<std::string>
TypePrefix("json-type-prefix",
cl::desc("Use a specific prefix for the types of the json prog "
"model builtin output. Default: __clc."),
cl::value_desc("prefix"), cl::init("__clc"), cl::Hidden);

// A list of signatures that are shared by one or more builtin functions.
struct BuiltinTableEntries {
SmallVector<StringRef, 4> Names;
Expand All @@ -92,7 +100,7 @@ class BuiltinNameEmitter {
// whether a function is a builtin function.
void Emit();

private:
protected:
// A list of indices into the builtin function table.
using BuiltinIndexListTy = SmallVector<unsigned, 11>;

Expand Down Expand Up @@ -237,6 +245,66 @@ class BuiltinNameEmitter {
// same entry (<I1, I2, I3>).
MapVector<BuiltinIndexListTy *, BuiltinTableEntries> SignatureListMap;
};

/// Generate a json file representing the builtins as understood by clang.
/// This allows the generation of a library interface.
class JSONBuiltinInterfaceEmitter : public BuiltinNameEmitter {
struct TypeDesc {
TypeDesc(const Record *T);
// Return a string representing the "base" type (ignores pointers).
std::string GetBaseTypeAsStr() const;

enum class UnderlyingType { BOOL, INT, UINT, FLOAT, BUILTIN };

llvm::StringRef Name;
// Size of the vector (if applicable).
int VecWidth;
// Size of the element in bits.
int ElementSize;
// The underlying type represented by the record.
UnderlyingType ElementType;
// Is a pointer.
bool IsPointer;
// "const" qualifier.
bool IsConst;
// "volatile" qualifier.
bool IsVolatile;
llvm::StringRef AddrSpace;
};

public:
JSONBuiltinInterfaceEmitter(RecordKeeper &Records, raw_ostream &OS,
llvm::StringRef Family)
: BuiltinNameEmitter(Records, OS, Family) {}

// Entrypoint to generate the functions and structures for checking
// whether a function is an builtin function.
void Emit();

private:
void EmitBuiltins();

// Expand all generic types into a list of types.
void ExpandTypes();
void ExpandType(const Record *Ty);

// For each Type records, map a vector of string representing a all types
// encoded in the record.
// Simple type will result in a 1 element vector, but generic type (e.g. all
// ints) will result in a n elements vector encoding concrete types (e.g. int8
// uint8, int16 etc.)
// The format use is:
// __clc_vec{len}_{type}_t
// where
// "_vec{len}" is added if the type is a vector.
// "{type}" is a normalized type:
// int{size}, uint{size}, fp{size}, size (for size_t), event, bool
// if the type is a pointer, "{const} {volatile} {asp} *" is appended with
// {asp}: <empty>, __private, __global, __local, __constant
//
// __clc can be changed using the CLI option TypePrefix
DenseMap<const Record *, llvm::SmallVector<std::string, 16>> ExpandedTypes;
};
} // namespace

void BuiltinNameEmitter::Emit() {
Expand Down Expand Up @@ -848,6 +916,233 @@ void BuiltinNameEmitter::EmitQualTypeFinder() {
OS << "\n} // Bultin2Qual\n";
}

JSONBuiltinInterfaceEmitter::TypeDesc::TypeDesc(const Record *T)
: Name(T->getValueAsString("Name")), VecWidth(T->getValueAsInt("VecWidth")),
ElementSize(T->getValueAsInt("ElementSize")),
IsPointer(T->getValueAsBit("IsPointer")),
IsConst(T->getValueAsBit("IsConst")),
IsVolatile(T->getValueAsBit("IsVolatile")) {
assert((!T->getValueAsBit("IsInteger") || !T->getValueAsBit("IsFloat")) &&
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Can you please add the isSubClassOf checks as well? It seems that it shouldn't be a BUILTIN either.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It is already handled by the unreachable.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I don't think it is. I'm talking about a type that IsInteger, IsFloat, and isSubClassOf(FundamentalType) AND isSubClassOf(VectorType).

At the moment you catch only isInteger and isFloat, but if you leave any of those off it isnt.

Additionally, the logic below allows a type to be both a fundamental and vector type.

"A type can't be both an int and a float");
if (T->getValueAsBit("IsInteger")) {
if (ElementSize == 1)
ElementType = UnderlyingType::BOOL;
else
ElementType = T->getValueAsBit("IsSigned") ? UnderlyingType::INT
: UnderlyingType::UINT;
} else {
if (T->getValueAsBit("IsFloat"))
ElementType = UnderlyingType::FLOAT;
else {
if (!T->isSubClassOf("FundamentalType") && !T->isSubClassOf("VectorType"))
ElementType = UnderlyingType::BUILTIN;
else
llvm_unreachable("Invalid type");
}
}
AddrSpace = StringSwitch<StringRef>(T->getValueAsString("AddrSpace"))
.Case("clang::LangAS::Default", "")
.Case("clang::LangAS::opencl_private", " __private")
.Case("clang::LangAS::opencl_global", " __global")
.Case("clang::LangAS::opencl_constant", " __constant")
.Case("clang::LangAS::opencl_local", " __local")
.Default(" __generic");
}

std::string JSONBuiltinInterfaceEmitter::TypeDesc::GetBaseTypeAsStr() const {
if (Name == "void")
return Name.str();

llvm::SmallString<32> Buffer;
{
llvm::raw_svector_ostream TypeStr(Buffer);
TypeStr << TypePrefix;
if (VecWidth != 1)
TypeStr << "_vec" << VecWidth;
switch (ElementType) {
case UnderlyingType::BOOL:
TypeStr << "_bool_t";
break;
case UnderlyingType::INT:
TypeStr << "_int" << ElementSize << "_t";
break;
case UnderlyingType::UINT:
TypeStr << "_uint" << ElementSize << "_t";
break;
case UnderlyingType::FLOAT:
TypeStr << "_fp" << ElementSize << "_t";
break;
case UnderlyingType::BUILTIN:
TypeStr << "_" << Name;
break;
}
}

return static_cast<std::string>(Buffer);
}

void JSONBuiltinInterfaceEmitter::Emit() {

// Parse the Records to populate the internal lists.
GetOverloads();
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I suspect calling these directly is against the original author's abstraction (which is why they were private in the first place). How do other emitters handle this?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

There is no other emitter

GroupBySignature();
ExpandTypes();

EmitBuiltins();
}

// Generate all possible types represented by the input record.
void JSONBuiltinInterfaceEmitter::ExpandType(const Record *Ty) {
auto ExpansionIt = ExpandedTypes.find(Ty);

if (ExpansionIt != ExpandedTypes.end())
return;
llvm::SmallVectorImpl<std::string> &Expansion = ExpandedTypes[Ty];

JSONBuiltinInterfaceEmitter::TypeDesc TypeDesc(Ty);

bool IsCompound = Ty->isSubClassOf("CompoundType");
if (IsCompound) {
const Record *EltRecord = Ty->getValueAsDef("ElementType");
ExpandType(EltRecord);
Expansion = ExpandedTypes[EltRecord];
for (std::string &TypeStr : Expansion) {
if (TypeDesc.IsPointer) {
if (TypeDesc.IsConst)
TypeStr += " const";
if (TypeDesc.IsVolatile)
TypeStr += " volatile";
TypeStr += TypeDesc.AddrSpace;
TypeStr += " *";
}
}

return;
}

if (Ty->isSubClassOf("GenericType")) {
assert(!Expansion.size() && "Types already expanded");
std::vector<long int> NVec =
Ty->getValueAsDef("VectorList")->getValueAsListOfInts("List");
Comment on lines +1025 to +1026
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Sorry for the delay.
Out internal pre-commit check failed to build this patch on Windows with the following message:

clang\utils\TableGen\ClangProgModelBuiltinEmitter.cpp(1026): error C2440: 'initializing': cannot convert from 'std::vector<int64_t,std::allocator<_Ty>>' to 'std::vector<long,std::allocator<_Ty>>'
         with
         [
             _Ty=int64_t
         ]
         and
         [
             _Ty=long
         ]
clang\utils\TableGen\ClangProgModelBuiltinEmitter.cpp(1026): note: No constructor could take the source type, or constructor overload resolution was ambiguous

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

On Windows, "long" is 32 bit, so int64_t is long long. This likely should be std::vector<int64_t> instead.

std::vector<Record *> NTypes =
Ty->getValueAsDef("TypeList")->getValueAsListOfDefs("List");
for (const Record *SubT : NTypes) {
for (long int VecWidth : NVec) {
JSONBuiltinInterfaceEmitter::TypeDesc TypeDesc(SubT);
TypeDesc.VecWidth = VecWidth;
Expansion.emplace_back(TypeDesc.GetBaseTypeAsStr());
}
}

return;
}

Expansion.emplace_back(TypeDesc.GetBaseTypeAsStr());
}

void JSONBuiltinInterfaceEmitter::ExpandTypes() {
for (const Record *T : Records.getAllDerivedDefinitions("Type")) {
ExpandType(T);
}
}

void JSONBuiltinInterfaceEmitter::EmitBuiltins() {
struct FnDesc {
FnDesc(const Record *Rec)
: IsPure(Rec->getValueAsBit("IsPure")),
IsConst(Rec->getValueAsBit("IsConst")),
IsConv(Rec->getValueAsBit("IsConv")),
IsVariadic(Rec->getValueAsBit("IsVariadic")) {}

bool IsPure;
bool IsConst;
bool IsConv;
bool IsVariadic;
SmallVector<llvm::StringRef, 8> Args;
};
StringMap<SmallVector<FnDesc, 16>> NameToProtoList;

// For each function names, gather the list of overloads
for (const std::pair<BuiltinIndexListTy *, BuiltinTableEntries> &SLM :
SignatureListMap) {
for (StringRef Name : SLM.second.Names) {
SmallVectorImpl<FnDesc> &PrototypeList = NameToProtoList[Name];

for (const std::pair<const Record *, unsigned> &Overload :
SLM.second.Signatures) {
std::vector<Record *> Signature =
Overload.first->getValueAsListOfDefs("Signature");
auto SignatureTypesIt = llvm::map_range(
Signature, [this](const Record *Ty) -> llvm::ArrayRef<std::string> {
return ExpandedTypes[Ty];
});
auto SignatureLenIt = llvm::map_range(
SignatureTypesIt, [](llvm::ArrayRef<std::string> ExpandedType) {
return ExpandedType.size();
});
assert(SignatureLenIt.begin() != SignatureLenIt.end());
// Find out how many overloads this signature holds.
size_t NbOverload = 1;
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Suggested change
size_t NbOverload = 1;
size_t NbOverload = *std::max_element(std::begin(SignatureLenIt), std::end(SignatureLenIt));

for (size_t Len : SignatureLenIt)
NbOverload = std::max(NbOverload, Len);
// assert the signature is consistent
assert(llvm::all_of(SignatureLenIt, [NbOverload](size_t NbTypes) {
return NbTypes == NbOverload || NbTypes == 1;
}));
for (size_t idx = 0; idx < NbOverload; idx++) {
FnDesc Proto(Overload.first);
for (llvm::ArrayRef<std::string> Types : SignatureTypesIt) {
Proto.Args.emplace_back(Types[idx >= Types.size() ? 0 : idx]);
}
PrototypeList.emplace_back(std::move(Proto));
}
}
}
}
OS << "{\n";
// emit name : ["ty1", "ty2", "ty3"], ["pure"]
llvm::interleave(
NameToProtoList.keys(), OS,
[&](llvm::StringRef FnName) {
OS << " \"" << FnName << "\" : [\n";
SmallVectorImpl<FnDesc> &PrototypeList = NameToProtoList[FnName];
// emit ["ty1", "ty2", "ty3"], ["pure"]
llvm::interleave(
PrototypeList, OS,
[&](const FnDesc &Prototype) {
OS << " [[";
// emit "ty1", "ty2", "ty3"
llvm::interleave(
Prototype.Args, OS,
[&](llvm::StringRef Type) { OS << "\"" << Type << "\""; },
",");
OS << "], [";
bool HasPrev = false;
// Helper, skip the comma for the first attribute to comply with
// JSON.
auto EmitWithPrefix = [&](llvm::StringRef Attr) {
if (HasPrev)
OS << ", ";
OS << "\"" << Attr << "\"";
HasPrev = true;
};
if (Prototype.IsPure)
EmitWithPrefix("pure");
if (Prototype.IsConst)
EmitWithPrefix("const");
if (Prototype.IsConv)
EmitWithPrefix("convergent");
if (Prototype.IsVariadic)
EmitWithPrefix("variadic");
OS << "]]";
},
",\n");
OS << "\n ]";
},
",\n");
OS << "\n}\n";
}

void clang::EmitClangOpenCLBuiltins(RecordKeeper &Records, raw_ostream &OS) {
BuiltinNameEmitter NameChecker(Records, OS, "OpenCL");
NameChecker.Emit();
Expand All @@ -857,3 +1152,9 @@ void clang::EmitClangSPIRVBuiltins(RecordKeeper &Records, raw_ostream &OS) {
BuiltinNameEmitter NameChecker(Records, OS, "SPIRV");
NameChecker.Emit();
}

void clang::EmitProgModelBuiltinsAsJSON(RecordKeeper &Records,
raw_ostream &OS) {
JSONBuiltinInterfaceEmitter NameChecker(Records, OS, "");
NameChecker.Emit();
}
8 changes: 8 additions & 0 deletions clang/utils/TableGen/TableGen.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -62,6 +62,7 @@ enum ActionType {
GenClangCommentCommandList,
GenClangOpenCLBuiltins,
GenClangSPIRVBuiltins,
GenClangProgModelBuiltinsAsJSON,
GenArmNeon,
GenArmFP16,
GenArmNeonSema,
Expand Down Expand Up @@ -187,6 +188,10 @@ cl::opt<ActionType> Action(
"Generate OpenCL builtin declaration handlers"),
clEnumValN(GenClangSPIRVBuiltins, "gen-clang-spirv-builtins",
"Generate SPIR-V builtin declaration handlers"),
clEnumValN(
GenClangProgModelBuiltinsAsJSON,
"gen-clang-progmodel-builtins-as-json",
"Generate programming model builtins declaration as a JSON file"),
clEnumValN(GenArmNeon, "gen-arm-neon", "Generate arm_neon.h for clang"),
clEnumValN(GenArmFP16, "gen-arm-fp16", "Generate arm_fp16.h for clang"),
clEnumValN(GenArmNeonSema, "gen-arm-neon-sema",
Expand Down Expand Up @@ -360,6 +365,9 @@ bool ClangTableGenMain(raw_ostream &OS, RecordKeeper &Records) {
case GenClangSPIRVBuiltins:
EmitClangSPIRVBuiltins(Records, OS);
break;
case GenClangProgModelBuiltinsAsJSON:
EmitProgModelBuiltinsAsJSON(Records, OS);
break;
case GenArmNeon:
EmitNeon(Records, OS);
break;
Expand Down
2 changes: 2 additions & 0 deletions clang/utils/TableGen/TableGenBackends.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -116,6 +116,8 @@ void EmitClangOptDocs(llvm::RecordKeeper &Records, llvm::raw_ostream &OS);
void EmitClangOpenCLBuiltins(llvm::RecordKeeper &Records,
llvm::raw_ostream &OS);
void EmitClangSPIRVBuiltins(llvm::RecordKeeper &Records, llvm::raw_ostream &OS);
void EmitProgModelBuiltinsAsJSON(llvm::RecordKeeper &Records,
llvm::raw_ostream &OS);

void EmitClangDataCollectors(llvm::RecordKeeper &Records,
llvm::raw_ostream &OS);
Expand Down
Loading