[Mlir-commits] [mlir] e7019f8 - [MLIR][IRDL] Instantiate "." in IRDL-defined names as nested namespaces (#207035)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Sat Sep 26 03:56:31 PDT 2026
Author: Ivan Ho
Date: 2026-09-26T18:56:24+08:00
New Revision: e7019f85692bab245c02f4fa433b37105f2d9c0d
URL: https://github.com/llvm/llvm-project/commit/e7019f85692bab245c02f4fa433b37105f2d9c0d
DIFF: https://github.com/llvm/llvm-project/commit/e7019f85692bab245c02f4fa433b37105f2d9c0d.diff
LOG: [MLIR][IRDL] Instantiate "." in IRDL-defined names as nested namespaces (#207035)
In PR #187911, it was revealed that the constraints on IRDL op names
placed by irdl-to-cpp were too restrictive.
This PR will address that by allowing "." in operation names to denote
nested namespaces, but not "$".
To discuss:
- ops like `wmma.f16.16x16x64.bf8_bf` cannot be registered as C++
identifiers cannot begin with a number. In discussions, we have opted to
prefix the generated namespaces with _.
Codex was used to generate the lit tests.
Close #187911.
Assisted-by: OpenAI Codex (GPT-5 and 6)
Added:
Modified:
mlir/lib/Dialect/IRDL/IR/IRDL.cpp
mlir/lib/Target/IRDLToCpp/IRDLToCpp.cpp
mlir/test/lib/Dialect/TestIRDLToCpp/test_irdl_to_cpp.irdl.mlir
Removed:
################################################################################
diff --git a/mlir/lib/Dialect/IRDL/IR/IRDL.cpp b/mlir/lib/Dialect/IRDL/IR/IRDL.cpp
index 278113fc6e966..10319687cd134 100644
--- a/mlir/lib/Dialect/IRDL/IR/IRDL.cpp
+++ b/mlir/lib/Dialect/IRDL/IR/IRDL.cpp
@@ -77,25 +77,30 @@ static llvm::LogicalResult isValidName(llvm::StringRef in, mlir::Operation *loc,
if (in.empty())
return loc->emitError("name of ") << label << " is empty";
- bool allowUnderscore = false;
- for (auto &elem : in) {
+ char prev = '_'; // treat the initial as _ to eliminate leading underscores
+ for (const auto &elem : in) {
if (elem == '_') {
- if (!allowUnderscore)
+ if (prev == '_')
return loc->emitError("name of ")
<< label << " should not contain leading or double underscores";
} else {
- if (!isalnum(elem))
- return loc->emitError("name of ")
- << label
- << " must contain only lowercase letters, digits and "
- "underscores";
-
- if (llvm::isUpper(elem))
- return loc->emitError("name of ")
- << label << " should not contain uppercase letters";
+ if (elem == '.') {
+ if (prev == '.')
+ return loc->emitError("empty namespace not allowed");
+ } else {
+ if (!isalnum(elem))
+ return loc->emitError("name of ")
+ << label
+ << " must contain only lowercase letters, digits and "
+ "underscores";
+
+ if (llvm::isUpper(elem))
+ return loc->emitError("name of ")
+ << label << " should not contain uppercase letters";
+ }
}
- allowUnderscore = elem != '_';
+ prev = elem;
}
return success();
diff --git a/mlir/lib/Target/IRDLToCpp/IRDLToCpp.cpp b/mlir/lib/Target/IRDLToCpp/IRDLToCpp.cpp
index 2bbd9de0ec51b..53d19fb5b8e80 100644
--- a/mlir/lib/Target/IRDLToCpp/IRDLToCpp.cpp
+++ b/mlir/lib/Target/IRDLToCpp/IRDLToCpp.cpp
@@ -13,6 +13,7 @@
#include "llvm/ADT/SmallString.h"
#include "llvm/ADT/SmallVector.h"
#include "llvm/ADT/StringExtras.h"
+#include "llvm/ADT/StringRef.h"
#include "llvm/ADT/TypeSwitch.h"
#include "llvm/Support/FormatVariadic.h"
#include "llvm/Support/raw_ostream.h"
@@ -52,6 +53,8 @@ struct TypeStrings {
struct OpStrings {
StringRef opName;
std::string opCppName;
+ std::string opScopedCppName;
+ SmallVector<std::string> opNameSpaces;
SmallVector<std::string> opResultNames;
SmallVector<std::string> opOperandNames;
SmallVector<std::string> opRegionNames;
@@ -65,16 +68,42 @@ static std::string joinNameList(llvm::ArrayRef<std::string> names) {
return nameArray;
}
+/// Prefix identifiers that start with a digit to make them valid C++ names.
+static std::string legalizeCppName(StringRef name) {
+ if (!name.empty() && llvm::isDigit(name.front()))
+ return ("_" + name).str();
+ return name.str();
+}
+
/// Generates the C++ type name for a TypeOp
static std::string typeToCppName(irdl::TypeOp type) {
- return llvm::formatv("{0}Type",
- convertToCamelFromSnakeCase(type.getSymName(), true));
+ return llvm::formatv("{0}Type", legalizeCppName(convertToCamelFromSnakeCase(
+ type.getSymName(), true)));
}
/// Generates the C++ class name for an OperationOp
static std::string opToCppName(irdl::OperationOp op) {
- return llvm::formatv("{0}Op",
- convertToCamelFromSnakeCase(op.getSymName(), true));
+ const auto opName = op.getSymName();
+ const auto periodIndex = opName.find_last_of(".");
+ const auto nameSubstr = periodIndex == std::string::npos
+ ? opName
+ : opName.substr(periodIndex + 1);
+ return llvm::formatv(
+ "{0}Op", legalizeCppName(convertToCamelFromSnakeCase(nameSubstr, true)));
+}
+
+/// Generates the C++ namespace components for an OperationOp.
+static SmallVector<std::string> opToCppNamespaces(irdl::OperationOp op) {
+ auto parts = SmallVector<StringRef>(llvm::split(op.getSymName(), "."));
+ parts.pop_back();
+ return llvm::map_to_vector(parts, legalizeCppName);
+}
+
+// Generates the C++ class name for an OperationOp, scoped to the namespace
+static std::string opToScopedCppName(irdl::OperationOp op) {
+ auto names = opToCppNamespaces(op);
+ names.push_back(opToCppName(op));
+ return llvm::join(names, "::");
}
/// Generates TypeStrings from a TypeOp
@@ -93,7 +122,9 @@ static OpStrings getStrings(irdl::OperationOp op) {
OpStrings strings;
strings.opName = op.getSymName();
+ strings.opNameSpaces = opToCppNamespaces(op);
strings.opCppName = opToCppName(op);
+ strings.opScopedCppName = opToScopedCppName(op);
if (operandOp) {
strings.opOperandNames = SmallVector<std::string>(
@@ -134,6 +165,7 @@ static void fillDict(irdl::detail::dictionary &dict, const OpStrings &strings) {
dict["OP_NAME"] = strings.opName;
dict["OP_CPP_NAME"] = strings.opCppName;
+ dict["OP_SCOPED_CPP_NAME"] = strings.opScopedCppName;
dict["OP_OPERAND_COUNT"] = std::to_string(strings.opOperandNames.size());
dict["OP_RESULT_COUNT"] = std::to_string(strings.opResultNames.size());
dict["OP_OPERAND_INITIALIZER_LIST"] =
@@ -141,6 +173,32 @@ static void fillDict(irdl::detail::dictionary &dict, const OpStrings &strings) {
dict["OP_RESULT_INITIALIZER_LIST"] =
resultCount ? joinNameList(strings.opResultNames) : "{\"\"}";
dict["OP_REGION_COUNT"] = std::to_string(regionCount);
+ dict["NAMESPACE_OPEN"] =
+ (dict["NAMESPACE_OPEN"] +
+ llvm::join(llvm::map_range(strings.opNameSpaces,
+ [](llvm::StringRef ref) -> std::string {
+ return llvm::formatv("namespace {0} {{",
+ ref);
+ }),
+ "\n"))
+ .str();
+ dict["NAMESPACE_PATH"] =
+ (dict["NAMESPACE_PATH"] +
+ llvm::join(llvm::map_range(strings.opNameSpaces,
+ [](llvm::StringRef ref) -> std::string {
+ return llvm::formatv("::{0}", ref);
+ }),
+ ""))
+ .str();
+ dict["NAMESPACE_CLOSE"] =
+ (llvm::join(llvm::map_range(llvm::reverse(strings.opNameSpaces),
+ [](llvm::StringRef ref) -> std::string {
+ return llvm::formatv("} // namespace {0}\n",
+ ref);
+ }),
+ "") +
+ dict["NAMESPACE_CLOSE"])
+ .str();
}
/// Fills a dictionary with values from DialectStrings
@@ -166,7 +224,7 @@ static LogicalResult generateTypedefList(irdl::DialectOp &dialect,
static LogicalResult generateOpList(irdl::DialectOp &dialect,
SmallVector<std::string> &opNames) {
auto operationOps = dialect.getOps<irdl::OperationOp>();
- auto range = llvm::map_range(operationOps, opToCppName);
+ auto range = llvm::map_range(operationOps, opToScopedCppName);
opNames = SmallVector<std::string>(range);
return success();
}
@@ -281,26 +339,27 @@ static SmallVector<std::string> generateTraits(irdl::OperationOp op,
return cppTraitNames;
}
-static LogicalResult generateOperationInclude(irdl::OperationOp op,
- raw_ostream &output,
- irdl::detail::dictionary &dict) {
+static LogicalResult
+generateOperationInclude(irdl::OperationOp op, raw_ostream &output,
+ const irdl::detail::dictionary &dict) {
static const auto perOpDeclTemplate = irdl::detail::Template(
#include "Templates/PerOperationDecl.txt"
);
const auto opStrings = getStrings(op);
- fillDict(dict, opStrings);
+ auto opDict = dict;
+ fillDict(opDict, opStrings);
SmallVector<std::string> traitNames = generateTraits(op, opStrings);
if (traitNames.empty())
- dict["OP_TEMPLATE_ARGS"] = opStrings.opCppName;
+ opDict["OP_TEMPLATE_ARGS"] = opStrings.opCppName;
else
- dict["OP_TEMPLATE_ARGS"] = llvm::formatv("{0}, {1}", opStrings.opCppName,
- llvm::join(traitNames, ", "));
+ opDict["OP_TEMPLATE_ARGS"] = llvm::formatv("{0}, {1}", opStrings.opCppName,
+ llvm::join(traitNames, ", "));
- generateOpGetterDeclarations(dict, opStrings);
- generateOpBuilderDeclarations(dict, opStrings);
+ generateOpGetterDeclarations(opDict, opStrings);
+ generateOpBuilderDeclarations(opDict, opStrings);
- perOpDeclTemplate.render(output, dict);
+ perOpDeclTemplate.render(output, opDict);
return success();
}
@@ -333,10 +392,16 @@ static LogicalResult generateInclude(irdl::DialectOp dialect,
return failure();
auto classDeclarations =
- llvm::join(llvm::map_range(opNames,
- [](llvm::StringRef name) -> std::string {
- return llvm::formatv("class {0};", name);
- }),
+ llvm::join(llvm::map_range(
+ opNames,
+ [](llvm::StringRef name) -> std::string {
+ if (name.contains("::")) {
+ auto [scope, className] = name.rsplit("::");
+ return llvm::formatv("namespace {0} {{\nclass {1};\n}",
+ scope, className);
+ }
+ return llvm::formatv("class {0};", name);
+ }),
"\n");
const auto forwardDeclarations = llvm::formatv(
"{1}\n{0}\n{2}", std::move(classDeclarations),
@@ -462,7 +527,8 @@ static std::string generateOpDefinition(irdl::detail::dictionary &dict,
};
auto opStrings = getStrings(op);
- fillDict(dict, opStrings);
+ auto opDict = dict;
+ fillDict(opDict, opStrings);
auto resultTypes = llvm::join(
llvm::map_range(opStrings.opResultNames,
@@ -516,13 +582,13 @@ void {0}::build(::mlir::OpBuilder &opBuilder, ::mlir::OperationState &opState, {
llvm::join(opStrings.opOperandNames, ",") +
(!opStrings.opOperandNames.empty() ? "," : ""));
- dict["OP_BUILD_DEFS"] = buildDefinition;
+ opDict["OP_BUILD_DEFS"] = buildDefinition;
- generateVerifiers(dict, op, opStrings);
+ generateVerifiers(opDict, op, opStrings);
std::string str;
llvm::raw_string_ostream stream{str};
- perOpDefTemplate.render(stream, dict);
+ perOpDefTemplate.render(stream, opDict);
return str;
}
diff --git a/mlir/test/lib/Dialect/TestIRDLToCpp/test_irdl_to_cpp.irdl.mlir b/mlir/test/lib/Dialect/TestIRDLToCpp/test_irdl_to_cpp.irdl.mlir
index 770cd36cdee33..35a973656d999 100644
--- a/mlir/test/lib/Dialect/TestIRDLToCpp/test_irdl_to_cpp.irdl.mlir
+++ b/mlir/test/lib/Dialect/TestIRDLToCpp/test_irdl_to_cpp.irdl.mlir
@@ -11,6 +11,26 @@ irdl.dialect @test_irdl_to_cpp {
// CHECK: class FooType
irdl.type @foo
+ // CHECK: class _8Type
+ irdl.type @"8"
+
+ // CHECK: class BarOp;
+ // CHECK-NEXT: namespace nested {
+ // CHECK-NEXT: class NamespacedOp;
+ // CHECK-NEXT: }
+ // CHECK-NEXT: namespace nested::namespaced {
+ // CHECK-NEXT: class MoreOp;
+ // CHECK-NEXT: }
+ // CHECK-NEXT: class BeefOp;
+
+ // CHECK: namespace wmma::f16::_16x16x64 {
+ // CHECK-NEXT: class Bf8BfOp;
+ // CHECK-NEXT: }
+ // CHECK-NEXT: namespace _7::_8 {
+ // CHECK-NEXT: class _9Op;
+ // CHECK-NEXT: }
+ // CHECK-NEXT: class _10Op;
+
// CHECK: class BarOp
// CHECK: ::mlir::Value getRes()
irdl.operation @bar {
@@ -18,6 +38,32 @@ irdl.dialect @test_irdl_to_cpp {
irdl.results(res: %0)
}
+ // CHECK: // ::mlir::test_irdl_to_cpp::nested::NamespacedOp declarations
+ // CHECK: namespace nested {
+ // CHECK: class NamespacedOp : public ::mlir::Op<NamespacedOp
+ // CHECK: static constexpr ::llvm::StringLiteral getOperationName()
+ // CHECK-NEXT: return ::llvm::StringLiteral("test_irdl_to_cpp.nested.namespaced");
+ // CHECK: } // namespace nested
+ // CHECK: MLIR_DECLARE_EXPLICIT_TYPE_ID(::mlir::test_irdl_to_cpp::nested::NamespacedOp)
+ irdl.operation @nested.namespaced {
+ %0 = irdl.any
+ irdl.results(res: %0)
+ }
+
+ // CHECK: // ::mlir::test_irdl_to_cpp::nested::namespaced::MoreOp declarations
+ // CHECK: namespace nested {
+ // CHECK: namespace namespaced {
+ // CHECK: class MoreOp : public ::mlir::Op<MoreOp
+ // CHECK: static constexpr ::llvm::StringLiteral getOperationName()
+ // CHECK-NEXT: return ::llvm::StringLiteral("test_irdl_to_cpp.nested.namespaced.more");
+ // CHECK: } // namespace namespaced
+ // CHECK: } // namespace nested
+ // CHECK: MLIR_DECLARE_EXPLICIT_TYPE_ID(::mlir::test_irdl_to_cpp::nested::namespaced::MoreOp)
+ irdl.operation @nested.namespaced.more {
+ %0 = irdl.any
+ irdl.results(res: %0)
+ }
+
// CHECK: class BeefOp
// CHECK: ::mlir::Value getLhs()
// CHECK: ::mlir::Value getRhs()
@@ -49,6 +95,56 @@ irdl.dialect @test_irdl_to_cpp {
// CHECK: ::mlir::Region &getThen() { return (*this)->getRegion(1); }
// CHECK: ::mlir::Region &getElse() { return (*this)->getRegion(2); }
+ irdl.operation @conditional {
+ %r0 = irdl.region // Unconstrained region
+ %r1 = irdl.region() // Region with no entry block arguments
+
+ // TODO(#161018): support irdl.is in irdl-to-cpp
+ // %v0 = irdl.is i1 // Type constraint: i1 (boolean)
+ %v0 = irdl.any
+ %r2 = irdl.region(%v0) // Region with one i1 entry block argument
+ irdl.regions(cond: %r2, then: %r0, else: %r1)
+
+ %0 = irdl.any
+ irdl.operands(input: %0)
+ }
+
+ // CHECK: // ::mlir::test_irdl_to_cpp::wmma::f16::_16x16x64::Bf8BfOp declarations
+ // CHECK: namespace wmma {
+ // CHECK-NEXT: namespace f16 {
+ // CHECK-NEXT: namespace _16x16x64 {
+ // CHECK: class Bf8BfOp : public ::mlir::Op<Bf8BfOp
+ // CHECK: return ::llvm::StringLiteral("test_irdl_to_cpp.wmma.f16.16x16x64.bf8_bf");
+ // CHECK: } // namespace _16x16x64
+ // CHECK: MLIR_DECLARE_EXPLICIT_TYPE_ID(::mlir::test_irdl_to_cpp::wmma::f16::_16x16x64::Bf8BfOp)
+ irdl.operation @wmma.f16.16x16x64.bf8_bf
+
+ // CHECK: // ::mlir::test_irdl_to_cpp::_7::_8::_9Op declarations
+ // CHECK: namespace _7 {
+ // CHECK-NEXT: namespace _8 {
+ // CHECK: class _9Op : public ::mlir::Op<_9Op
+ // CHECK: return ::llvm::StringLiteral("test_irdl_to_cpp.7.8.9");
+ // CHECK: MLIR_DECLARE_EXPLICIT_TYPE_ID(::mlir::test_irdl_to_cpp::_7::_8::_9Op)
+ irdl.operation @"7.8.9"
+
+ // CHECK: class _10Op : public ::mlir::Op<_10Op
+ // CHECK: return ::llvm::StringLiteral("test_irdl_to_cpp.10");
+ irdl.operation @"10"
+
+ // CHECK: // ::mlir::test_irdl_to_cpp::nested::NamespacedOp definitions
+ // CHECK: namespace nested {
+ // CHECK: NamespacedOp::build
+ // CHECK: } // namespace nested
+ // CHECK: MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::test_irdl_to_cpp::nested::NamespacedOp)
+
+ // CHECK: // ::mlir::test_irdl_to_cpp::nested::namespaced::MoreOp definitions
+ // CHECK: namespace nested {
+ // CHECK: namespace namespaced {
+ // CHECK: MoreOp::build
+ // CHECK: } // namespace namespaced
+ // CHECK: } // namespace nested
+ // CHECK: MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::test_irdl_to_cpp::nested::namespaced::MoreOp)
+
// CHECK: ConditionalOp definitions
// CHECK: __mlir_irdl_local_region_constraint_ConditionalOp_cond
// CHECK: if (!(region.getNumArguments() == 1)) {
@@ -73,17 +169,19 @@ irdl.dialect @test_irdl_to_cpp {
// CHECK: __mlir_irdl_local_region_constraint_ConditionalOp_else
// CHECK: failure
// CHECK: success
- irdl.operation @conditional {
- %r0 = irdl.region // Unconstrained region
- %r1 = irdl.region() // Region with no entry block arguments
-
- // TODO(#161018): support irdl.is in irdl-to-cpp
- // %v0 = irdl.is i1 // Type constraint: i1 (boolean)
- %v0 = irdl.any
- %r2 = irdl.region(%v0) // Region with one i1 entry block argument
- irdl.regions(cond: %r2, then: %r0, else: %r1)
- %0 = irdl.any
- irdl.operands(input: %0)
- }
+ // CHECK: // ::mlir::test_irdl_to_cpp::wmma::f16::_16x16x64::Bf8BfOp definitions
+ // CHECK: namespace _16x16x64 {
+ // CHECK: Bf8BfOp::build
+ // CHECK: } // namespace _16x16x64
+ // CHECK: MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::test_irdl_to_cpp::wmma::f16::_16x16x64::Bf8BfOp)
+ // CHECK: _9Op::build
+ // CHECK: MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::test_irdl_to_cpp::_7::_8::_9Op)
+ // CHECK: _10Op::build
+ // CHECK: void TestIrdlToCppDialect::initialize()
+ // CHECK: ::mlir::test_irdl_to_cpp::nested::NamespacedOp
+ // CHECK: ::mlir::test_irdl_to_cpp::nested::namespaced::MoreOp
+ // CHECK: ::mlir::test_irdl_to_cpp::wmma::f16::_16x16x64::Bf8BfOp,
+ // CHECK-NEXT: ::mlir::test_irdl_to_cpp::_7::_8::_9Op,
+ // CHECK-NEXT: ::mlir::test_irdl_to_cpp::_10Op
}
More information about the Mlir-commits
mailing list