Skip to content

[StrTable] Mechanically convert NVPTX builtins to use TableGen #122873

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

Merged
merged 1 commit into from
Jan 28, 2025
Merged

[StrTable] Mechanically convert NVPTX builtins to use TableGen #122873

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
1,119 changes: 0 additions & 1,119 deletions clang/include/clang/Basic/BuiltinsNVPTX.def
View file
Open in desktop

This file was deleted.

1,078 changes: 1,078 additions & 0 deletions clang/include/clang/Basic/BuiltinsNVPTX.td
View file
Open in desktop

Large diffs are not rendered by default.

4 changes: 4 additions & 0 deletions clang/include/clang/Basic/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -72,6 +72,10 @@ clang_tablegen(BuiltinsBPF.inc -gen-clang-builtins
SOURCE BuiltinsBPF.td
TARGET ClangBuiltinsBPF)

clang_tablegen(BuiltinsNVPTX.inc -gen-clang-builtins
SOURCE BuiltinsNVPTX.td
TARGET ClangBuiltinsNVPTX)

clang_tablegen(BuiltinsRISCV.inc -gen-clang-builtins
SOURCE BuiltinsRISCV.td
TARGET ClangBuiltinsRISCV)
Expand Down
10 changes: 5 additions & 5 deletions clang/include/clang/Basic/TargetBuiltins.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -101,12 +101,12 @@ namespace clang {

/// NVPTX builtins
namespace NVPTX {
enum {
LastTIBuiltin = clang::Builtin::FirstTSBuiltin-1,
enum {
LastTIBuiltin = clang::Builtin::FirstTSBuiltin - 1,
#define BUILTIN(ID, TYPE, ATTRS) BI##ID,
#include "clang/Basic/BuiltinsNVPTX.def"
LastTSBuiltin
};
#include "clang/Basic/BuiltinsNVPTX.inc"
LastTSBuiltin
};
}

/// AMDGPU builtins
Expand Down
1 change: 0 additions & 1 deletion clang/include/module.modulemap
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -53,7 +53,6 @@ module Clang_Basic {
textual header "clang/Basic/BuiltinsLoongArchLSX.def"
textual header "clang/Basic/BuiltinsMips.def"
textual header "clang/Basic/BuiltinsNEON.def"
textual header "clang/Basic/BuiltinsNVPTX.def"
textual header "clang/Basic/BuiltinsPPC.def"
textual header "clang/Basic/BuiltinsRISCVVector.def"
textual header "clang/Basic/BuiltinsSME.def"
Expand Down
6 changes: 1 addition & 5 deletions clang/lib/Basic/Targets/NVPTX.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,13 +21,9 @@ using namespace clang;
using namespace clang::targets;

static constexpr Builtin::Info BuiltinInfo[] = {
#define BUILTIN(ID, TYPE, ATTRS) \
{#ID, TYPE, ATTRS, nullptr, HeaderDesc::NO_HEADER, ALL_LANGUAGES},
#define LIBBUILTIN(ID, TYPE, ATTRS, HEADER) \
{#ID, TYPE, ATTRS, nullptr, HeaderDesc::HEADER, ALL_LANGUAGES},
#define TARGET_BUILTIN(ID, TYPE, ATTRS, FEATURE) \
{#ID, TYPE, ATTRS, FEATURE, HeaderDesc::NO_HEADER, ALL_LANGUAGES},
#include "clang/Basic/BuiltinsNVPTX.def"
#include "clang/Basic/BuiltinsNVPTX.inc"
};

const char *const NVPTXTargetInfo::GCCRegNames[] = {"r0"};
Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGen/builtins-nvptx.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -202,7 +202,7 @@ __device__ void exit() {
// NVVM intrinsics

// The idea is not to test all intrinsics, just that Clang is recognizing the
// builtins defined in BuiltinsNVPTX.def
// builtins defined in BuiltinsNVPTX.td
__device__ void nvvm_math(float f1, float f2, double d1, double d2) {
// CHECK: call float @llvm.nvvm.fmax.f
float t1 = __nvvm_fmax_f(f1, f2);
Expand Down
37 changes: 37 additions & 0 deletions clang/utils/TableGen/ClangBuiltinsEmitter.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -104,9 +104,39 @@ class PrototypeParser {

void ParseType(StringRef T) {
T = T.trim();

auto ConsumeAddrSpace = [&]() -> std::optional<unsigned> {
T = T.trim();
if (!T.consume_back(">"))
return std::nullopt;

auto Open = T.find_last_of('<');
if (Open == StringRef::npos)
PrintFatalError(Loc, "Mismatched angle-brackets in type");

StringRef ArgStr = T.substr(Open + 1);
T = T.slice(0, Open);
if (!T.consume_back("address_space"))
PrintFatalError(Loc,
"Only `address_space<N>` supported as a parameterized "
"pointer or reference type qualifier");

unsigned Number = 0;
if (ArgStr.getAsInteger(10, Number))
PrintFatalError(
Loc, "Expected an integer argument to the address_space qualifier");
if (Number == 0)
PrintFatalError(Loc, "No need for a qualifier for address space `0`");
return Number;
};

if (T.consume_back("*")) {
// Pointers may have an address space qualifier immediately before them.
std::optional<unsigned> AS = ConsumeAddrSpace();
Copy link
Member

Choose a reason for hiding this comment

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

Do we have any tests for the tablegen parser changes?

Copy link
Member Author

Choose a reason for hiding this comment

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

I'm not an expert at any of this... I'm just muddling my way through to try and make some tactical improvements.

There is a file target-builtins-prototype-parser.td, but it doesn't test many aspects of this code... Given that it didn't seem to try to be comprehensive, and that NVPTX provided some pretty reasonable tests, I just focused there. That's also what I did with the analogous change to x86. 🤷

Let me know if you'd like me to add more tests to cover this.

Copy link
Member

Choose a reason for hiding this comment

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

I'm not an expert at any of this... I'm just muddling my way through to try and make some tactical improvements.

Welcome to the club. This matches my own thoughts. :-)

I agree that it would not improve effective test coverage much, but it would be useful for whoever may need to tinker with these parts of tablegen in the future. Having wider range of inputs elsewhere is great, but it is way less convenient to work with, considering that it tends to be much larger and has additional irrelevant moving parts.

ParseType(T);
Type += "*";
if (AS)
Type += std::to_string(*AS);
} else if (T.consume_back("const")) {
ParseType(T);
Type += "C";
Expand All @@ -117,6 +147,13 @@ class PrototypeParser {
ParseType(T);
Type += "R";
} else if (T.consume_back("&")) {
// References may have an address space qualifier immediately before them.
std::optional<unsigned> AS = ConsumeAddrSpace();
ParseType(T);
Type += "&";
if (AS)
Type += std::to_string(*AS);
} else if (T.consume_back(")")) {
ParseType(T);
Type += "&";
} else if (EnableOpenCLLong && T.consume_front("long long")) {
Expand Down
Loading