From bf92affa0a6c50cb8973cb3060f4fcf4bc6b6ad6 Mon Sep 17 00:00:00 2001 From: max Date: Wed, 1 Jan 2025 21:54:55 -0500 Subject: [PATCH 1/6] [eudsl] factor out eudsl-nbgen --- projects/eudsl-nbgen/CMakeLists.txt | 89 ++++ projects/eudsl-nbgen/eudsl-nbgen.cpp | 722 +++++++++++++++++++++++++++ 2 files changed, 811 insertions(+) create mode 100644 projects/eudsl-nbgen/CMakeLists.txt create mode 100644 projects/eudsl-nbgen/eudsl-nbgen.cpp diff --git a/projects/eudsl-nbgen/CMakeLists.txt b/projects/eudsl-nbgen/CMakeLists.txt new file mode 100644 index 00000000..72fa12cb --- /dev/null +++ b/projects/eudsl-nbgen/CMakeLists.txt @@ -0,0 +1,89 @@ +cmake_minimum_required(VERSION 3.29) +set(CMAKE_CXX_STANDARD 17) +set(CMAKE_EXPORT_COMPILE_COMMANDS ON) + +set(LLVM_SUBPROJECT_TITLE "EUDSLPY") +set(EUDSLPY_BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR}) + +if (CMAKE_SOURCE_DIR STREQUAL CMAKE_CURRENT_SOURCE_DIR) + message("Building ${LLVM_SUBPROJECT_TITLE} as a standalone project.") + project(${LLVM_SUBPROJECT_TITLE} CXX C) + set(EUDSLPY_STANDALONE_BUILD ON) +else() + enable_language(CXX C) + set(EUDSLPY_STANDALONE_BUILD OFF) +endif() + +if(EUDSLPY_STANDALONE_BUILD) + find_package(LLVM REQUIRED CONFIG) + find_package(MLIR REQUIRED CONFIG) + find_package(Clang REQUIRED CONFIG PATHS "${LLVM_BINARY_DIR}" NO_DEFAULT_PATH NO_CMAKE_FIND_ROOT_PATH) + + message(STATUS "Using LLVMConfig.cmake in: ${LLVM_DIR}") + message(STATUS "Using MLIRConfig.cmake in: ${MLIR_DIR}") + message(STATUS "Using ClangConfig.cmake in: ${Clang_DIR}") + + set(LLVM_RUNTIME_OUTPUT_INTDIR ${CMAKE_BINARY_DIR}/bin) + set(LLVM_LIBRARY_OUTPUT_INTDIR ${CMAKE_BINARY_DIR}/lib) + set(MLIR_BINARY_DIR ${CMAKE_BINARY_DIR}) + + list(APPEND CMAKE_MODULE_PATH "${MLIR_CMAKE_DIR}") + list(APPEND CMAKE_MODULE_PATH "${LLVM_CMAKE_DIR}") + list(APPEND CMAKE_MODULE_PATH "${CLANG_CMAKE_DIR}") + + include(TableGen) + include(AddLLVM) + include(AddMLIR) + include(AddClang) + # TODO(max): probably don't need this anymore after landing the nanobind fix? + # technically we need this on windows too but our LLVM is compiled without exception handling + # and that breaks windows + if(NOT WIN32) + include(HandleLLVMOptions) + endif() + # for out-of-tree MLIR_INCLUDE_DIR points to the build dir by default + # and MLIR_INCLUDE_DIRS points to the correct place + set(MLIR_INCLUDE_DIR ${MLIR_INCLUDE_DIRS}) +else() + # turning LLVM -DLLVM_OPTIMIZED_TABLEGEN=ON builds some stuff in the NATIVE dir + # but not everything so LLVM_BINARY_DIR isn't correct + string(REPLACE "NATIVE" "" LLVM_BINARY_DIR "${LLVM_BINARY_DIR}") + # Build via external projects mechanism + set(LLVM_INCLUDE_DIR ${LLVM_MAIN_SRC_DIR}/include) + set(LLVM_GENERATED_INCLUDE_DIR ${LLVM_BINARY_DIR}/include) + set(LLVM_INCLUDE_DIRS "${LLVM_INCLUDE_DIR};${LLVM_GENERATED_INCLUDE_DIR}") + + set(MLIR_MAIN_SRC_DIR ${LLVM_MAIN_SRC_DIR}/../mlir) + set(MLIR_INCLUDE_DIR ${MLIR_MAIN_SRC_DIR}/include) + set(MLIR_GENERATED_INCLUDE_DIR ${LLVM_BINARY_DIR}/tools/mlir/include) + set(MLIR_INCLUDE_DIRS "${MLIR_INCLUDE_DIR};${MLIR_GENERATED_INCLUDE_DIR}") + + set(CLANG_MAIN_SRC_DIR ${LLVM_MAIN_SRC_DIR}/../clang) + set(CLANG_INCLUDE_DIR ${CLANG_MAIN_SRC_DIR}/include) + set(CLANG_GENERATED_INCLUDE_DIR ${LLVM_BINARY_DIR}/tools/clang/include) + set(CLANG_INCLUDE_DIRS "${CLANG_INCLUDE_DIR};${CLANG_GENERATED_INCLUDE_DIR}") +endif() + +include_directories(${LLVM_INCLUDE_DIRS}) +include_directories(${MLIR_INCLUDE_DIRS}) +include_directories(${CLANG_INCLUDE_DIRS}) +link_directories(${LLVM_BUILD_LIBRARY_DIR}) +add_definitions(${LLVM_DEFINITIONS}) + +if(NOT TARGET LLVMSupport) + message(FATAL_ERROR "LLVMSupport not found") +endif() + +add_llvm_executable(eudsl-nbgen + eudsl-nbgen.cpp DISABLE_LLVM_LINK_LLVM_DYLIB PARTIAL_SOURCES_INTENDED) +target_link_libraries(eudsl-nbgen + PRIVATE + clangAST + clangBasic + clangFrontend + clangSerialization + clangTooling + LLVMSupport +) + +install(TARGETS eudslpy-gen RUNTIME DESTINATION "${CMAKE_INSTALL_PREFIX}/bin") \ No newline at end of file diff --git a/projects/eudsl-nbgen/eudsl-nbgen.cpp b/projects/eudsl-nbgen/eudsl-nbgen.cpp new file mode 100644 index 00000000..0b65ab64 --- /dev/null +++ b/projects/eudsl-nbgen/eudsl-nbgen.cpp @@ -0,0 +1,722 @@ +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// Copyright (c) 2024. + +#include "clang/AST/ASTConsumer.h" +#include "clang/AST/LexicallyOrderedRecursiveASTVisitor.h" +#include "clang/AST/Mangle.h" +#include "clang/AST/RecursiveASTVisitor.h" +#include "clang/Frontend/CompilerInstance.h" +#include "clang/Frontend/FrontendAction.h" +#include "clang/Sema/Lookup.h" +#include "clang/Sema/Sema.h" +#include "clang/Tooling/Tooling.h" +#include "llvm/Support/FormatVariadic.h" +#include "llvm/Support/ToolOutputFile.h" + +#include + +static llvm::cl::OptionCategory EUDSLPYGenCat("Options for eudslpy-gen"); +static llvm::cl::opt InputFilename(llvm::cl::Positional, + llvm::cl::desc(""), + llvm::cl::Required, + llvm::cl::cat(EUDSLPYGenCat)); + +static llvm::cl::list + IncludeDirs("I", llvm::cl::desc("Directory of include files"), + llvm::cl::value_desc("directory"), llvm::cl::Prefix, + llvm::cl::cat(EUDSLPYGenCat)); + +static llvm::cl::opt + OutputFilename("o", llvm::cl::desc("Output filename"), + llvm::cl::value_desc("filename"), llvm::cl::Required, + llvm::cl::cat(EUDSLPYGenCat)); + +static llvm::cl::list + Namespaces("namespaces", llvm::cl::desc("Namespaces to generate from"), + llvm::cl::CommaSeparated, llvm::cl::cat(EUDSLPYGenCat)); + +static llvm::cl::list + Defines("D", llvm::cl::desc("Name of the macro to be defined"), + llvm::cl::value_desc("macro name"), llvm::cl::Prefix, + llvm::cl::cat(EUDSLPYGenCat)); + +static llvm::cl::opt ShardSize("shard-size", llvm::cl::desc("Shard size"), + llvm::cl::value_desc("shard size"), + llvm::cl::cat(EUDSLPYGenCat), + llvm::cl::init(10)); + +static bool filterInNamespace(const std::string &s) { + if (Namespaces.empty()) + return true; + for (auto ns : Namespaces) + if (ns == s || ("::" + ns == s)) + return true; + return false; +} + +static std::string +getNBBindClassName(const std::string &qualifiedNameAsString) { + std::string s = qualifiedNameAsString; + s = std::regex_replace(s, std::regex(R"(\s+)"), ""); + s = std::regex_replace(s, std::regex("::"), "_"); + s = std::regex_replace(s, std::regex("[<|>]"), "__"); + s = std::regex_replace(s, std::regex(R"(\*)"), "___"); + s = std::regex_replace(s, std::regex(","), "____"); + return s; +} + +static std::string getPyClassName(const std::string &qualifiedNameAsString) { + std::string s = qualifiedNameAsString; + s = std::regex_replace(s, std::regex(R"(\s+)"), ""); + s = std::regex_replace(s, std::regex(R"(\*)"), ""); + s = std::regex_replace(s, std::regex("<"), "["); + s = std::regex_replace(s, std::regex(">"), "]"); + return s; +} + +static std::string snakeCase(const std::string &name) { + std::string s = name; + s = std::regex_replace(s, std::regex(R"(([A-Z]+)([A-Z][a-z]))"), "$1_$2"); + s = std::regex_replace(s, std::regex(R"(([a-z\d])([A-Z]))"), "$1_$2"); + s = std::regex_replace(s, std::regex("-"), "_"); + std::transform(s.begin(), s.end(), s.begin(), + [](unsigned char c) { return std::tolower(c); }); + return s; +} + +static clang::PrintingPolicy getPrintingPolicy(bool canonical = true) { + clang::LangOptions lo; + lo.CPlusPlus = true; + clang::PrintingPolicy p(lo); + // TODO(max): none of this really does anything except PrintCanonical + // keep namespaces + p.FullyQualifiedName = true; + p.SuppressScope = false; + p.SuppressInlineNamespace = false; + p.PrintCanonicalTypes = canonical; + p.Bool = true; + + return p; +} + +// stolen from +// https://github.com/llvm/llvm-project/blob/99dddef340e566e9d303010f1219f7d7d6d37a11/clang/lib/Sema/SemaChecking.cpp#L7055 +// Determines if the specified is a C++ class or struct containing +// a member with the specified name and kind (e.g. a CXXMethodDecl named +// "c_str()"). +template +static llvm::SmallPtrSet findOverloads(clang::FunctionDecl *decl, + clang::Sema &s) { + llvm::SmallPtrSet results; + clang::LookupResult r(s, &s.Context.Idents.get(decl->getNameAsString()), + decl->getLocation(), clang::Sema::LookupOrdinaryName); + r.suppressDiagnostics(); + if (s.LookupQualifiedName(r, decl->getDeclContext())) + for (clang::LookupResult::iterator i = r.begin(), e = r.end(); i != e; + ++i) { + clang::NamedDecl *namedDecl = (*i)->getUnderlyingDecl(); + if (T *fk = llvm::dyn_cast(namedDecl)) + results.insert(fk); + } + return results; +} + +// TODO(max): split this into two functions (one for names and one for types) +static std::string sanitizeNameOrType(std::string nameOrType, + int emptyIdx = 0) { + if (nameOrType == "from") + nameOrType = "from_"; + else if (nameOrType == "except") + nameOrType = "except_"; + else if (nameOrType == "") + nameOrType = std::string(emptyIdx + 1, '_'); + else if (nameOrType.rfind("ArrayRef", 0) == 0) + nameOrType = "llvm::" + nameOrType; + if (std::regex_search(nameOrType, std::regex(R"(std::__1)"))) + nameOrType = std::regex_replace(nameOrType, std::regex("std::__1"), "std"); + return nameOrType; +} + +// emit a lambda body to disambiguate/break ties amongst overloads +// TODO(max):: overloadimpl or whatever should work but it doesn't... +std::string emitNBLambdaBody(clang::FunctionDecl *decl, + llvm::SmallVector paramNames, + llvm::SmallVector paramTypes) { + std::string typedParamsStr; + if (decl->getNumParams()) { + llvm::SmallVector typedParams = + llvm::to_vector(llvm::map_range( + llvm::zip(paramTypes, paramNames), + [](std::tuple item) -> std::string { + auto [t, n] = item; + return llvm::formatv("{0} {1}", t, n); + })); + typedParamsStr = llvm::join(typedParams, ", "); + } + // since we're emitting a body, we need to do std::move for some + // unique_ptrs + llvm::SmallVector newParamNames(paramNames); + for (auto [idx, item] : + llvm::enumerate(llvm::zip(paramTypes, newParamNames))) { + // TODO(max): characterize this condition better... + auto [t, n] = item; + if ((t.rfind("std::unique_ptr") != std::string::npos && t.back() != '&') || + t.rfind("&&") != std::string::npos) + n = llvm::formatv("std::move({0})", n); + } + std::string newParamNamesStr = llvm::join(newParamNames, ", "); + std::string funcRef; + std::string returnType = sanitizeNameOrType( + decl->getReturnType().getAsString(getPrintingPolicy())); + if (decl->isStatic() || !decl->isCXXClassMember()) { + funcRef = llvm::formatv("\n []({0}) -> {1} {{\n return {2}({3});\n }", + typedParamsStr, returnType, + decl->getQualifiedNameAsString(), newParamNamesStr); + } else { + assert(decl->isCXXClassMember() && "expected class member"); + if (decl->getNumParams()) + typedParamsStr = llvm::formatv("self, {0}", typedParamsStr); + else + typedParamsStr = "self"; + const clang::CXXRecordDecl *parentRecord = + llvm::cast(decl->getParent()); + funcRef = llvm::formatv( + "\n []({0}& {1}) -> {2} {{\n return self.{3}({4});\n }", + parentRecord->getQualifiedNameAsString(), typedParamsStr, returnType, + decl->getNameAsString(), newParamNamesStr); + } + return funcRef; +} + +static bool +emitClassMethodOrFunction(clang::FunctionDecl *decl, + clang::CompilerInstance &ci, + std::shared_ptr outputFile) { + llvm::SmallVector paramNames; + llvm::SmallVector paramTypes; + for (unsigned i = 0; i < decl->getNumParams(); ++i) { + clang::ParmVarDecl *param = decl->getParamDecl(i); + std::string name = param->getNameAsString(); + auto t = param->getType(); + bool canonical = true; + // TODO(max): this is dumb... (maybe there's a way to check where the + // typedef is defined...) + // word boundary excludes x86_amx_tdpbf16ps + if (std::regex_search(t.getAsString(), std::regex(R"(_t\b)"))) + canonical = false; + std::string paramType = t.getAsString(getPrintingPolicy(canonical)); + paramTypes.push_back(sanitizeNameOrType(paramType)); + paramNames.push_back(sanitizeNameOrType(name, i)); + } + + llvm::SmallPtrSet funcOverloads = + findOverloads(decl, ci.getSema()); + llvm::SmallPtrSet funcTemplOverloads = + findOverloads(decl, ci.getSema()); + std::string funcRef, nbFnName; + + if (auto ctor = llvm::dyn_cast(decl)) { + if (ctor->isDeleted()) + return false; + funcRef = llvm::formatv("nb::init<{0}>()", llvm::join(paramTypes, ", ")); + } else { + if (funcOverloads.size() == 1 && funcTemplOverloads.empty()) { + funcRef = llvm::formatv("&{0}", decl->getQualifiedNameAsString()); + } else { + funcRef = emitNBLambdaBody(decl, paramNames, paramTypes); + } + + nbFnName = snakeCase(decl->getNameAsString()); + if (decl->isOverloadedOperator()) { + // TODO(max): handle overloaded operators + // nbFnName = nbFnName; + } else if (decl->isStatic() && funcOverloads.size() > 1 && + llvm::any_of(funcOverloads, [](clang::FunctionDecl *m) { + return !m->isStatic(); + })) { + // disambiguate static method with non-static overloads (nanobind doesn't + // let you overload static with non-static) see mlir::ElementsAttr + nbFnName += "_static"; + } + } + + std::string paramNamesStr; + if (decl->getNumParams()) { + paramNames = llvm::to_vector( + llvm::map_range(paramNames, [](std::string s) -> std::string { + return llvm::formatv("\"{0}\"_a", snakeCase(s)); + })); + paramNamesStr = ", " + llvm::join(paramNames, ", "); + } + + std::string refInternal, defStr = "def"; + if (decl->isCXXClassMember()) { + if (decl->isStatic()) + defStr = "def_static"; + if (decl->getReturnType()->hasPointerRepresentation()) + refInternal = ", nb::rv_policy::reference_internal"; + } + + std::string sig; + if (!nbFnName.empty()) { + // no clue why but nb has trouble inferring the signature + // (and this causes and assert failure in nb_func_render_signature + // https://github.com/wjakob/nanobind/blob/c2e394eee5d19816871151de43c29b4051fbf9ff/src/nb_func.cpp#L1020 + if (decl->isStatic() && !decl->getNumParams()) + sig = + llvm::formatv(", nb::sig(\"def {0}(/) -> {1}\") ", nbFnName, + decl->getReturnType().getAsString(getPrintingPolicy())); + // uncomment this to debug nb which doesn't tell where the method actually + // is when it fails if (parentRecord) { + // nbFnName += parentRecord->getNameAsString(); + // } + nbFnName = llvm::formatv("\"{0}\", ", nbFnName); + } + + std::string scope = "m"; + if (decl->isCXXClassMember()) { + const clang::CXXRecordDecl *parentRecord = + llvm::cast(decl->getParent()); + scope = getNBBindClassName(parentRecord->getQualifiedNameAsString()); + } + + outputFile->os() << llvm::formatv("{0}.{1}({2}{3}{4}{5}{6});\n", scope, + defStr, nbFnName, funcRef, paramNamesStr, + refInternal, sig); + + return true; +} + +std::string getNBScope(clang::TagDecl *decl) { + std::string scope = "m"; + const clang::DeclContext *declContext = decl->getDeclContext(); + if (declContext->isRecord()) { + const clang::CXXRecordDecl *ctx = + llvm::cast(declContext); + scope = getNBBindClassName(ctx->getQualifiedNameAsString()); + } + return scope; +} + +static bool emitClass(clang::CXXRecordDecl *decl, clang::CompilerInstance &ci, + std::shared_ptr outputFile) { + if (decl->isTemplated()) { + clang::DiagnosticBuilder builder = ci.getDiagnostics().Report( + decl->getLocation(), ci.getDiagnostics().getCustomDiagID( + clang::DiagnosticsEngine::Note, + "template classes not supported yet")); + return false; + } + + std::string scope = getNBScope(decl); + std::string additional = ""; + std::string className = decl->getQualifiedNameAsString(); + if (decl->getNumBases() > 1) { + clang::DiagnosticBuilder builder = ci.getDiagnostics().Report( + decl->getLocation(), ci.getDiagnostics().getCustomDiagID( + clang::DiagnosticsEngine::Note, + "multiple base classes not supported")); + } else if (decl->getNumBases() == 1) { + // handle some known bases that we've already found a wap to bind + clang::CXXBaseSpecifier baseClass = *decl->bases_begin(); + std::string baseName = baseClass.getType().getAsString(getPrintingPolicy()); + // TODO(max): these could be lookups on the corresponding recorddecls using + // sema... + if (baseName.rfind("mlir::Op<", 0) == 0) { + className = llvm::formatv("{0}, mlir::OpState", className); + } else if (baseName.rfind("mlir::detail::StorageUserBase<", 0) == 0) { + llvm::SmallVector templParams; + llvm::StringRef{baseName}.split(templParams, ","); + className = llvm::formatv("{0}, {1}", className, templParams[1]); + } else if (baseName.rfind("mlir::Dialect", 0) == 0 && + className.rfind("mlir::ExtensibleDialect") == + std::string::npos) { + // clang-format off + additional += llvm::formatv("\n .def_static(\"insert_into_registry\", [](mlir::DialectRegistry ®istry) {{ registry.insert<{0}>(); })", className); + additional += llvm::formatv("\n .def_static(\"load_into_context\", [](mlir::MLIRContext &context) {{ return context.getOrLoadDialect<{0}>(); })", className); + // clang-format on + } else if (!llvm::isa( + baseClass.getType()->getAsCXXRecordDecl())) { + className = llvm::formatv("{0}, {1}", className, baseName); + } else { + assert(llvm::isa( + baseClass.getType()->getAsCXXRecordDecl()) && + "expected class template specialization"); + clang::DiagnosticBuilder builder = ci.getDiagnostics().Report( + baseClass.getBeginLoc(), ci.getDiagnostics().getCustomDiagID( + clang::DiagnosticsEngine::Note, + "unknown base templated base class: ")); + builder << baseName << "\n"; + } + } + + std::string autoVar = llvm::formatv( + "auto {0}", getNBBindClassName(decl->getQualifiedNameAsString())); + + outputFile->os() << llvm::formatv( + "\n{0} = nb::class_<{1}>({2}, \"{3}\"){4};\n", autoVar, className, scope, + getPyClassName(decl->getNameAsString()), additional); + + return true; +} + +static bool emitEnum(clang::EnumDecl *decl, clang::CompilerInstance &ci, + std::shared_ptr outputFile) { + outputFile->os() << llvm::formatv("nb::enum_<{0}>({1}, \"{2}\")\n", + decl->getQualifiedNameAsString(), + getNBScope(decl), decl->getNameAsString()); + + int i = 0, nDecls = std::distance(decl->decls_begin(), decl->decls_end()); + for (clang::Decl *cst : decl->decls()) { + clang::EnumConstantDecl *cstDecl = llvm::cast(cst); + outputFile->os() << llvm::formatv(" .value(\"{0}\", {1})", + cstDecl->getNameAsString(), + cstDecl->getQualifiedNameAsString()); + if (i++ < nDecls - 1) + outputFile->os() << "\n"; + else + outputFile->os() << ";\n"; + } + outputFile->os() << "\n"; + return true; +} + +static bool emitField(clang::DeclaratorDecl *field, clang::CompilerInstance &ci, + std::shared_ptr outputFile) { + const clang::CXXRecordDecl *parentRecord = + llvm::cast(field->getLexicalDeclContext()); + + std::string defStr = "def_rw"; + if (clang::VarDecl *vard = llvm::dyn_cast(field)) { + if (vard->isStaticDataMember()) { + if (vard->isConstexpr()) + defStr = "def_ro_static"; + else + defStr = "def_rw_static"; + } + } + + std::string refInternal; + if (field->getType()->hasPointerRepresentation()) + refInternal = ", nb::rv_policy::reference_internal"; + + std::string scope = + getNBBindClassName(parentRecord->getQualifiedNameAsString()); + std::string nbFnName = + llvm::formatv("\"{0}\"", snakeCase(field->getNameAsString())); + outputFile->os() << llvm::formatv("{0}.{1}({2}, &{3}{4});\n", scope, defStr, + nbFnName, field->getQualifiedNameAsString(), + refInternal); + return true; +} + +template +static bool shouldSkip(T *decl) { + auto *encl = llvm::dyn_cast( + decl->getEnclosingNamespaceContext()); + if (!encl) + return true; + // this isn't redundant - filter might be empty but we still don't want to + // bind std:: + if (encl->isStdNamespace() || encl->isInStdNamespace()) + return true; + if (!filterInNamespace(encl->getQualifiedNameAsString())) + return true; + if constexpr (std::is_same_v || + std::is_same_v) { + if (!decl->isCompleteDefinition()) + return true; + } + if (decl->getAccess() == clang::AS_private || + decl->getAccess() == clang::AS_protected) + return true; + if (decl->isImplicit()) + return true; + + return false; +} + +struct HackDeclContext : clang::DeclContext { + bool islastDecl(clang::Decl *d) const { return d == LastDecl; } + clang::Decl *getLastDecl() const { return LastDecl; } +}; + +struct BindingsVisitor + : clang::LexicallyOrderedRecursiveASTVisitor { + BindingsVisitor(clang::CompilerInstance &ci, + std::shared_ptr outputFile) + : LexicallyOrderedRecursiveASTVisitor( + ci.getSourceManager()), + ci(ci), outputFile(outputFile), + mc(clang::ItaniumMangleContext::create(ci.getASTContext(), + ci.getDiagnostics())) {} + + bool VisitCXXRecordDecl(clang::CXXRecordDecl *decl) { + if (shouldSkip(decl)) + return true; + if (decl->isClass() || decl->isStruct()) { + if (emitClass(decl, ci, outputFile)) + visitedRecords.insert(decl); + } + return true; + } + + // clang-format off + /* method (member function) -> CXXMethodDecl (isStatic() tells statis/non-static) + * non-static data member -> FieldDecl + * static data member -> VarDecl with isStaticDataMember() == true + * static local variable -> VarDecl with isStaticLocal() == true + * non-static local variable -> VarDecl with isLocalVarDecl() == true + * there's also a isLocalVarDeclOrParm if you want local or parameter, etc + */ + // clang-format on + + bool VisitFieldDecl(clang::FieldDecl *decl) { + if (decl->isFunctionOrFunctionTemplate() || decl->isFunctionPointerType()) + return true; + // fields can be methods??? + if (llvm::isa(decl)) + return true; + if (!visitedRecords.contains(decl->getParent())) + return true; + if (decl->getAccess() == clang::AS_private || + decl->getAccess() == clang::AS_protected) + return true; + + if (decl->isAnonymousStructOrUnion()) { + clang::DiagnosticBuilder builder = ci.getDiagnostics().Report( + decl->getLocation(), ci.getDiagnostics().getCustomDiagID( + clang::DiagnosticsEngine::Note, + "anon structs/union fields not supported")); + return true; + } + if (decl->isBitField()) + return true; + + emitField(decl, ci, outputFile); + return true; + } + + bool VisitVarDecl(clang::VarDecl *decl) { + if (decl->isFunctionOrMethodVarDecl()) + return true; + if (llvm::isa(decl)) + return true; + if (auto parent = llvm::dyn_cast( + decl->getLexicalDeclContext())) { + if (visitedRecords.contains(parent)) { + if (decl->getAccess() == clang::AS_private || + decl->getAccess() == clang::AS_protected) + return true; + emitField(decl, ci, outputFile); + } + } + + return true; + } + + bool VisitCXXMethodDecl(clang::CXXMethodDecl *decl) { + if (shouldSkip(decl) || llvm::isa(decl) || + !visitedRecords.contains(decl->getParent())) + return true; + if (decl->isTemplated() || decl->isTemplateDecl() || + decl->isTemplateInstantiation() || + decl->isFunctionTemplateSpecialization()) { + clang::DiagnosticBuilder builder = ci.getDiagnostics().Report( + decl->getLocation(), ci.getDiagnostics().getCustomDiagID( + clang::DiagnosticsEngine::Note, + "template methods not supported yet")); + return true; + } + if (decl->getFriendObjectKind()) { + clang::DiagnosticBuilder builder = ci.getDiagnostics().Report( + decl->getLocation(), ci.getDiagnostics().getCustomDiagID( + clang::DiagnosticsEngine::Note, + "friend functions not supported")); + return true; + } + emitClassMethodOrFunction(decl, ci, outputFile); + return true; + } + + bool VisitFunctionDecl(clang::FunctionDecl *decl) { + if (shouldSkip(decl) || decl->isCXXClassMember()) + return true; + // clang-format off + // this + // template ::std::optional symbolizeEnum(::llvm::StringRef); + // is not a `TemplateDecl` but it is `Templated`... + // on the other hand every method in a template class `isTemplated` even + // if the template params don't play in the method decl (which is why this visitor can't be combined with VisitCXXMethodDecl) + // clang-format on + if (decl->isTemplated() || decl->isTemplateDecl() || + decl->isTemplateInstantiation() || + decl->isFunctionTemplateSpecialization()) { + clang::DiagnosticBuilder builder = ci.getDiagnostics().Report( + decl->getLocation(), ci.getDiagnostics().getCustomDiagID( + clang::DiagnosticsEngine::Note, + "template functions not supported yet")); + return true; + } + emitClassMethodOrFunction(decl, ci, outputFile); + return true; + } + + bool VisitEnumDecl(clang::EnumDecl *decl) { + if (shouldSkip(decl)) + return true; + if (decl->getQualifiedNameAsString().rfind("unnamed enum") != + std::string::npos) + return true; + emitEnum(llvm::cast(decl), ci, outputFile); + return true; + } + + void maybeEmitShard(const clang::CXXRecordDecl *decl) { + clang::CXXBaseSpecifier baseClass = *decl->bases_begin(); + std::string baseName = baseClass.getType().getAsString(getPrintingPolicy()); + if (baseName.rfind("mlir::Op<", 0) == 0) { + ++opClassesSeen; + if (opClassesSeen % ShardSize == 0) + outputFile->os() << llvm::formatv("// eudslpy-gen-shard {}\n", + int(opClassesSeen / ShardSize)); + } + } + + // implicit methods are the "last decls" + bool shouldVisitImplicitCode() const { return true; } + + // TODO(max): this is a hack and not stable + bool VisitDecl(clang::Decl *decl) { + const clang::DeclContext *declContext = decl->getDeclContext(); + HackDeclContext *ctx = + static_cast(decl->getDeclContext()); + if (declContext && declContext->isRecord()) { + const clang::CXXRecordDecl *recordDecl = + llvm::cast(declContext); + if (visitedRecords.contains(recordDecl) && ctx->islastDecl(decl)) { + outputFile->os() << "// " << recordDecl->getQualifiedNameAsString() + << "\n\n"; + if (recordDecl->getNumBases() == 1) + maybeEmitShard(recordDecl); + } + } + return true; + } + + clang::CompilerInstance &ci; + std::shared_ptr outputFile; + std::unique_ptr mc; + llvm::DenseSet visitedRecords; + int opClassesSeen{0}; +}; + +struct ClassStructEnumConsumer : clang::ASTConsumer { + ClassStructEnumConsumer( + clang::CompilerInstance &ci, + const std::shared_ptr &outputFile) + : visitor(ci, outputFile) {} + + void HandleTranslationUnit(clang::ASTContext &context) override { + if (!visitor.ci.hasSema()) + visitor.ci.createSema(clang::TU_Prefix, + /*CompletionConsumer*/ nullptr); + visitor.TraverseDecl(context.getTranslationUnitDecl()); + } + bool shouldSkipFunctionBody(clang::Decl *) override { return true; } + BindingsVisitor visitor; +}; + +struct GenerateBindingsAction : clang::ASTFrontendAction { + explicit GenerateBindingsAction( + const std::shared_ptr &outputFile) + : outputFile(outputFile) {} + std::unique_ptr + CreateASTConsumer(clang::CompilerInstance &compiler, + llvm::StringRef inFile) override { + compiler.getDiagnosticOpts().ShowLevel = true; + return std::make_unique(compiler, outputFile); + } + + // PARSE_INCOMPLETE + clang::TranslationUnitKind getTranslationUnitKind() override { + return clang::TU_Prefix; + } + + std::shared_ptr outputFile; +}; + +int main(int argc, char **argv) { + llvm::cl::ParseCommandLineOptions(argc, argv); + + llvm::ErrorOr> fileOrErr = + llvm::MemoryBuffer::getFileOrSTDIN(InputFilename); + if (std::error_code ec = fileOrErr.getError()) { + llvm::errs() << "Could not open input file: " << ec.message() << "\n"; + return -1; + } + llvm::StringRef buffer = fileOrErr.get()->getBuffer(); + + if (OutputFilename.empty()) { + llvm::errs() << "-o outputfile must be set.\n"; + return -1; + } + + std::error_code error; + auto outputFile = std::make_shared( + OutputFilename, error, llvm::sys::fs::OF_None); + if (error) { + llvm::errs() << "cannot open output file '" + OutputFilename + + "': " + error.message(); + return -1; + } + std::vector args{ + "-E", + "-xc++", + "-std=c++17", + "-fdirectives-only", + "-fkeep-system-includes", + "-fdelayed-template-parsing", + "-Wno-unused-command-line-argument", + // "-v", + // annoyingly clang will insert -internal-isystem with relative paths and + // those could hit on the build dir (which have headers and relationships + // amongst them that won't necessarily be valid) more annoyingly different + // toolchains decide differently which flag determines whether to include + // the tell-tale sign is if you uncomment -v and see "ignoring + // non-existant directory..." in the output + "-nobuiltininc", + "-nostdinc", + "-nostdinc++", + "-nostdlibinc", + }; + int clangArgs = args.size(); + for (const auto &includeDir : IncludeDirs) + args.emplace_back(llvm::formatv("-I{0}", includeDir)); + for (const auto &define : Defines) + args.emplace_back(llvm::formatv("-D{0}", define)); + + outputFile->os() << "// Generated with eudslpy-gen args:\n"; + outputFile->os() << "// " << InputFilename << " "; + auto arg = args.begin(); + for (std::advance(arg, clangArgs); arg != args.end(); ++arg) { + outputFile->os() << *arg << ' '; + } + for (auto ns : Namespaces) + outputFile->os() << "-namespaces=" << ns << " "; + outputFile->os() << "-o " << OutputFilename << "\n"; + outputFile->os() << "\n"; + + if (!clang::tooling::runToolOnCodeWithArgs( + std::make_unique(outputFile), buffer, args, + InputFilename)) { + llvm::errs() << "bindings generation failed.\n"; + return -1; + } + + outputFile->keep(); + + return 0; +} From 497314ffa3328a368a8d99fd3bcb48f18a029aa4 Mon Sep 17 00:00:00 2001 From: Maksim Levental Date: Thu, 2 Jan 2025 13:13:05 -0500 Subject: [PATCH 2/6] [eudsl][nbgen] start refactor --- .github/actions/setup_base/action.yml | 2 + .../workflows/build_test_release_eudsl.yml | 36 +++++++------ projects/eudsl-nbgen/CMakeLists.txt | 49 +++++++++-------- .../cmake/eudsl-nbgen-config.cmake} | 54 ++++++++++++++----- .../cmake/make_generated_shards.py} | 0 projects/eudsl-nbgen/pyproject.toml | 35 ++++++++++++ projects/eudsl-nbgen/src/__init__.py | 13 +++++ projects/eudsl-nbgen/src/__main__.py | 28 ++++++++++ .../eudsl-nbgen/{ => src}/eudsl-nbgen.cpp | 0 projects/eudsl-py/CMakeLists.txt | 30 +---------- projects/eudsl-py/pyproject.toml | 18 +++++-- .../eudsl-tblgen/{src => }/CMakeLists.txt | 6 ++- projects/eudsl-tblgen/pyproject.toml | 20 +++++-- 13 files changed, 200 insertions(+), 91 deletions(-) rename projects/{eudsl-py/cmake/EUDSLPYConfig.cmake => eudsl-nbgen/cmake/eudsl-nbgen-config.cmake} (85%) rename projects/{eudsl-py/cmake/make_generated_registration.py => eudsl-nbgen/cmake/make_generated_shards.py} (100%) create mode 100644 projects/eudsl-nbgen/pyproject.toml create mode 100644 projects/eudsl-nbgen/src/__init__.py create mode 100644 projects/eudsl-nbgen/src/__main__.py rename projects/eudsl-nbgen/{ => src}/eudsl-nbgen.cpp (100%) rename projects/eudsl-tblgen/{src => }/CMakeLists.txt (94%) diff --git a/.github/actions/setup_base/action.yml b/.github/actions/setup_base/action.yml index dfaeb918..9aabec35 100644 --- a/.github/actions/setup_base/action.yml +++ b/.github/actions/setup_base/action.yml @@ -150,6 +150,8 @@ runs: echo "CCACHE_SLOPPINESS=include_file_ctime,include_file_mtime,time_macros" >> $GITHUB_ENV echo "CCACHE_CPP2=true" >> $GITHUB_ENV echo "CCACHE_UMASK=002" >> $GITHUB_ENV + + ccache -z echo "CMAKE_GENERATOR=Ninja" >> $GITHUB_ENV echo "CMAKE_C_COMPILER_LAUNCHER=ccache" >> $GITHUB_ENV diff --git a/.github/workflows/build_test_release_eudsl.yml b/.github/workflows/build_test_release_eudsl.yml index 8e1dc211..7b00e3be 100644 --- a/.github/workflows/build_test_release_eudsl.yml +++ b/.github/workflows/build_test_release_eudsl.yml @@ -104,17 +104,6 @@ jobs: curl -sLO $RELEASE_URL tar xf $RELEASE_PREFIX*.tar.gz - if [[ "${{ matrix.os }}" == "ubuntu" ]]; then - echo "LLVM_DIR=/host/$PWD/llvm-install/lib/cmake/llvm" >> $GITHUB_ENV - echo "MLIR_DIR=/host/$PWD/llvm-install/lib/cmake/mlir" >> $GITHUB_ENV - echo "Clang_DIR=/host/$PWD/llvm-install/lib/cmake/clang" >> $GITHUB_ENV - echo "CCACHE_DIR=/host/$CCACHE_DIR" >> $GITHUB_ENV - else - echo "LLVM_DIR=$PWD/llvm-install/lib/cmake/llvm" >> $GITHUB_ENV - echo "MLIR_DIR=$PWD/llvm-install/lib/cmake/mlir" >> $GITHUB_ENV - echo "Clang_DIR=$PWD/llvm-install/lib/cmake/clang" >> $GITHUB_ENV - fi - # since linux builds execute in the cibuildwheel almalinux container if [[ "${{ matrix.os }}" == "ubuntu" ]]; then echo CC=clang >> $GITHUB_ENV @@ -134,19 +123,36 @@ jobs: - name: "Build eudsl-tblgen" run: | - ccache -z + if [[ "${{ matrix.os }}" == "ubuntu" ]]; then + export CMAKE_PREFIX_PATH="/host/$PWD/llvm-install" + export CCACHE_DIR="/host/$CCACHE_DIR" + else + export CMAKE_PREFIX_PATH="$PWD/llvm-install" + fi + $python3_command -m cibuildwheel "$PWD/projects/eudsl-tblgen" --output-dir wheelhouse - ccache -s + + - name: "Build eudsl-nbgen" + run: | + + # this is not run in cibuildwheel container so its different + export CMAKE_PREFIX_PATH="$PWD/llvm-install" + $python3_command -m pip wheel "$PWD/projects/eudsl-nbgen" -w wheelhouse - name: "Build eudsl-py" if: ${{ ! startsWith(matrix.os, 'windows') }} run: | + if [[ "${{ matrix.os }}" == "ubuntu" ]]; then + export CMAKE_PREFIX_PATH="/host/$PWD/llvm-install" + export CCACHE_DIR="/host/$CCACHE_DIR" + else + export CMAKE_PREFIX_PATH="$PWD/llvm-install" + fi + # prevent OOM on free GHA export DISABLE_COMPILE_OPT="${{ inputs.debug_with_tmate }}" - ccache -z $python3_command -m cibuildwheel "$PWD/projects/eudsl-py" --output-dir wheelhouse - ccache -s - name: "Save cache" uses: actions/cache/save@v3 diff --git a/projects/eudsl-nbgen/CMakeLists.txt b/projects/eudsl-nbgen/CMakeLists.txt index 72fa12cb..82936d78 100644 --- a/projects/eudsl-nbgen/CMakeLists.txt +++ b/projects/eudsl-nbgen/CMakeLists.txt @@ -2,38 +2,33 @@ cmake_minimum_required(VERSION 3.29) set(CMAKE_CXX_STANDARD 17) set(CMAKE_EXPORT_COMPILE_COMMANDS ON) -set(LLVM_SUBPROJECT_TITLE "EUDSLPY") -set(EUDSLPY_BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR}) +set(LLVM_SUBPROJECT_TITLE "EUDSL_NBGEN") +set(EUDSL_NBGEN_BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR}) if (CMAKE_SOURCE_DIR STREQUAL CMAKE_CURRENT_SOURCE_DIR) message("Building ${LLVM_SUBPROJECT_TITLE} as a standalone project.") project(${LLVM_SUBPROJECT_TITLE} CXX C) - set(EUDSLPY_STANDALONE_BUILD ON) + set(EUDSL_NBGEN_STANDALONE_BUILD ON) else() enable_language(CXX C) - set(EUDSLPY_STANDALONE_BUILD OFF) + set(EUDSL_NBGEN_STANDALONE_BUILD OFF) endif() -if(EUDSLPY_STANDALONE_BUILD) +if(EUDSL_NBGEN_STANDALONE_BUILD) find_package(LLVM REQUIRED CONFIG) - find_package(MLIR REQUIRED CONFIG) find_package(Clang REQUIRED CONFIG PATHS "${LLVM_BINARY_DIR}" NO_DEFAULT_PATH NO_CMAKE_FIND_ROOT_PATH) message(STATUS "Using LLVMConfig.cmake in: ${LLVM_DIR}") - message(STATUS "Using MLIRConfig.cmake in: ${MLIR_DIR}") message(STATUS "Using ClangConfig.cmake in: ${Clang_DIR}") set(LLVM_RUNTIME_OUTPUT_INTDIR ${CMAKE_BINARY_DIR}/bin) set(LLVM_LIBRARY_OUTPUT_INTDIR ${CMAKE_BINARY_DIR}/lib) - set(MLIR_BINARY_DIR ${CMAKE_BINARY_DIR}) - list(APPEND CMAKE_MODULE_PATH "${MLIR_CMAKE_DIR}") list(APPEND CMAKE_MODULE_PATH "${LLVM_CMAKE_DIR}") list(APPEND CMAKE_MODULE_PATH "${CLANG_CMAKE_DIR}") include(TableGen) include(AddLLVM) - include(AddMLIR) include(AddClang) # TODO(max): probably don't need this anymore after landing the nanobind fix? # technically we need this on windows too but our LLVM is compiled without exception handling @@ -41,9 +36,6 @@ if(EUDSLPY_STANDALONE_BUILD) if(NOT WIN32) include(HandleLLVMOptions) endif() - # for out-of-tree MLIR_INCLUDE_DIR points to the build dir by default - # and MLIR_INCLUDE_DIRS points to the correct place - set(MLIR_INCLUDE_DIR ${MLIR_INCLUDE_DIRS}) else() # turning LLVM -DLLVM_OPTIMIZED_TABLEGEN=ON builds some stuff in the NATIVE dir # but not everything so LLVM_BINARY_DIR isn't correct @@ -53,11 +45,6 @@ else() set(LLVM_GENERATED_INCLUDE_DIR ${LLVM_BINARY_DIR}/include) set(LLVM_INCLUDE_DIRS "${LLVM_INCLUDE_DIR};${LLVM_GENERATED_INCLUDE_DIR}") - set(MLIR_MAIN_SRC_DIR ${LLVM_MAIN_SRC_DIR}/../mlir) - set(MLIR_INCLUDE_DIR ${MLIR_MAIN_SRC_DIR}/include) - set(MLIR_GENERATED_INCLUDE_DIR ${LLVM_BINARY_DIR}/tools/mlir/include) - set(MLIR_INCLUDE_DIRS "${MLIR_INCLUDE_DIR};${MLIR_GENERATED_INCLUDE_DIR}") - set(CLANG_MAIN_SRC_DIR ${LLVM_MAIN_SRC_DIR}/../clang) set(CLANG_INCLUDE_DIR ${CLANG_MAIN_SRC_DIR}/include) set(CLANG_GENERATED_INCLUDE_DIR ${LLVM_BINARY_DIR}/tools/clang/include) @@ -65,7 +52,6 @@ else() endif() include_directories(${LLVM_INCLUDE_DIRS}) -include_directories(${MLIR_INCLUDE_DIRS}) include_directories(${CLANG_INCLUDE_DIRS}) link_directories(${LLVM_BUILD_LIBRARY_DIR}) add_definitions(${LLVM_DEFINITIONS}) @@ -74,8 +60,8 @@ if(NOT TARGET LLVMSupport) message(FATAL_ERROR "LLVMSupport not found") endif() -add_llvm_executable(eudsl-nbgen - eudsl-nbgen.cpp DISABLE_LLVM_LINK_LLVM_DYLIB PARTIAL_SOURCES_INTENDED) +# this will set the rpath of the exe to be `../lib`, where is where we deposit libLLVM.so below +add_llvm_executable(eudsl-nbgen src/eudsl-nbgen.cpp) target_link_libraries(eudsl-nbgen PRIVATE clangAST @@ -83,7 +69,24 @@ target_link_libraries(eudsl-nbgen clangFrontend clangSerialization clangTooling - LLVMSupport ) -install(TARGETS eudslpy-gen RUNTIME DESTINATION "${CMAKE_INSTALL_PREFIX}/bin") \ No newline at end of file +string(TOLOWER ${LLVM_SUBPROJECT_TITLE} EUDSL_NBGEN_INSTALL_DATADIR) +# https://github.com/scikit-build/scikit-build-core/blob/a887a9b6c057b4ce9d3cfd53ae24e73caf1395a2/docs/build.md?plain=1#L139-L148 +# actually installs to venv/bin +install(TARGETS eudsl-nbgen RUNTIME DESTINATION "${EUDSL_NBGEN_INSTALL_DATADIR}.data/scripts") +# this actually installs into venv/lib +install(IMPORTED_RUNTIME_ARTIFACTS LLVM LIBRARY DESTINATION "${EUDSL_NBGEN_INSTALL_DATADIR}.data/data/lib") + +install( + DIRECTORY src/ + DESTINATION "${EUDSL_NBGEN_INSTALL_DATADIR}" + FILES_MATCHING PATTERN "*\.py" +) + +install( + FILES + cmake/eudsl-nbgen-config.cmake + cmake/make_generated_shards.py + DESTINATION "${EUDSL_NBGEN_INSTALL_DATADIR}/cmake" +) diff --git a/projects/eudsl-py/cmake/EUDSLPYConfig.cmake b/projects/eudsl-nbgen/cmake/eudsl-nbgen-config.cmake similarity index 85% rename from projects/eudsl-py/cmake/EUDSLPYConfig.cmake rename to projects/eudsl-nbgen/cmake/eudsl-nbgen-config.cmake index 6ff8617b..fa853ee8 100644 --- a/projects/eudsl-py/cmake/EUDSLPYConfig.cmake +++ b/projects/eudsl-nbgen/cmake/eudsl-nbgen-config.cmake @@ -5,20 +5,6 @@ # copy-pasta from AddMLIR.cmake/AddLLVM.cmake/TableGen.cmake -# no clue why but with LLVM_LINK_LLVM_DYLIB even static libs depend on LLVM -get_property(MLIR_ALL_LIBS GLOBAL PROPERTY MLIR_ALL_LIBS) -foreach(_lib ${MLIR_ALL_LIBS}) - get_target_property(_interface_link_libraries ${_lib} INTERFACE_LINK_LIBRARIES) - if(NOT _interface_link_libraries) - continue() - endif() - list(REMOVE_DUPLICATES _interface_link_libraries) - list(REMOVE_ITEM _interface_link_libraries LLVM) - # for some reason, explicitly adding below as a link library doesn't work - missing symbols... - list(APPEND _interface_link_libraries LLVMSupport) - set_target_properties(${_lib} PROPERTIES INTERFACE_LINK_LIBRARIES "${_interface_link_libraries}") -endforeach() - function(eudslpygen target inputFile) set(EUDSLPYGEN_TARGET_DEFINITIONS ${inputFile}) cmake_parse_arguments(ARG "" "" "DEPENDS;EXTRA_INCLUDES;NAMESPACES" ${ARGN}) @@ -241,3 +227,43 @@ macro(add_eudslpygen target project) set_property(GLOBAL APPEND PROPERTY ${export_upper}_EXPORTS ${target}) endif() endmacro() + +function(patch_mlir_llvm_rpath target) + # hack so we can move libMLIR and libLLVM into the wheel + # see AddLLVM.cmake#llvm_setup_rpath + if(APPLE OR UNIX) + set(_origin_prefix "\$ORIGIN") + if(APPLE) + set(_origin_prefix "@loader_path") + endif() + if (EUDSLPY_STANDALONE_BUILD) + get_target_property(_mlir_loc MLIR LOCATION) + get_target_property(_llvm_loc LLVM LOCATION) + else() + set(_mlir_loc "$") + set(_llvm_loc "$") + endif() + set(_old_rpath "${_origin_prefix}/../lib${LLVM_LIBDIR_SUFFIX}") + if(APPLE) + if (EXISTS ${_mlir_loc}) + execute_process(COMMAND install_name_tool -rpath "${_old_rpath}" ${_origin_prefix} "${_mlir_loc}" ERROR_VARIABLE rpath_err) + endif() + if (EXISTS ${_llvm_loc}) + execute_process(COMMAND install_name_tool -rpath "${_old_rpath}" ${_origin_prefix} "${_llvm_loc}" ERROR_VARIABLE rpath_err) + endif() + # maybe already updated... + if (rpath_err AND NOT rpath_err MATCHES "no LC_RPATH load command with path: ${_old_rpath}") + message(FATAL_ERROR "couldn't update rpath because: ${rpath_err}") + endif() + else() + # sneaky sneaky - undocumented + if (EXISTS ${_mlir_loc}) + file(RPATH_CHANGE FILE "${_mlir_loc}" OLD_RPATH "${_old_rpath}" NEW_RPATH "${_origin_prefix}") + endif() + if (EXISTS ${_llvm_loc}) + file(RPATH_CHANGE FILE "${_llvm_loc}" OLD_RPATH "${_old_rpath}" NEW_RPATH "${_origin_prefix}") + endif() + endif() + set_target_properties(${target} PROPERTIES INSTALL_RPATH "${_origin_prefix}") + endif() +endfunction() \ No newline at end of file diff --git a/projects/eudsl-py/cmake/make_generated_registration.py b/projects/eudsl-nbgen/cmake/make_generated_shards.py similarity index 100% rename from projects/eudsl-py/cmake/make_generated_registration.py rename to projects/eudsl-nbgen/cmake/make_generated_shards.py diff --git a/projects/eudsl-nbgen/pyproject.toml b/projects/eudsl-nbgen/pyproject.toml new file mode 100644 index 00000000..f1e58aae --- /dev/null +++ b/projects/eudsl-nbgen/pyproject.toml @@ -0,0 +1,35 @@ +# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +# See https://llvm.org/LICENSE.txt for license information. +# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +# Copyright (c) 2024. + +[project] +name = "eudsl-nbgen" +version = "0.0.1" +requires-python = ">=3.9" +[project.urls] +Homepage = "https://github.com/llvm/eudsl" + +[build-system] +requires = [ + "scikit-build-core==0.10.7", + "nanobind==2.4.0", + "typing_extensions==4.12.2", +] +build-backend = "scikit_build_core.build" + +[tool.scikit-build] +minimum-version = "0.4" +build-dir = "build/{wheel_tag}" +# gives you a wheel tagged py3-none +wheel.py-api = "py3" +cmake.source-dir = "." + +[tool.scikit-build.cmake.define] +LLVM_DIR = { env = "LLVM_DIR", default = "EMPTY" } +Clang_DIR = { env = "Clang_DIR", default = "EMPTY" } +CMAKE_PREFIX_PATH = { env = "CMAKE_PREFIX_PATH", default = "EMPTY" } +CMAKE_C_COMPILER_LAUNCHER = { env = "CMAKE_C_COMPILER_LAUNCHER", default = "" } +CMAKE_CXX_COMPILER_LAUNCHER = { env = "CMAKE_CXX_COMPILER_LAUNCHER", default = "" } +CMAKE_CXX_VISIBILITY_PRESET = "hidden" +CMAKE_VERBOSE_MAKEFILE = "ON" diff --git a/projects/eudsl-nbgen/src/__init__.py b/projects/eudsl-nbgen/src/__init__.py new file mode 100644 index 00000000..ef47aa59 --- /dev/null +++ b/projects/eudsl-nbgen/src/__init__.py @@ -0,0 +1,13 @@ +# Copyright (c) 2025. +# +# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +# See https://llvm.org/LICENSE.txt for license information. +# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +from pathlib import Path + + +def cmake_dir() -> str: + return str(Path(__file__).parent / "cmake") + + +__all__ = ["cmake_dir"] diff --git a/projects/eudsl-nbgen/src/__main__.py b/projects/eudsl-nbgen/src/__main__.py new file mode 100644 index 00000000..2fa71d3b --- /dev/null +++ b/projects/eudsl-nbgen/src/__main__.py @@ -0,0 +1,28 @@ +# Copyright (c) 2025. +# +# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +# See https://llvm.org/LICENSE.txt for license information. +# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +import argparse +import sys + +from . import cmake_dir + + +def main() -> None: + parser = argparse.ArgumentParser("eudsl-nbgen") + parser.add_argument( + "--cmake_dir", + action="store_true", + help="Print the path to the eudsl-nbgen CMake module directory.", + ) + args = parser.parse_args() + if not sys.argv[1:]: + parser.print_help() + if args.cmake_dir: + print(cmake_dir()) + + +if __name__ == "__main__": + main() diff --git a/projects/eudsl-nbgen/eudsl-nbgen.cpp b/projects/eudsl-nbgen/src/eudsl-nbgen.cpp similarity index 100% rename from projects/eudsl-nbgen/eudsl-nbgen.cpp rename to projects/eudsl-nbgen/src/eudsl-nbgen.cpp diff --git a/projects/eudsl-py/CMakeLists.txt b/projects/eudsl-py/CMakeLists.txt index 91156d42..245784c6 100644 --- a/projects/eudsl-py/CMakeLists.txt +++ b/projects/eudsl-py/CMakeLists.txt @@ -435,35 +435,7 @@ if ("$ENV{DISABLE_COMPILE_OPT}" MATCHES "true") target_compile_options(eudslpy_ext PRIVATE -O0) endif() -# hack so we can move libMLIR and libLLVM into the wheel -# see AddLLVM.cmake#llvm_setup_rpath -if(APPLE OR UNIX) - set(_origin_prefix "\$ORIGIN") - if(APPLE) - set(_origin_prefix "@loader_path") - endif() - if (EUDSLPY_STANDALONE_BUILD) - get_target_property(_mlir_loc MLIR LOCATION) - get_target_property(_llvm_loc LLVM LOCATION) - else() - set(_mlir_loc "$") - set(_llvm_loc "$") - endif() - set(_old_rpath "${_origin_prefix}/../lib${LLVM_LIBDIR_SUFFIX}") - if(APPLE) - execute_process(COMMAND install_name_tool -rpath "${_old_rpath}" ${_origin_prefix} "${_mlir_loc}" ERROR_VARIABLE rpath_err) - execute_process(COMMAND install_name_tool -rpath "${_old_rpath}" ${_origin_prefix} "${_llvm_loc}" ERROR_VARIABLE rpath_err) - # maybe already updated... - if (rpath_err AND NOT rpath_err MATCHES "no LC_RPATH load command with path: ${_old_rpath}") - message(FATAL_ERROR "couldn't update rpath because: ${rpath_err}") - endif() - else() - # sneaky sneaky - undocumented - file(RPATH_CHANGE FILE "${_mlir_loc}" OLD_RPATH "${_old_rpath}" NEW_RPATH "${_origin_prefix}") - file(RPATH_CHANGE FILE "${_llvm_loc}" OLD_RPATH "${_old_rpath}" NEW_RPATH "${_origin_prefix}") - endif() - set_target_properties(eudslpy_ext PROPERTIES INSTALL_RPATH "${_origin_prefix}") -endif() +patch_mlir_llvm_rpath(eudslpy_ext) # copy libMLIR into the ext dir for wheels install(IMPORTED_RUNTIME_ARTIFACTS MLIR LLVM LIBRARY DESTINATION eudsl) diff --git a/projects/eudsl-py/pyproject.toml b/projects/eudsl-py/pyproject.toml index f23a9e39..c2070deb 100644 --- a/projects/eudsl-py/pyproject.toml +++ b/projects/eudsl-py/pyproject.toml @@ -33,7 +33,7 @@ cmake.source-dir = "." LLVM_DIR = { env = "LLVM_DIR", default = "EMPTY" } MLIR_DIR = { env = "MLIR_DIR", default = "EMPTY" } Clang_DIR = { env = "Clang_DIR", default = "EMPTY" } -CMAKE_PREFIX_PATH = { env = "CMAKE_PREFIX_PATH", default = "EMPTY" } +CMAKE_PREFIX_PATH = { env = "CMAKE_PREFIX_PATH", default = "" } CMAKE_C_COMPILER_LAUNCHER = { env = "CMAKE_C_COMPILER_LAUNCHER", default = "" } CMAKE_CXX_COMPILER_LAUNCHER = { env = "CMAKE_CXX_COMPILER_LAUNCHER", default = "" } CMAKE_CXX_VISIBILITY_PRESET = "hidden" @@ -64,10 +64,11 @@ environment-pass = [ ] before-build = [ "export CCACHE_DIR=$CCACHE_DIR/$(python -c 'import platform; print(platform.python_version())')", - "mkdir -p $CCACHE_DIR" + "mkdir -p $CCACHE_DIR", + "ccache -z" ] # uncomment to make sure ccache is working inside containers -# test-command = "ccache -s" +test-command = "ccache -s" [tool.cibuildwheel.linux] before-all = [ @@ -77,6 +78,15 @@ before-all = [ "tar -xf ccache-4.10.2-linux-x86_64.tar.xz", "pushd ccache-4.10.2-linux-x86_64 && make install && popd" ] +# synchronize TZ with host so ccache files have correct timestamp +container-engine = { name = "docker", create-args = ["-v", "/etc/timezone:/etc/timezone:ro", "-v", "/etc/localtime:/etc/localtime:ro"] } + +[tool.cibuildwheel.macos] +before-build = [ + "ccache -z" +] [tool.cibuildwheel.windows] -before-build = [] +before-build = [ + "ccache -z" +] diff --git a/projects/eudsl-tblgen/src/CMakeLists.txt b/projects/eudsl-tblgen/CMakeLists.txt similarity index 94% rename from projects/eudsl-tblgen/src/CMakeLists.txt rename to projects/eudsl-tblgen/CMakeLists.txt index c8fef467..d8ff2530 100644 --- a/projects/eudsl-tblgen/src/CMakeLists.txt +++ b/projects/eudsl-tblgen/CMakeLists.txt @@ -29,7 +29,11 @@ find_package(nanobind CONFIG REQUIRED) set(CMAKE_LIBRARY_OUTPUT_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/eudsl_tblgen) set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/bin) -nanobind_add_module(eudsl_tblgen_ext NB_STATIC STABLE_ABI eudsl_tblgen_ext.cpp TGParser.cpp TGLexer.cpp) +nanobind_add_module(eudsl_tblgen_ext NB_STATIC STABLE_ABI + src/eudsl_tblgen_ext.cpp + src/TGParser.cpp + src/TGLexer.cpp +) target_link_libraries(eudsl_tblgen_ext PRIVATE LLVMTableGenCommon LLVMTableGen) target_compile_options(eudsl_tblgen_ext PUBLIC diff --git a/projects/eudsl-tblgen/pyproject.toml b/projects/eudsl-tblgen/pyproject.toml index 8061b32f..4a8d241d 100644 --- a/projects/eudsl-tblgen/pyproject.toml +++ b/projects/eudsl-tblgen/pyproject.toml @@ -18,14 +18,14 @@ Homepage = "https://github.com/llvm/eudsl" [tool.scikit-build] minimum-version = "0.4" build-dir = "build/{wheel_tag}" -wheel.py-api = "cp312" -cmake.source-dir = "src" +cmake.source-dir = "." [tool.scikit-build.cmake.define] LLVM_DIR = { env = "LLVM_DIR", default = "EMPTY" } CMAKE_CXX_VISIBILITY_PRESET = "hidden" CMAKE_C_COMPILER_LAUNCHER = { env = "CMAKE_C_COMPILER_LAUNCHER", default = "" } CMAKE_CXX_COMPILER_LAUNCHER = { env = "CMAKE_CXX_COMPILER_LAUNCHER", default = "" } +CMAKE_PREFIX_PATH = { env = "CMAKE_PREFIX_PATH", default = "" } [tool.cibuildwheel] build-verbosity = 1 @@ -49,10 +49,11 @@ environment-pass = [ ] before-build = [ "export CCACHE_DIR=$CCACHE_DIR/$(python -c 'import platform; print(platform.python_version())')", - "mkdir -p $CCACHE_DIR" + "mkdir -p $CCACHE_DIR", + "ccache -z" ] # uncomment to make sure ccache is working inside containers -# test-command = "ccache -s" +test-command = "ccache -s" [tool.cibuildwheel.linux] before-all = [ @@ -62,6 +63,15 @@ before-all = [ "tar -xf ccache-4.10.2-linux-x86_64.tar.xz", "pushd ccache-4.10.2-linux-x86_64 && make install && popd" ] +# synchronize TZ with host so ccache files have correct timestamp +container-engine = { name = "docker", create-args = ["-v", "/etc/timezone:/etc/timezone:ro", "-v", "/etc/localtime:/etc/localtime:ro"] } + +[tool.cibuildwheel.macos] +before-build = [ + "ccache -z" +] [tool.cibuildwheel.windows] -before-build = [] +before-build = [ + "ccache -z" +] From acc8d3da53dc864c68c7e1093e131926ae499282 Mon Sep 17 00:00:00 2001 From: Maksim Levental Date: Thu, 2 Jan 2025 17:07:34 -0500 Subject: [PATCH 3/6] [eudsl][nbgen] eudsl-nbgen into eudsl-py working with pip wheel --- .../workflows/build_test_release_eudsl.yml | 4 +- projects/eudsl-nbgen/CMakeLists.txt | 2 +- .../cmake/eudsl-nbgen-config.cmake | 269 ------- .../cmake/eudsl_nbgen-config.cmake | 151 ++++ .../cmake/make_generated_shards.py | 18 +- projects/eudsl-py/CMakeLists.txt | 195 +++-- projects/eudsl-py/pyproject.toml | 4 +- projects/eudsl-py/src/eudslpy-gen.cpp | 722 ------------------ projects/eudsl-py/src/eudslpy_ext.cpp | 189 ++--- 9 files changed, 359 insertions(+), 1195 deletions(-) delete mode 100644 projects/eudsl-nbgen/cmake/eudsl-nbgen-config.cmake create mode 100644 projects/eudsl-nbgen/cmake/eudsl_nbgen-config.cmake delete mode 100644 projects/eudsl-py/src/eudslpy-gen.cpp diff --git a/.github/workflows/build_test_release_eudsl.yml b/.github/workflows/build_test_release_eudsl.yml index 7b00e3be..e631e829 100644 --- a/.github/workflows/build_test_release_eudsl.yml +++ b/.github/workflows/build_test_release_eudsl.yml @@ -137,7 +137,7 @@ jobs: # this is not run in cibuildwheel container so its different export CMAKE_PREFIX_PATH="$PWD/llvm-install" - $python3_command -m pip wheel "$PWD/projects/eudsl-nbgen" -w wheelhouse + $python3_command -m pip wheel "$PWD/projects/eudsl-nbgen" -w wheelhouse -v - name: "Build eudsl-py" if: ${{ ! startsWith(matrix.os, 'windows') }} @@ -146,8 +146,10 @@ jobs: if [[ "${{ matrix.os }}" == "ubuntu" ]]; then export CMAKE_PREFIX_PATH="/host/$PWD/llvm-install" export CCACHE_DIR="/host/$CCACHE_DIR" + export PIP_FIND_LINKS="/host/$PWD/wheelhouse" else export CMAKE_PREFIX_PATH="$PWD/llvm-install" + export PIP_FIND_LINKS="$PWD/wheelhouse" fi # prevent OOM on free GHA diff --git a/projects/eudsl-nbgen/CMakeLists.txt b/projects/eudsl-nbgen/CMakeLists.txt index 82936d78..b1e92720 100644 --- a/projects/eudsl-nbgen/CMakeLists.txt +++ b/projects/eudsl-nbgen/CMakeLists.txt @@ -86,7 +86,7 @@ install( install( FILES - cmake/eudsl-nbgen-config.cmake + cmake/eudsl_nbgen-config.cmake cmake/make_generated_shards.py DESTINATION "${EUDSL_NBGEN_INSTALL_DATADIR}/cmake" ) diff --git a/projects/eudsl-nbgen/cmake/eudsl-nbgen-config.cmake b/projects/eudsl-nbgen/cmake/eudsl-nbgen-config.cmake deleted file mode 100644 index fa853ee8..00000000 --- a/projects/eudsl-nbgen/cmake/eudsl-nbgen-config.cmake +++ /dev/null @@ -1,269 +0,0 @@ -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# Copyright (c) 2024. - -# copy-pasta from AddMLIR.cmake/AddLLVM.cmake/TableGen.cmake - -function(eudslpygen target inputFile) - set(EUDSLPYGEN_TARGET_DEFINITIONS ${inputFile}) - cmake_parse_arguments(ARG "" "" "DEPENDS;EXTRA_INCLUDES;NAMESPACES" ${ARGN}) - if (IS_ABSOLUTE ${EUDSLPYGEN_TARGET_DEFINITIONS}) - set(EUDSLPYGEN_TARGET_DEFINITIONS_ABSOLUTE ${inputFile}) - else() - set(EUDSLPYGEN_TARGET_DEFINITIONS_ABSOLUTE ${CMAKE_CURRENT_SOURCE_DIR}/${inputFile}) - endif() - - if(NOT EXISTS "${EUDSLPYGEN_TARGET_DEFINITIONS_ABSOLUTE}") - message(FATAL_ERROR "${inputFile} does not exist") - endif() - - get_directory_property(eudslpygen_includes INCLUDE_DIRECTORIES) - list(APPEND eudslpygen_includes ${CMAKE_CXX_IMPLICIT_INCLUDE_DIRECTORIES} ${ARG_EXTRA_INCLUDES}) - list(REMOVE_ITEM eudslpygen_includes "") - list(TRANSFORM eudslpygen_includes PREPEND -I) - - set(_gen_target_dir "${CMAKE_CURRENT_BINARY_DIR}/generated/${target}") - file(MAKE_DIRECTORY ${_gen_target_dir}) - set(fullGenFile "${_gen_target_dir}/${target}.cpp.gen") - file(RELATIVE_PATH fullGenFile_rel "${CMAKE_BINARY_DIR}" "${fullGenFile}") - set(_depfile ${fullGenFile}.d) - - # hack but most of the time we're loading headers that are downstream of tds anyway - # this could be smarter by asking people to list td targets or something but that's too onerous - file(GLOB_RECURSE global_tds "${MLIR_INCLUDE_DIR}/mlir/*.td") - # use cc -MM to collect all transitive headers - set(clang_command ${CMAKE_CXX_COMPILER} - # -v - -xc++ "-std=c++${CMAKE_CXX_STANDARD}" - -MM ${EUDSLPYGEN_TARGET_DEFINITIONS_ABSOLUTE} - -MT ${fullGenFile_rel} - ${eudslpygen_includes} - -o ${_depfile}) - execute_process(COMMAND ${clang_command} RESULT_VARIABLE had_error - # COMMAND_ECHO STDERR - ) - if(had_error OR NOT EXISTS "${_depfile}") - set(additional_cmdline -o "${fullGenFile_rel}") - else() - # Use depfile instead of globbing arbitrary *.td(s) for Ninja. - if(CMAKE_GENERATOR MATCHES "Ninja") - # Make output path relative to build.ninja, assuming located on ${CMAKE_BINARY_DIR}. - # CMake emits build targets as relative paths but Ninja doesn't identify - # absolute path (in *.d) as relative path (in build.ninja) - # Note that eudslpygen is executed on ${CMAKE_BINARY_DIR} as working directory. - set(additional_cmdline -o "${fullGenFile_rel}" DEPFILE "${_depfile}") - else() - # the length of the first line in the depfile... - string(LENGTH "${fullGenFile_rel}: \\" depfile_offset) - file(READ ${_depfile} local_headers OFFSET ${depfile_offset}) - string(REPLACE "\\" ";" local_headers "${local_headers}") - string(REGEX REPLACE "[ \t\r\n]" "" local_headers "${local_headers}") - list(REMOVE_ITEM local_headers "") - set(additional_cmdline -o "${fullGenFile_rel}") - endif() - endif() - - if (IS_ABSOLUTE ${EUDSLPYGEN_TARGET_DEFINITIONS}) - set(EUDSLPYGEN_TARGET_DEFINITIONS_ABSOLUTE ${EUDSLPYGEN_TARGET_DEFINITIONS}) - else() - set(EUDSLPYGEN_TARGET_DEFINITIONS_ABSOLUTE - ${CMAKE_CURRENT_SOURCE_DIR}/${EUDSLPYGEN_TARGET_DEFINITIONS}) - endif() - - # We need both _EUDSLPYGEN_TARGET and _EUDSLPYGEN_EXE in the DEPENDS list - # (both the target and the file) to have .inc files rebuilt on - # a eudslpygen change, as cmake does not propagate file-level dependencies - # of custom targets. See the following ticket for more information: - # https://cmake.org/Bug/view.php?id=15858 - # The dependency on both, the target and the file, produces the same - # dependency twice in the result file when - # ("${EUDSLPY_EUDSLPYGEN_TARGET}" STREQUAL "${EUDSLPY_EUDSLPYGEN_EXE}") - # but lets us having smaller and cleaner code here. - set(eudslpygen_exe ${EUDSLPY_EUDSLPYGEN_EXE}) - set(eudslpygen_depends ${EUDSLPY_EUDSLPYGEN_TARGET} ${eudslpygen_exe}) - - string(REPLACE " " ";" eudslpygen_defines "${LLVM_DEFINITIONS}") - list(JOIN ARG_NAMESPACES "," namespaces) - - add_custom_command(OUTPUT ${fullGenFile} - COMMAND ${eudslpygen_exe} ${EUDSLPYGEN_TARGET_DEFINITIONS_ABSOLUTE} - -I${CMAKE_CURRENT_SOURCE_DIR} - -namespaces=${namespaces} - ${eudslpygen_includes} - ${eudslpygen_defines} - ${additional_cmdline} - WORKING_DIRECTORY ${CMAKE_BINARY_DIR} - # The file in EUDSLPYGEN_TARGET_DEFINITIONS may be not in the current - # directory and local_headers may not contain it, so we must - # explicitly list it here: - DEPENDS ${ARG_DEPENDS} ${eudslpygen_depends} ${local_headers} ${global_tds} - COMMENT "EUDSLPY: Generating ${fullGenFile}..." - ) - # this is the specific thing connected the dependencies... - set_property(DIRECTORY APPEND PROPERTY CMAKE_CONFIGURE_DEPENDS ${fullGenFile}) - - # epic hack to specify all shards that will be generated even though we don't know them before hand - # TODO(max): refactor eudslpy-gen into its own subproject so that we can do execute_process(CMAKE_COMMAND... ) - set(_byproducts) - # lol spirv has 260 ops - set(_max_num_shards 30) - foreach(i RANGE ${_max_num_shards}) - list(APPEND _byproducts "${fullGenFile}.shard.${i}.cpp") - endforeach() - - add_custom_command(OUTPUT "${fullGenFile}.sharded.cpp" - COMMAND ${Python_EXECUTABLE} ${CMAKE_CURRENT_FUNCTION_LIST_DIR}/make_generated_registration.py - ${fullGenFile} -t ${target} -I ${ARG_EXTRA_INCLUDES} ${EUDSLPYGEN_TARGET_DEFINITIONS_ABSOLUTE} - -m ${_max_num_shards} - BYPRODUCTS ${_byproducts} - WORKING_DIRECTORY ${CMAKE_BINARY_DIR} - DEPENDS ${fullGenFile} - COMMENT "EUDSLPY: Generating ${fullGenFile}.sharded.cpp..." - ) - set_property(DIRECTORY APPEND PROPERTY CMAKE_CONFIGURE_DEPENDS "${fullGenFile}.sharded.cpp") - - add_library(${target} STATIC "${fullGenFile}.sharded.cpp" ${_byproducts}) - execute_process( - COMMAND "${Python_EXECUTABLE}" -m nanobind --include_dir - OUTPUT_STRIP_TRAILING_WHITESPACE OUTPUT_VARIABLE nanobind_include_dir) - target_include_directories(${target} PRIVATE ${eudslpygen_includes} - ${Python_INCLUDE_DIRS} ${nanobind_include_dir}) - - # `make clean' must remove all those generated files: - # TODO(max): clean up dep files - set_property(DIRECTORY ${CMAKE_CURRENT_LIST_DIR} - APPEND PROPERTY ADDITIONAL_MAKE_CLEAN_FILES ${_byproducts}) - set_source_files_properties(${_byproducts} PROPERTIES GENERATED 1) -endfunction() - -macro(add_eudslpygen target project) - set(options) - set(oneValueArgs DESTINATION EXPORT) - set(multiValueArgs) - # When used inside a macro, arg might not be a suitable prefix because the code will affect the calling scope. - cmake_parse_arguments(ARG_EUDSLPYGEN - "${options}" "${oneValueArgs}" "${multiValueArgs}" - ${ARGN} - ) - - add_llvm_executable(${target} DISABLE_LLVM_LINK_LLVM_DYLIB - ${ARG_EUDSLPYGEN_UNPARSED_ARGUMENTS} PARTIAL_SOURCES_INTENDED) - target_link_libraries(${target} - PRIVATE - clangAST - clangBasic - clangFrontend - clangSerialization - clangTooling - LLVMSupport - ) - - set(${project}_EUDSLPYGEN_DEFAULT "${target}") - if (LLVM_NATIVE_TOOL_DIR) - if (EXISTS "${LLVM_NATIVE_TOOL_DIR}/${target}${LLVM_HOST_EXECUTABLE_SUFFIX}") - set(${project}_EUDSLPYGEN_DEFAULT "${LLVM_NATIVE_TOOL_DIR}/${target}${LLVM_HOST_EXECUTABLE_SUFFIX}") - endif() - endif() - - if(ARG_EUDSLPYGEN_EXPORT) - set(${project}_EUDSLPYGEN "${${project}_EUDSLPYGEN_DEFAULT}" CACHE - STRING "Native eudslpy-gen executable. Saves building one when cross-compiling.") - else() - set(${project}_EUDSLPYGEN "${${project}_EUDSLPYGEN_DEFAULT}") - set_target_properties(${target} PROPERTIES EXCLUDE_FROM_ALL ON) - endif() - - if(PROJECT_IS_TOP_LEVEL) - set(_parent_scope) - else() - set(_parent_scope "PARENT_SCOPE") - endif() - - # Effective eudslpygen executable to be used: - set(${project}_EUDSLPYGEN_EXE ${${project}_EUDSLPYGEN} ${_parent_scope}) - set(${project}_EUDSLPYGEN_TARGET ${${project}_EUDSLPYGEN} ${_parent_scope}) - - if(LLVM_USE_HOST_TOOLS) - if( ${${project}_EUDSLPYGEN} STREQUAL "${target}" ) - # The NATIVE eudslpygen executable *must* depend on the current target one - # otherwise the native one won't get rebuilt when the tablgen sources - # change, and we end up with incorrect builds. - build_native_tool(${target} ${project}_EUDSLPYGEN_EXE DEPENDS ${target}) - set(${project}_EUDSLPYGEN_EXE ${${project}_EUDSLPYGEN_EXE} PARENT_SCOPE) - - add_custom_target(${target}-host DEPENDS ${${project}_EUDSLPYGEN_EXE}) - get_subproject_title(subproject_title) - set_target_properties(${target}-host PROPERTIES FOLDER "${subproject_title}/Native") - set(${project}_EUDSLPYGEN_TARGET ${target}-host PARENT_SCOPE) - - # If we're using the host eudslpygen, and utils were not requested, we have no - # need to build this eudslpygen. - if (NOT LLVM_BUILD_UTILS) - set_target_properties(${target} PROPERTIES EXCLUDE_FROM_ALL ON) - endif() - endif() - endif() - - if (ARG_EUDSLPYGEN_DESTINATION AND NOT LLVM_INSTALL_TOOLCHAIN_ONLY AND - (LLVM_BUILD_UTILS OR ${target} IN_LIST LLVM_DISTRIBUTION_COMPONENTS)) - set(export_arg) - if(ARG_EUDSLPYGEN_EXPORT) - get_target_export_arg(${target} ${ARG_EUDSLPYGEN_EXPORT} export_arg) - endif() - install(TARGETS ${target} - ${export_arg} - COMPONENT ${target} - RUNTIME DESTINATION "${ARG_EUDSLPYGEN_DESTINATION}") - if(NOT LLVM_ENABLE_IDE) - # TODO(max): need my own one of these... - add_llvm_install_targets("install-${target}" - DEPENDS ${target} - COMPONENT ${target}) - endif() - endif() - if(ARG_EUDSLPYGEN_EXPORT) - string(TOUPPER ${ARG_EUDSLPYGEN_EXPORT} export_upper) - set_property(GLOBAL APPEND PROPERTY ${export_upper}_EXPORTS ${target}) - endif() -endmacro() - -function(patch_mlir_llvm_rpath target) - # hack so we can move libMLIR and libLLVM into the wheel - # see AddLLVM.cmake#llvm_setup_rpath - if(APPLE OR UNIX) - set(_origin_prefix "\$ORIGIN") - if(APPLE) - set(_origin_prefix "@loader_path") - endif() - if (EUDSLPY_STANDALONE_BUILD) - get_target_property(_mlir_loc MLIR LOCATION) - get_target_property(_llvm_loc LLVM LOCATION) - else() - set(_mlir_loc "$") - set(_llvm_loc "$") - endif() - set(_old_rpath "${_origin_prefix}/../lib${LLVM_LIBDIR_SUFFIX}") - if(APPLE) - if (EXISTS ${_mlir_loc}) - execute_process(COMMAND install_name_tool -rpath "${_old_rpath}" ${_origin_prefix} "${_mlir_loc}" ERROR_VARIABLE rpath_err) - endif() - if (EXISTS ${_llvm_loc}) - execute_process(COMMAND install_name_tool -rpath "${_old_rpath}" ${_origin_prefix} "${_llvm_loc}" ERROR_VARIABLE rpath_err) - endif() - # maybe already updated... - if (rpath_err AND NOT rpath_err MATCHES "no LC_RPATH load command with path: ${_old_rpath}") - message(FATAL_ERROR "couldn't update rpath because: ${rpath_err}") - endif() - else() - # sneaky sneaky - undocumented - if (EXISTS ${_mlir_loc}) - file(RPATH_CHANGE FILE "${_mlir_loc}" OLD_RPATH "${_old_rpath}" NEW_RPATH "${_origin_prefix}") - endif() - if (EXISTS ${_llvm_loc}) - file(RPATH_CHANGE FILE "${_llvm_loc}" OLD_RPATH "${_old_rpath}" NEW_RPATH "${_origin_prefix}") - endif() - endif() - set_target_properties(${target} PROPERTIES INSTALL_RPATH "${_origin_prefix}") - endif() -endfunction() \ No newline at end of file diff --git a/projects/eudsl-nbgen/cmake/eudsl_nbgen-config.cmake b/projects/eudsl-nbgen/cmake/eudsl_nbgen-config.cmake new file mode 100644 index 00000000..bfdc4968 --- /dev/null +++ b/projects/eudsl-nbgen/cmake/eudsl_nbgen-config.cmake @@ -0,0 +1,151 @@ +# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +# See https://llvm.org/LICENSE.txt for license information. +# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +# Copyright (c) 2024. + +# copy-pasta from AddMLIR.cmake/AddLLVM.cmake/TableGen.cmake + +function(eudslpygen target inputFile) + set(EUDSLPYGEN_TARGET_DEFINITIONS ${inputFile}) + cmake_parse_arguments(ARG "" "" "DEPENDS;EXTRA_INCLUDES;NAMESPACES" ${ARGN}) + if (IS_ABSOLUTE ${EUDSLPYGEN_TARGET_DEFINITIONS}) + set(EUDSLPYGEN_TARGET_DEFINITIONS_ABSOLUTE ${inputFile}) + else() + set(EUDSLPYGEN_TARGET_DEFINITIONS_ABSOLUTE ${CMAKE_CURRENT_SOURCE_DIR}/${inputFile}) + endif() + + if(NOT EXISTS "${EUDSLPYGEN_TARGET_DEFINITIONS_ABSOLUTE}") + message(FATAL_ERROR "${inputFile} does not exist") + endif() + + get_directory_property(eudslpygen_includes INCLUDE_DIRECTORIES) + list(APPEND eudslpygen_includes ${CMAKE_CXX_IMPLICIT_INCLUDE_DIRECTORIES} ${ARG_EXTRA_INCLUDES}) + list(REMOVE_ITEM eudslpygen_includes "") + list(TRANSFORM eudslpygen_includes PREPEND -I) + + set(_gen_target_dir "${CMAKE_CURRENT_BINARY_DIR}/generated/${target}") + file(MAKE_DIRECTORY ${_gen_target_dir}) + set(_full_gen_file "${_gen_target_dir}/${target}.cpp.gen") + set(_depfile ${_full_gen_file}.d) + + # hack but most of the time we're loading headers that are downstream of tds anyway + # this could be smarter by asking people to list td targets or something but that's too onerous + file(GLOB_RECURSE global_tds "${MLIR_INCLUDE_DIR}/mlir/*.td") + # use cc -MM to collect all transitive headers + set(clang_command ${CMAKE_CXX_COMPILER} + # -v + -xc++ "-std=c++${CMAKE_CXX_STANDARD}" + -MM ${EUDSLPYGEN_TARGET_DEFINITIONS_ABSOLUTE} + -MT ${_full_gen_file} + ${eudslpygen_includes} + -o ${_depfile} + ) + execute_process(COMMAND ${clang_command} RESULT_VARIABLE _had_error_depfile + COMMAND_ECHO STDOUT + ) + + if (IS_ABSOLUTE ${EUDSLPYGEN_TARGET_DEFINITIONS}) + set(EUDSLPYGEN_TARGET_DEFINITIONS_ABSOLUTE ${EUDSLPYGEN_TARGET_DEFINITIONS}) + else() + set(EUDSLPYGEN_TARGET_DEFINITIONS_ABSOLUTE + ${CMAKE_CURRENT_SOURCE_DIR}/${EUDSLPYGEN_TARGET_DEFINITIONS}) + endif() + + string(REPLACE " " ";" eudslpygen_defines "${LLVM_DEFINITIONS}") + list(JOIN ARG_NAMESPACES "," namespaces) + + find_program(EUDSLPY_EUDSLPYGEN_EXE "eudsl-nbgen" REQUIRED) + execute_process( + COMMAND ${EUDSLPY_EUDSLPYGEN_EXE} ${EUDSLPYGEN_TARGET_DEFINITIONS_ABSOLUTE} + -I${CMAKE_CURRENT_SOURCE_DIR} + -namespaces=${namespaces} + ${eudslpygen_includes} + ${eudslpygen_defines} + -o "${_full_gen_file}" + WORKING_DIRECTORY ${CMAKE_BINARY_DIR} + RESULT_VARIABLE _had_error_gen_cpp + COMMAND_ECHO STDOUT + ) + if((_had_error_gen_cpp AND NOT _had_error_gen_cpp EQUAL 0) OR NOT EXISTS "${_full_gen_file}") + message(FATAL_ERROR "failed to create ${_full_gen_file}: ${_had_error_gen_cpp}") + endif() + # this is the specific thing connected the dependencies... + set_property(DIRECTORY APPEND PROPERTY CMAKE_CONFIGURE_DEPENDS ${_full_gen_file}) + execute_process( + COMMAND ${Python_EXECUTABLE} ${CMAKE_CURRENT_FUNCTION_LIST_DIR}/make_generated_shards.py + ${_full_gen_file} -t ${target} -I ${ARG_EXTRA_INCLUDES} ${EUDSLPYGEN_TARGET_DEFINITIONS_ABSOLUTE} + WORKING_DIRECTORY ${CMAKE_BINARY_DIR} + RESULT_VARIABLE _had_error_gen_sharded + COMMAND_ECHO STDOUT + ) + if((_had_error_gen_sharded AND NOT _had_error_gen_sharded EQUAL 0) OR NOT EXISTS "${_full_gen_file}.sharded.cpp") + message(FATAL_ERROR "failed to create ${_full_gen_file}.sharded.cpp: ${_had_error_gen_sharded}") + endif() + set_property(DIRECTORY APPEND PROPERTY CMAKE_CONFIGURE_DEPENDS "${_full_gen_file}.sharded.cpp") + file(GLOB _shards CONFIGURE_DEPENDS "${_gen_target_dir}/*shard*cpp") + if(NOT _shards) + message(FATAL_ERROR "no shards created") + endif() + + add_library(${target} STATIC "${_full_gen_file}.sharded.cpp" ${_shards}) + execute_process( + COMMAND "${Python_EXECUTABLE}" -m nanobind --include_dir + OUTPUT_STRIP_TRAILING_WHITESPACE OUTPUT_VARIABLE nanobind_include_dir + RESULT_VARIABLE _has_err_find_nanobind + ) + if((_has_err_find_nanobind AND NOT _has_err_find_nanobind EQUAL 0) OR NOT EXISTS "${nanobind_include_dir}") + message(FATAL_ERROR "couldn't find nanobind include dir: ${_has_err_find_nanobind}") + endif() + target_include_directories(${target} PRIVATE + ${eudslpygen_includes} + ${Python_INCLUDE_DIRS} + ${nanobind_include_dir} + ) + target_compile_options(${target} PRIVATE -Wno-cast-qual) + + # `make clean' must remove all those generated files: + # TODO(max): clean up dep files + set_property(DIRECTORY ${CMAKE_CURRENT_LIST_DIR} + APPEND PROPERTY ADDITIONAL_MAKE_CLEAN_FILES ${_shards}) + set_source_files_properties(${_shards} PROPERTIES GENERATED 1) +endfunction() + +function(patch_mlir_llvm_rpath target) + # hack so we can move libMLIR and libLLVM into the wheel + # see AddLLVM.cmake#llvm_setup_rpath + if(APPLE OR UNIX) + set(_origin_prefix "\$ORIGIN") + if(APPLE) + set(_origin_prefix "@loader_path") + endif() + if (EUDSLPY_STANDALONE_BUILD) + get_target_property(_mlir_loc MLIR LOCATION) + get_target_property(_llvm_loc LLVM LOCATION) + else() + set(_mlir_loc "$") + set(_llvm_loc "$") + endif() + set(_old_rpath "${_origin_prefix}/../lib${LLVM_LIBDIR_SUFFIX}") + if(APPLE) + if (EXISTS ${_mlir_loc}) + execute_process(COMMAND install_name_tool -rpath "${_old_rpath}" ${_origin_prefix} "${_mlir_loc}" ERROR_VARIABLE rpath_err) + endif() + if (EXISTS ${_llvm_loc}) + execute_process(COMMAND install_name_tool -rpath "${_old_rpath}" ${_origin_prefix} "${_llvm_loc}" ERROR_VARIABLE rpath_err) + endif() + # maybe already updated... + if (rpath_err AND NOT rpath_err MATCHES "no LC_RPATH load command with path: ${_old_rpath}") + message(FATAL_ERROR "couldn't update rpath because: ${rpath_err}") + endif() + else() + # sneaky sneaky - undocumented + if (EXISTS ${_mlir_loc}) + file(RPATH_CHANGE FILE "${_mlir_loc}" OLD_RPATH "${_old_rpath}" NEW_RPATH "${_origin_prefix}") + endif() + if (EXISTS ${_llvm_loc}) + file(RPATH_CHANGE FILE "${_llvm_loc}" OLD_RPATH "${_old_rpath}" NEW_RPATH "${_origin_prefix}") + endif() + endif() + set_target_properties(${target} PROPERTIES INSTALL_RPATH "${_origin_prefix}") + endif() +endfunction() \ No newline at end of file diff --git a/projects/eudsl-nbgen/cmake/make_generated_shards.py b/projects/eudsl-nbgen/cmake/make_generated_shards.py index 0a2d63f2..207fe4f4 100644 --- a/projects/eudsl-nbgen/cmake/make_generated_shards.py +++ b/projects/eudsl-nbgen/cmake/make_generated_shards.py @@ -5,7 +5,6 @@ import argparse import re -import sys from pathlib import Path from textwrap import dedent @@ -39,11 +38,12 @@ def make_source_shards(filename: Path, target, extra_includes, max_num_shards): print(shar, file=f) print("}", file=f) - if len(shards) > max_num_shards: - raise RuntimeError("expected less than 20 shards") - for i in range(len(shards), max_num_shards): - with open(f"{filename}.shard.{i}.cpp", "w") as f: - print(f"// dummy shard {i}", file=f) + if max_num_shards is not None: + if len(shards) > max_num_shards: + raise RuntimeError(f"expected less than {max_num_shards} shards") + for i in range(len(shards), max_num_shards): + with open(f"{filename}.shard.{i}.cpp", "w") as f: + print(f"// dummy shard {i}", file=f) with open(f"{filename}.sharded.cpp", "w") as f: print( @@ -72,8 +72,10 @@ def make_source_shards(filename: Path, target, extra_includes, max_num_shards): parser.add_argument("filename") parser.add_argument("-t", "--target") parser.add_argument("-I", "--extra_includes", nargs="*") - parser.add_argument("-m", "--max-num-shards", type=int, default=20) + parser.add_argument("-m", "--max-num-shards", type=int) args = parser.parse_args() + if args.max_num_shards: + args.max_num_shards += 1 make_source_shards( - Path(args.filename), args.target, args.extra_includes, args.max_num_shards + 1 + Path(args.filename), args.target, args.extra_includes, args.max_num_shards ) diff --git a/projects/eudsl-py/CMakeLists.txt b/projects/eudsl-py/CMakeLists.txt index 245784c6..c212ea2b 100644 --- a/projects/eudsl-py/CMakeLists.txt +++ b/projects/eudsl-py/CMakeLists.txt @@ -102,24 +102,20 @@ set(nanobind_options ) set(EUDSLPY_SRC_DIR "${CMAKE_CURRENT_LIST_DIR}/src") -set(EUDSLPY_DIR "${CMAKE_CURRENT_LIST_DIR}/cmake") include_directories(${EUDSLPY_BINARY_DIR}) include_directories(${EUDSLPY_SRC_DIR}) -find_package(EUDSLPY CONFIG REQUIRED) - -add_eudslpygen(eudslpy-gen EUDSLPY - DESTINATION "${CMAKE_INSTALL_BINDIR}" - EXPORT ${LLVM_SUBPROJECT_TITLE} - src/eudslpy-gen.cpp -) +execute_process( + COMMAND "${Python_EXECUTABLE}" -m eudsl_nbgen --cmake_dir + OUTPUT_STRIP_TRAILING_WHITESPACE OUTPUT_VARIABLE eudsl_nbgen_DIR) +find_package(eudsl_nbgen CONFIG REQUIRED) # too big -# eudslpygen(EUDSLGenacc +# eudslpygen(EUDSLGen_acc # ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h # NAMESPACES mlir::acc mlir::acc::detail # ) -eudslpygen(EUDSLGenaffine +eudslpygen(EUDSLGen_affine ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::affine mlir::affine::detail EXTRA_INCLUDES @@ -127,50 +123,50 @@ eudslpygen(EUDSLGenaffine mlir/Dialect/Affine/IR/AffineValueMap.h ) -eudslpygen(EUDSLGenamdgpu +eudslpygen(EUDSLGen_amdgpu ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::amdgpu mlir::amdgpu::detail ) -eudslpygen(EUDSLGenamx +eudslpygen(EUDSLGen_amx ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::amx mlir::amx::detail ) -eudslpygen(EUDSLGenarith +eudslpygen(EUDSLGen_arith ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::arith mlir::arith::detail ) -eudslpygen(EUDSLGenarm_neon +eudslpygen(EUDSLGen_arm_neon ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::arm_neon mlir::arm_neon::detail ) # too big -# eudslpygen(EUDSLGenarm_sme +# eudslpygen(EUDSLGen_arm_sme # ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h # NAMESPACES mlir::arm_sme mlir::arm_sme::detail # ) -eudslpygen(EUDSLGenarm_sve +eudslpygen(EUDSLGen_arm_sve ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::arm_sve mlir::arm_sve::detail ) -eudslpygen(EUDSLGenasync +eudslpygen(EUDSLGen_async ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::async mlir::async::detail ) -eudslpygen(EUDSLGenbufferization +eudslpygen(EUDSLGen_bufferization ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::bufferization mlir::bufferization::detail EXTRA_INCLUDES mlir/Dialect/Bufferization/Transforms/Bufferize.h ) -eudslpygen(EUDSLGencf +eudslpygen(EUDSLGen_cf ${MLIR_INCLUDE_DIR}/mlir/Dialect/ControlFlow/IR/ControlFlowOps.h NAMESPACES mlir::cf mlir::cf::detail EXTRA_INCLUDES @@ -178,27 +174,27 @@ eudslpygen(EUDSLGencf mlir/IR/PatternMatch.h ) -eudslpygen(EUDSLGencomplex +eudslpygen(EUDSLGen_complex ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::complex mlir::complex::detail ) -eudslpygen(EUDSLGenDLTIDialect +eudslpygen(EUDSLGen_DLTIDialect ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::DLTIDialect mlir::DLTIDialect::detail ) -eudslpygen(EUDSLGenemitc +eudslpygen(EUDSLGen_emitc ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::emitc mlir::emitc::detail ) -eudslpygen(EUDSLGenfunc +eudslpygen(EUDSLGen_func ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::func mlir::func::detail ) -eudslpygen(EUDSLGengpu +eudslpygen(EUDSLGen_gpu ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::gpu mlir::gpu::detail EXTRA_INCLUDES @@ -206,7 +202,7 @@ eudslpygen(EUDSLGengpu llvm/IR/IRBuilder.h ) -eudslpygen(EUDSLGenindex +eudslpygen(EUDSLGen_index ${MLIR_INCLUDE_DIR}/mlir/Dialect/Index/IR/IndexOps.h NAMESPACES mlir::index mlir::index::detail EXTRA_INCLUDES @@ -214,76 +210,76 @@ eudslpygen(EUDSLGenindex mlir/IR/PatternMatch.h ) -#eudslpygen(EUDSLGenirdl +#eudslpygen(EUDSLGen_irdl # ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h # NAMESPACES mlir::irdl mlir::irdl::detail #) -eudslpygen(EUDSLGenlinalg +eudslpygen(EUDSLGen_linalg ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::linalg mlir::linalg::detail ) -eudslpygen(EUDSLGenLLVM +eudslpygen(EUDSLGen_LLVM ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::LLVM mlir::LLVM::detail ) -eudslpygen(EUDSLGenmath +eudslpygen(EUDSLGen_math ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::math mlir::math::detail ) -eudslpygen(EUDSLGenmemref +eudslpygen(EUDSLGen_memref ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::memref mlir::memref::detail ) -eudslpygen(EUDSLGenmesh +eudslpygen(EUDSLGen_mesh ${MLIR_INCLUDE_DIR}/mlir/Dialect/Mesh/IR/MeshOps.h NAMESPACES mlir::mesh mlir::mesh::detail EXTRA_INCLUDES mlir/Dialect/Mesh/IR/MeshOps.h ) -eudslpygen(EUDSLGenml_program +eudslpygen(EUDSLGen_ml_program ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::ml_program mlir::ml_program::detail ) -eudslpygen(EUDSLGenmpi +eudslpygen(EUDSLGen_mpi ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::mpi mlir::mpi::detail ) -eudslpygen(EUDSLGennvgpu +eudslpygen(EUDSLGen_nvgpu ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::nvgpu mlir::nvgpu::detail ) -eudslpygen(EUDSLGenNVVM +eudslpygen(EUDSLGen_NVVM ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::NVVM mlir::NVVM::detail ) -#eudslpygen(EUDSLGenomp +#eudslpygen(EUDSLGen_omp # ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h # NAMESPACES mlir::omp mlir::omp::detail #) -eudslpygen(EUDSLGenpdl +eudslpygen(EUDSLGen_pdl ${MLIR_INCLUDE_DIR}/mlir/Dialect/PDL/IR/PDLOps.h NAMESPACES mlir::pdl mlir::pdl::detail EXTRA_INCLUDES mlir/Dialect/PDL/IR/PDLOps.h ) -eudslpygen(EUDSLGenpdl_interp +eudslpygen(EUDSLGen_pdl_interp ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::pdl_interp mlir::pdl_interp::detail ) -eudslpygen(EUDSLGenpolynomial +eudslpygen(EUDSLGen_polynomial ${MLIR_INCLUDE_DIR}/mlir/Dialect/Polynomial/IR/PolynomialOps.h NAMESPACES mlir::polynomial mlir::polynomial::detail EXTRA_INCLUDES @@ -291,7 +287,7 @@ eudslpygen(EUDSLGenpolynomial mlir/IR/PatternMatch.h ) -eudslpygen(EUDSLGenptr +eudslpygen(EUDSLGen_ptr ${MLIR_INCLUDE_DIR}/mlir/Dialect/Ptr/IR/PtrOps.h NAMESPACES mlir::ptr mlir::ptr::detail EXTRA_INCLUDES @@ -299,35 +295,35 @@ eudslpygen(EUDSLGenptr mlir/IR/DialectImplementation.h ) -eudslpygen(EUDSLGenquant +eudslpygen(EUDSLGen_quant ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::quant mlir::quant::detail EXTRA_INCLUDES mlir/Dialect/Quant/IR/QuantTypes.h ) -eudslpygen(EUDSLGenROCDL +eudslpygen(EUDSLGen_ROCDL ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::ROCDL mlir::ROCDL::detail ) -eudslpygen(EUDSLGenscf +eudslpygen(EUDSLGen_scf ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::scf mlir::scf::detail ) -eudslpygen(EUDSLGenshape +eudslpygen(EUDSLGen_shape ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::shape mlir::shape::detail ) -eudslpygen(EUDSLGensparse_tensor +eudslpygen(EUDSLGen_sparse_tensor ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::sparse_tensor mlir::sparse_tensor::detail ) # nb::detail::nb_func_new("get_vce_triple_attr_name"): mismatched static/instance method flags in function overloads! -# eudslpygen(EUDSLGenspirv +# eudslpygen(EUDSLGen_spirv # ${MLIR_INCLUDE_DIR}/mlir/Dialect/SPIRV/IR/SPIRVOps.h # NAMESPACES mlir::spirv mlir::spirv::detail # EXTRA_INCLUDES @@ -335,37 +331,37 @@ eudslpygen(EUDSLGensparse_tensor # mlir/IR/PatternMatch.h # ) -eudslpygen(EUDSLGentensor +eudslpygen(EUDSLGen_tensor ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::tensor mlir::tensor::detail ) -eudslpygen(EUDSLGentosa +eudslpygen(EUDSLGen_tosa ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::tosa mlir::tosa::detail ) -eudslpygen(EUDSLGentransform +eudslpygen(EUDSLGen_transform ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::transform mlir::transform::detail ) -eudslpygen(EUDSLGenub +eudslpygen(EUDSLGen_ub ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::ub mlir::ub::detail ) -#eudslpygen(EUDSLGenvector +#eudslpygen(EUDSLGen_vector # ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h # NAMESPACES mlir::vector mlir::vector::detail #) -eudslpygen(EUDSLGenx86vector +eudslpygen(EUDSLGen_x86vector ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::x86vector mlir::x86vector::detail ) -eudslpygen(EUDSLGenxegpu +eudslpygen(EUDSLGen_xegpu ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::xegpu mlir::xegpu::detail ) @@ -377,53 +373,53 @@ nanobind_add_module(eudslpy_ext ) target_link_libraries(eudslpy_ext PRIVATE # broken - # EUDSLGenirdl - # EUDSLGenomp - # EUDSLGenvector + # EUDSLGen_irdl + # EUDSLGen_omp + # EUDSLGen_vector # weight reduction - # EUDSLGenacc - # EUDSLGenamx - # EUDSLGenarm_neon - # EUDSLGenarm_sme - # EUDSLGenarm_sve - # EUDSLGenDLTIDialect - # EUDSLGenmesh - # EUDSLGenml_program - # EUDSLGenmpi - # EUDSLGenptr - # EUDSLGenquant - # EUDSLGensparse_tensor - # EUDSLGenspirv - # EUDSLGentransform - # EUDSLGenub - # EUDSLGenx86vector - # EUDSLGenxegpu - - EUDSLGenaffine - EUDSLGenamdgpu - EUDSLGenarith - EUDSLGenasync - EUDSLGenbufferization - EUDSLGencf - EUDSLGencomplex - EUDSLGenemitc - EUDSLGenfunc - EUDSLGengpu - EUDSLGenindex - EUDSLGenlinalg - EUDSLGenLLVM - EUDSLGenmath - EUDSLGenmemref - EUDSLGennvgpu - EUDSLGenNVVM - EUDSLGenpdl - EUDSLGenpdl_interp - EUDSLGenpolynomial - EUDSLGenROCDL - EUDSLGenscf - EUDSLGenshape - EUDSLGentensor - EUDSLGentosa + # EUDSLGen_acc + # EUDSLGen_amx + # EUDSLGen_arm_neon + # EUDSLGen_arm_sme + # EUDSLGen_arm_sve + # EUDSLGen_DLTIDialect + # EUDSLGen_mesh + # EUDSLGen_ml_program + # EUDSLGen_mpi + # EUDSLGen_ptr + # EUDSLGen_quant + # EUDSLGen_sparse_tensor + # EUDSLGen_spirv + # EUDSLGen_transform + # EUDSLGen_ub + # EUDSLGen_x86vector + # EUDSLGen_xegpu + + EUDSLGen_affine + EUDSLGen_amdgpu + EUDSLGen_arith + EUDSLGen_async + EUDSLGen_bufferization + EUDSLGen_cf + EUDSLGen_complex + EUDSLGen_emitc + EUDSLGen_func + EUDSLGen_gpu + EUDSLGen_index + EUDSLGen_linalg + EUDSLGen_LLVM + EUDSLGen_math + EUDSLGen_memref + EUDSLGen_nvgpu + EUDSLGen_NVVM + EUDSLGen_pdl + EUDSLGen_pdl_interp + EUDSLGen_polynomial + EUDSLGen_ROCDL + EUDSLGen_scf + EUDSLGen_shape + EUDSLGen_tensor + EUDSLGen_tosa ) set_target_properties(eudslpy_ext PROPERTIES @@ -461,5 +457,6 @@ install( PATTERN "*.pyc" EXCLUDE PATTERN "*.so" EXCLUDE PATTERN "*.a" EXCLUDE + PATTERN "__pycache__" EXCLUDE PATTERN ".gitignore" EXCLUDE ) diff --git a/projects/eudsl-py/pyproject.toml b/projects/eudsl-py/pyproject.toml index c2070deb..60fb1786 100644 --- a/projects/eudsl-py/pyproject.toml +++ b/projects/eudsl-py/pyproject.toml @@ -8,7 +8,8 @@ requires = [ "scikit-build-core==0.10.7", "nanobind==2.4.0", "typing_extensions==4.12.2", - "numpy==2.0.2" + "numpy==2.0.2", + "eudsl-nbgen" ] build-backend = "scikit_build_core.build" @@ -53,6 +54,7 @@ environment-pass = [ "DISABLE_COMPILE_OPT", "CC", "CXX", + "PIP_FIND_LINKS", # ccache "CCACHE_DIR", "CCACHE_MAXSIZE=700M", diff --git a/projects/eudsl-py/src/eudslpy-gen.cpp b/projects/eudsl-py/src/eudslpy-gen.cpp deleted file mode 100644 index 0b65ab64..00000000 --- a/projects/eudsl-py/src/eudslpy-gen.cpp +++ /dev/null @@ -1,722 +0,0 @@ -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// Copyright (c) 2024. - -#include "clang/AST/ASTConsumer.h" -#include "clang/AST/LexicallyOrderedRecursiveASTVisitor.h" -#include "clang/AST/Mangle.h" -#include "clang/AST/RecursiveASTVisitor.h" -#include "clang/Frontend/CompilerInstance.h" -#include "clang/Frontend/FrontendAction.h" -#include "clang/Sema/Lookup.h" -#include "clang/Sema/Sema.h" -#include "clang/Tooling/Tooling.h" -#include "llvm/Support/FormatVariadic.h" -#include "llvm/Support/ToolOutputFile.h" - -#include - -static llvm::cl::OptionCategory EUDSLPYGenCat("Options for eudslpy-gen"); -static llvm::cl::opt InputFilename(llvm::cl::Positional, - llvm::cl::desc(""), - llvm::cl::Required, - llvm::cl::cat(EUDSLPYGenCat)); - -static llvm::cl::list - IncludeDirs("I", llvm::cl::desc("Directory of include files"), - llvm::cl::value_desc("directory"), llvm::cl::Prefix, - llvm::cl::cat(EUDSLPYGenCat)); - -static llvm::cl::opt - OutputFilename("o", llvm::cl::desc("Output filename"), - llvm::cl::value_desc("filename"), llvm::cl::Required, - llvm::cl::cat(EUDSLPYGenCat)); - -static llvm::cl::list - Namespaces("namespaces", llvm::cl::desc("Namespaces to generate from"), - llvm::cl::CommaSeparated, llvm::cl::cat(EUDSLPYGenCat)); - -static llvm::cl::list - Defines("D", llvm::cl::desc("Name of the macro to be defined"), - llvm::cl::value_desc("macro name"), llvm::cl::Prefix, - llvm::cl::cat(EUDSLPYGenCat)); - -static llvm::cl::opt ShardSize("shard-size", llvm::cl::desc("Shard size"), - llvm::cl::value_desc("shard size"), - llvm::cl::cat(EUDSLPYGenCat), - llvm::cl::init(10)); - -static bool filterInNamespace(const std::string &s) { - if (Namespaces.empty()) - return true; - for (auto ns : Namespaces) - if (ns == s || ("::" + ns == s)) - return true; - return false; -} - -static std::string -getNBBindClassName(const std::string &qualifiedNameAsString) { - std::string s = qualifiedNameAsString; - s = std::regex_replace(s, std::regex(R"(\s+)"), ""); - s = std::regex_replace(s, std::regex("::"), "_"); - s = std::regex_replace(s, std::regex("[<|>]"), "__"); - s = std::regex_replace(s, std::regex(R"(\*)"), "___"); - s = std::regex_replace(s, std::regex(","), "____"); - return s; -} - -static std::string getPyClassName(const std::string &qualifiedNameAsString) { - std::string s = qualifiedNameAsString; - s = std::regex_replace(s, std::regex(R"(\s+)"), ""); - s = std::regex_replace(s, std::regex(R"(\*)"), ""); - s = std::regex_replace(s, std::regex("<"), "["); - s = std::regex_replace(s, std::regex(">"), "]"); - return s; -} - -static std::string snakeCase(const std::string &name) { - std::string s = name; - s = std::regex_replace(s, std::regex(R"(([A-Z]+)([A-Z][a-z]))"), "$1_$2"); - s = std::regex_replace(s, std::regex(R"(([a-z\d])([A-Z]))"), "$1_$2"); - s = std::regex_replace(s, std::regex("-"), "_"); - std::transform(s.begin(), s.end(), s.begin(), - [](unsigned char c) { return std::tolower(c); }); - return s; -} - -static clang::PrintingPolicy getPrintingPolicy(bool canonical = true) { - clang::LangOptions lo; - lo.CPlusPlus = true; - clang::PrintingPolicy p(lo); - // TODO(max): none of this really does anything except PrintCanonical - // keep namespaces - p.FullyQualifiedName = true; - p.SuppressScope = false; - p.SuppressInlineNamespace = false; - p.PrintCanonicalTypes = canonical; - p.Bool = true; - - return p; -} - -// stolen from -// https://github.com/llvm/llvm-project/blob/99dddef340e566e9d303010f1219f7d7d6d37a11/clang/lib/Sema/SemaChecking.cpp#L7055 -// Determines if the specified is a C++ class or struct containing -// a member with the specified name and kind (e.g. a CXXMethodDecl named -// "c_str()"). -template -static llvm::SmallPtrSet findOverloads(clang::FunctionDecl *decl, - clang::Sema &s) { - llvm::SmallPtrSet results; - clang::LookupResult r(s, &s.Context.Idents.get(decl->getNameAsString()), - decl->getLocation(), clang::Sema::LookupOrdinaryName); - r.suppressDiagnostics(); - if (s.LookupQualifiedName(r, decl->getDeclContext())) - for (clang::LookupResult::iterator i = r.begin(), e = r.end(); i != e; - ++i) { - clang::NamedDecl *namedDecl = (*i)->getUnderlyingDecl(); - if (T *fk = llvm::dyn_cast(namedDecl)) - results.insert(fk); - } - return results; -} - -// TODO(max): split this into two functions (one for names and one for types) -static std::string sanitizeNameOrType(std::string nameOrType, - int emptyIdx = 0) { - if (nameOrType == "from") - nameOrType = "from_"; - else if (nameOrType == "except") - nameOrType = "except_"; - else if (nameOrType == "") - nameOrType = std::string(emptyIdx + 1, '_'); - else if (nameOrType.rfind("ArrayRef", 0) == 0) - nameOrType = "llvm::" + nameOrType; - if (std::regex_search(nameOrType, std::regex(R"(std::__1)"))) - nameOrType = std::regex_replace(nameOrType, std::regex("std::__1"), "std"); - return nameOrType; -} - -// emit a lambda body to disambiguate/break ties amongst overloads -// TODO(max):: overloadimpl or whatever should work but it doesn't... -std::string emitNBLambdaBody(clang::FunctionDecl *decl, - llvm::SmallVector paramNames, - llvm::SmallVector paramTypes) { - std::string typedParamsStr; - if (decl->getNumParams()) { - llvm::SmallVector typedParams = - llvm::to_vector(llvm::map_range( - llvm::zip(paramTypes, paramNames), - [](std::tuple item) -> std::string { - auto [t, n] = item; - return llvm::formatv("{0} {1}", t, n); - })); - typedParamsStr = llvm::join(typedParams, ", "); - } - // since we're emitting a body, we need to do std::move for some - // unique_ptrs - llvm::SmallVector newParamNames(paramNames); - for (auto [idx, item] : - llvm::enumerate(llvm::zip(paramTypes, newParamNames))) { - // TODO(max): characterize this condition better... - auto [t, n] = item; - if ((t.rfind("std::unique_ptr") != std::string::npos && t.back() != '&') || - t.rfind("&&") != std::string::npos) - n = llvm::formatv("std::move({0})", n); - } - std::string newParamNamesStr = llvm::join(newParamNames, ", "); - std::string funcRef; - std::string returnType = sanitizeNameOrType( - decl->getReturnType().getAsString(getPrintingPolicy())); - if (decl->isStatic() || !decl->isCXXClassMember()) { - funcRef = llvm::formatv("\n []({0}) -> {1} {{\n return {2}({3});\n }", - typedParamsStr, returnType, - decl->getQualifiedNameAsString(), newParamNamesStr); - } else { - assert(decl->isCXXClassMember() && "expected class member"); - if (decl->getNumParams()) - typedParamsStr = llvm::formatv("self, {0}", typedParamsStr); - else - typedParamsStr = "self"; - const clang::CXXRecordDecl *parentRecord = - llvm::cast(decl->getParent()); - funcRef = llvm::formatv( - "\n []({0}& {1}) -> {2} {{\n return self.{3}({4});\n }", - parentRecord->getQualifiedNameAsString(), typedParamsStr, returnType, - decl->getNameAsString(), newParamNamesStr); - } - return funcRef; -} - -static bool -emitClassMethodOrFunction(clang::FunctionDecl *decl, - clang::CompilerInstance &ci, - std::shared_ptr outputFile) { - llvm::SmallVector paramNames; - llvm::SmallVector paramTypes; - for (unsigned i = 0; i < decl->getNumParams(); ++i) { - clang::ParmVarDecl *param = decl->getParamDecl(i); - std::string name = param->getNameAsString(); - auto t = param->getType(); - bool canonical = true; - // TODO(max): this is dumb... (maybe there's a way to check where the - // typedef is defined...) - // word boundary excludes x86_amx_tdpbf16ps - if (std::regex_search(t.getAsString(), std::regex(R"(_t\b)"))) - canonical = false; - std::string paramType = t.getAsString(getPrintingPolicy(canonical)); - paramTypes.push_back(sanitizeNameOrType(paramType)); - paramNames.push_back(sanitizeNameOrType(name, i)); - } - - llvm::SmallPtrSet funcOverloads = - findOverloads(decl, ci.getSema()); - llvm::SmallPtrSet funcTemplOverloads = - findOverloads(decl, ci.getSema()); - std::string funcRef, nbFnName; - - if (auto ctor = llvm::dyn_cast(decl)) { - if (ctor->isDeleted()) - return false; - funcRef = llvm::formatv("nb::init<{0}>()", llvm::join(paramTypes, ", ")); - } else { - if (funcOverloads.size() == 1 && funcTemplOverloads.empty()) { - funcRef = llvm::formatv("&{0}", decl->getQualifiedNameAsString()); - } else { - funcRef = emitNBLambdaBody(decl, paramNames, paramTypes); - } - - nbFnName = snakeCase(decl->getNameAsString()); - if (decl->isOverloadedOperator()) { - // TODO(max): handle overloaded operators - // nbFnName = nbFnName; - } else if (decl->isStatic() && funcOverloads.size() > 1 && - llvm::any_of(funcOverloads, [](clang::FunctionDecl *m) { - return !m->isStatic(); - })) { - // disambiguate static method with non-static overloads (nanobind doesn't - // let you overload static with non-static) see mlir::ElementsAttr - nbFnName += "_static"; - } - } - - std::string paramNamesStr; - if (decl->getNumParams()) { - paramNames = llvm::to_vector( - llvm::map_range(paramNames, [](std::string s) -> std::string { - return llvm::formatv("\"{0}\"_a", snakeCase(s)); - })); - paramNamesStr = ", " + llvm::join(paramNames, ", "); - } - - std::string refInternal, defStr = "def"; - if (decl->isCXXClassMember()) { - if (decl->isStatic()) - defStr = "def_static"; - if (decl->getReturnType()->hasPointerRepresentation()) - refInternal = ", nb::rv_policy::reference_internal"; - } - - std::string sig; - if (!nbFnName.empty()) { - // no clue why but nb has trouble inferring the signature - // (and this causes and assert failure in nb_func_render_signature - // https://github.com/wjakob/nanobind/blob/c2e394eee5d19816871151de43c29b4051fbf9ff/src/nb_func.cpp#L1020 - if (decl->isStatic() && !decl->getNumParams()) - sig = - llvm::formatv(", nb::sig(\"def {0}(/) -> {1}\") ", nbFnName, - decl->getReturnType().getAsString(getPrintingPolicy())); - // uncomment this to debug nb which doesn't tell where the method actually - // is when it fails if (parentRecord) { - // nbFnName += parentRecord->getNameAsString(); - // } - nbFnName = llvm::formatv("\"{0}\", ", nbFnName); - } - - std::string scope = "m"; - if (decl->isCXXClassMember()) { - const clang::CXXRecordDecl *parentRecord = - llvm::cast(decl->getParent()); - scope = getNBBindClassName(parentRecord->getQualifiedNameAsString()); - } - - outputFile->os() << llvm::formatv("{0}.{1}({2}{3}{4}{5}{6});\n", scope, - defStr, nbFnName, funcRef, paramNamesStr, - refInternal, sig); - - return true; -} - -std::string getNBScope(clang::TagDecl *decl) { - std::string scope = "m"; - const clang::DeclContext *declContext = decl->getDeclContext(); - if (declContext->isRecord()) { - const clang::CXXRecordDecl *ctx = - llvm::cast(declContext); - scope = getNBBindClassName(ctx->getQualifiedNameAsString()); - } - return scope; -} - -static bool emitClass(clang::CXXRecordDecl *decl, clang::CompilerInstance &ci, - std::shared_ptr outputFile) { - if (decl->isTemplated()) { - clang::DiagnosticBuilder builder = ci.getDiagnostics().Report( - decl->getLocation(), ci.getDiagnostics().getCustomDiagID( - clang::DiagnosticsEngine::Note, - "template classes not supported yet")); - return false; - } - - std::string scope = getNBScope(decl); - std::string additional = ""; - std::string className = decl->getQualifiedNameAsString(); - if (decl->getNumBases() > 1) { - clang::DiagnosticBuilder builder = ci.getDiagnostics().Report( - decl->getLocation(), ci.getDiagnostics().getCustomDiagID( - clang::DiagnosticsEngine::Note, - "multiple base classes not supported")); - } else if (decl->getNumBases() == 1) { - // handle some known bases that we've already found a wap to bind - clang::CXXBaseSpecifier baseClass = *decl->bases_begin(); - std::string baseName = baseClass.getType().getAsString(getPrintingPolicy()); - // TODO(max): these could be lookups on the corresponding recorddecls using - // sema... - if (baseName.rfind("mlir::Op<", 0) == 0) { - className = llvm::formatv("{0}, mlir::OpState", className); - } else if (baseName.rfind("mlir::detail::StorageUserBase<", 0) == 0) { - llvm::SmallVector templParams; - llvm::StringRef{baseName}.split(templParams, ","); - className = llvm::formatv("{0}, {1}", className, templParams[1]); - } else if (baseName.rfind("mlir::Dialect", 0) == 0 && - className.rfind("mlir::ExtensibleDialect") == - std::string::npos) { - // clang-format off - additional += llvm::formatv("\n .def_static(\"insert_into_registry\", [](mlir::DialectRegistry ®istry) {{ registry.insert<{0}>(); })", className); - additional += llvm::formatv("\n .def_static(\"load_into_context\", [](mlir::MLIRContext &context) {{ return context.getOrLoadDialect<{0}>(); })", className); - // clang-format on - } else if (!llvm::isa( - baseClass.getType()->getAsCXXRecordDecl())) { - className = llvm::formatv("{0}, {1}", className, baseName); - } else { - assert(llvm::isa( - baseClass.getType()->getAsCXXRecordDecl()) && - "expected class template specialization"); - clang::DiagnosticBuilder builder = ci.getDiagnostics().Report( - baseClass.getBeginLoc(), ci.getDiagnostics().getCustomDiagID( - clang::DiagnosticsEngine::Note, - "unknown base templated base class: ")); - builder << baseName << "\n"; - } - } - - std::string autoVar = llvm::formatv( - "auto {0}", getNBBindClassName(decl->getQualifiedNameAsString())); - - outputFile->os() << llvm::formatv( - "\n{0} = nb::class_<{1}>({2}, \"{3}\"){4};\n", autoVar, className, scope, - getPyClassName(decl->getNameAsString()), additional); - - return true; -} - -static bool emitEnum(clang::EnumDecl *decl, clang::CompilerInstance &ci, - std::shared_ptr outputFile) { - outputFile->os() << llvm::formatv("nb::enum_<{0}>({1}, \"{2}\")\n", - decl->getQualifiedNameAsString(), - getNBScope(decl), decl->getNameAsString()); - - int i = 0, nDecls = std::distance(decl->decls_begin(), decl->decls_end()); - for (clang::Decl *cst : decl->decls()) { - clang::EnumConstantDecl *cstDecl = llvm::cast(cst); - outputFile->os() << llvm::formatv(" .value(\"{0}\", {1})", - cstDecl->getNameAsString(), - cstDecl->getQualifiedNameAsString()); - if (i++ < nDecls - 1) - outputFile->os() << "\n"; - else - outputFile->os() << ";\n"; - } - outputFile->os() << "\n"; - return true; -} - -static bool emitField(clang::DeclaratorDecl *field, clang::CompilerInstance &ci, - std::shared_ptr outputFile) { - const clang::CXXRecordDecl *parentRecord = - llvm::cast(field->getLexicalDeclContext()); - - std::string defStr = "def_rw"; - if (clang::VarDecl *vard = llvm::dyn_cast(field)) { - if (vard->isStaticDataMember()) { - if (vard->isConstexpr()) - defStr = "def_ro_static"; - else - defStr = "def_rw_static"; - } - } - - std::string refInternal; - if (field->getType()->hasPointerRepresentation()) - refInternal = ", nb::rv_policy::reference_internal"; - - std::string scope = - getNBBindClassName(parentRecord->getQualifiedNameAsString()); - std::string nbFnName = - llvm::formatv("\"{0}\"", snakeCase(field->getNameAsString())); - outputFile->os() << llvm::formatv("{0}.{1}({2}, &{3}{4});\n", scope, defStr, - nbFnName, field->getQualifiedNameAsString(), - refInternal); - return true; -} - -template -static bool shouldSkip(T *decl) { - auto *encl = llvm::dyn_cast( - decl->getEnclosingNamespaceContext()); - if (!encl) - return true; - // this isn't redundant - filter might be empty but we still don't want to - // bind std:: - if (encl->isStdNamespace() || encl->isInStdNamespace()) - return true; - if (!filterInNamespace(encl->getQualifiedNameAsString())) - return true; - if constexpr (std::is_same_v || - std::is_same_v) { - if (!decl->isCompleteDefinition()) - return true; - } - if (decl->getAccess() == clang::AS_private || - decl->getAccess() == clang::AS_protected) - return true; - if (decl->isImplicit()) - return true; - - return false; -} - -struct HackDeclContext : clang::DeclContext { - bool islastDecl(clang::Decl *d) const { return d == LastDecl; } - clang::Decl *getLastDecl() const { return LastDecl; } -}; - -struct BindingsVisitor - : clang::LexicallyOrderedRecursiveASTVisitor { - BindingsVisitor(clang::CompilerInstance &ci, - std::shared_ptr outputFile) - : LexicallyOrderedRecursiveASTVisitor( - ci.getSourceManager()), - ci(ci), outputFile(outputFile), - mc(clang::ItaniumMangleContext::create(ci.getASTContext(), - ci.getDiagnostics())) {} - - bool VisitCXXRecordDecl(clang::CXXRecordDecl *decl) { - if (shouldSkip(decl)) - return true; - if (decl->isClass() || decl->isStruct()) { - if (emitClass(decl, ci, outputFile)) - visitedRecords.insert(decl); - } - return true; - } - - // clang-format off - /* method (member function) -> CXXMethodDecl (isStatic() tells statis/non-static) - * non-static data member -> FieldDecl - * static data member -> VarDecl with isStaticDataMember() == true - * static local variable -> VarDecl with isStaticLocal() == true - * non-static local variable -> VarDecl with isLocalVarDecl() == true - * there's also a isLocalVarDeclOrParm if you want local or parameter, etc - */ - // clang-format on - - bool VisitFieldDecl(clang::FieldDecl *decl) { - if (decl->isFunctionOrFunctionTemplate() || decl->isFunctionPointerType()) - return true; - // fields can be methods??? - if (llvm::isa(decl)) - return true; - if (!visitedRecords.contains(decl->getParent())) - return true; - if (decl->getAccess() == clang::AS_private || - decl->getAccess() == clang::AS_protected) - return true; - - if (decl->isAnonymousStructOrUnion()) { - clang::DiagnosticBuilder builder = ci.getDiagnostics().Report( - decl->getLocation(), ci.getDiagnostics().getCustomDiagID( - clang::DiagnosticsEngine::Note, - "anon structs/union fields not supported")); - return true; - } - if (decl->isBitField()) - return true; - - emitField(decl, ci, outputFile); - return true; - } - - bool VisitVarDecl(clang::VarDecl *decl) { - if (decl->isFunctionOrMethodVarDecl()) - return true; - if (llvm::isa(decl)) - return true; - if (auto parent = llvm::dyn_cast( - decl->getLexicalDeclContext())) { - if (visitedRecords.contains(parent)) { - if (decl->getAccess() == clang::AS_private || - decl->getAccess() == clang::AS_protected) - return true; - emitField(decl, ci, outputFile); - } - } - - return true; - } - - bool VisitCXXMethodDecl(clang::CXXMethodDecl *decl) { - if (shouldSkip(decl) || llvm::isa(decl) || - !visitedRecords.contains(decl->getParent())) - return true; - if (decl->isTemplated() || decl->isTemplateDecl() || - decl->isTemplateInstantiation() || - decl->isFunctionTemplateSpecialization()) { - clang::DiagnosticBuilder builder = ci.getDiagnostics().Report( - decl->getLocation(), ci.getDiagnostics().getCustomDiagID( - clang::DiagnosticsEngine::Note, - "template methods not supported yet")); - return true; - } - if (decl->getFriendObjectKind()) { - clang::DiagnosticBuilder builder = ci.getDiagnostics().Report( - decl->getLocation(), ci.getDiagnostics().getCustomDiagID( - clang::DiagnosticsEngine::Note, - "friend functions not supported")); - return true; - } - emitClassMethodOrFunction(decl, ci, outputFile); - return true; - } - - bool VisitFunctionDecl(clang::FunctionDecl *decl) { - if (shouldSkip(decl) || decl->isCXXClassMember()) - return true; - // clang-format off - // this - // template ::std::optional symbolizeEnum(::llvm::StringRef); - // is not a `TemplateDecl` but it is `Templated`... - // on the other hand every method in a template class `isTemplated` even - // if the template params don't play in the method decl (which is why this visitor can't be combined with VisitCXXMethodDecl) - // clang-format on - if (decl->isTemplated() || decl->isTemplateDecl() || - decl->isTemplateInstantiation() || - decl->isFunctionTemplateSpecialization()) { - clang::DiagnosticBuilder builder = ci.getDiagnostics().Report( - decl->getLocation(), ci.getDiagnostics().getCustomDiagID( - clang::DiagnosticsEngine::Note, - "template functions not supported yet")); - return true; - } - emitClassMethodOrFunction(decl, ci, outputFile); - return true; - } - - bool VisitEnumDecl(clang::EnumDecl *decl) { - if (shouldSkip(decl)) - return true; - if (decl->getQualifiedNameAsString().rfind("unnamed enum") != - std::string::npos) - return true; - emitEnum(llvm::cast(decl), ci, outputFile); - return true; - } - - void maybeEmitShard(const clang::CXXRecordDecl *decl) { - clang::CXXBaseSpecifier baseClass = *decl->bases_begin(); - std::string baseName = baseClass.getType().getAsString(getPrintingPolicy()); - if (baseName.rfind("mlir::Op<", 0) == 0) { - ++opClassesSeen; - if (opClassesSeen % ShardSize == 0) - outputFile->os() << llvm::formatv("// eudslpy-gen-shard {}\n", - int(opClassesSeen / ShardSize)); - } - } - - // implicit methods are the "last decls" - bool shouldVisitImplicitCode() const { return true; } - - // TODO(max): this is a hack and not stable - bool VisitDecl(clang::Decl *decl) { - const clang::DeclContext *declContext = decl->getDeclContext(); - HackDeclContext *ctx = - static_cast(decl->getDeclContext()); - if (declContext && declContext->isRecord()) { - const clang::CXXRecordDecl *recordDecl = - llvm::cast(declContext); - if (visitedRecords.contains(recordDecl) && ctx->islastDecl(decl)) { - outputFile->os() << "// " << recordDecl->getQualifiedNameAsString() - << "\n\n"; - if (recordDecl->getNumBases() == 1) - maybeEmitShard(recordDecl); - } - } - return true; - } - - clang::CompilerInstance &ci; - std::shared_ptr outputFile; - std::unique_ptr mc; - llvm::DenseSet visitedRecords; - int opClassesSeen{0}; -}; - -struct ClassStructEnumConsumer : clang::ASTConsumer { - ClassStructEnumConsumer( - clang::CompilerInstance &ci, - const std::shared_ptr &outputFile) - : visitor(ci, outputFile) {} - - void HandleTranslationUnit(clang::ASTContext &context) override { - if (!visitor.ci.hasSema()) - visitor.ci.createSema(clang::TU_Prefix, - /*CompletionConsumer*/ nullptr); - visitor.TraverseDecl(context.getTranslationUnitDecl()); - } - bool shouldSkipFunctionBody(clang::Decl *) override { return true; } - BindingsVisitor visitor; -}; - -struct GenerateBindingsAction : clang::ASTFrontendAction { - explicit GenerateBindingsAction( - const std::shared_ptr &outputFile) - : outputFile(outputFile) {} - std::unique_ptr - CreateASTConsumer(clang::CompilerInstance &compiler, - llvm::StringRef inFile) override { - compiler.getDiagnosticOpts().ShowLevel = true; - return std::make_unique(compiler, outputFile); - } - - // PARSE_INCOMPLETE - clang::TranslationUnitKind getTranslationUnitKind() override { - return clang::TU_Prefix; - } - - std::shared_ptr outputFile; -}; - -int main(int argc, char **argv) { - llvm::cl::ParseCommandLineOptions(argc, argv); - - llvm::ErrorOr> fileOrErr = - llvm::MemoryBuffer::getFileOrSTDIN(InputFilename); - if (std::error_code ec = fileOrErr.getError()) { - llvm::errs() << "Could not open input file: " << ec.message() << "\n"; - return -1; - } - llvm::StringRef buffer = fileOrErr.get()->getBuffer(); - - if (OutputFilename.empty()) { - llvm::errs() << "-o outputfile must be set.\n"; - return -1; - } - - std::error_code error; - auto outputFile = std::make_shared( - OutputFilename, error, llvm::sys::fs::OF_None); - if (error) { - llvm::errs() << "cannot open output file '" + OutputFilename + - "': " + error.message(); - return -1; - } - std::vector args{ - "-E", - "-xc++", - "-std=c++17", - "-fdirectives-only", - "-fkeep-system-includes", - "-fdelayed-template-parsing", - "-Wno-unused-command-line-argument", - // "-v", - // annoyingly clang will insert -internal-isystem with relative paths and - // those could hit on the build dir (which have headers and relationships - // amongst them that won't necessarily be valid) more annoyingly different - // toolchains decide differently which flag determines whether to include - // the tell-tale sign is if you uncomment -v and see "ignoring - // non-existant directory..." in the output - "-nobuiltininc", - "-nostdinc", - "-nostdinc++", - "-nostdlibinc", - }; - int clangArgs = args.size(); - for (const auto &includeDir : IncludeDirs) - args.emplace_back(llvm::formatv("-I{0}", includeDir)); - for (const auto &define : Defines) - args.emplace_back(llvm::formatv("-D{0}", define)); - - outputFile->os() << "// Generated with eudslpy-gen args:\n"; - outputFile->os() << "// " << InputFilename << " "; - auto arg = args.begin(); - for (std::advance(arg, clangArgs); arg != args.end(); ++arg) { - outputFile->os() << *arg << ' '; - } - for (auto ns : Namespaces) - outputFile->os() << "-namespaces=" << ns << " "; - outputFile->os() << "-o " << OutputFilename << "\n"; - outputFile->os() << "\n"; - - if (!clang::tooling::runToolOnCodeWithArgs( - std::make_unique(outputFile), buffer, args, - InputFilename)) { - llvm::errs() << "bindings generation failed.\n"; - return -1; - } - - outputFile->keep(); - - return 0; -} diff --git a/projects/eudsl-py/src/eudslpy_ext.cpp b/projects/eudsl-py/src/eudslpy_ext.cpp index 0948a8e8..b527527a 100644 --- a/projects/eudsl-py/src/eudslpy_ext.cpp +++ b/projects/eudsl-py/src/eudslpy_ext.cpp @@ -1,7 +1,8 @@ +// Copyright (c) 2025. +// // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// Copyright (c) 2024. #include #include @@ -254,109 +255,109 @@ void populateIRModule(nb::module_ &m) { } // too big -// extern void populateEUDSLGenaccModule(nb::module_ &m); +// extern void populateEUDSLGen_accModule(nb::module_ &m); -extern void populateEUDSLGenaffineModule(nb::module_ &m); +extern void populateEUDSLGen_affineModule(nb::module_ &m); -extern void populateEUDSLGenamdgpuModule(nb::module_ &m); +extern void populateEUDSLGen_amdgpuModule(nb::module_ &m); -// extern void populateEUDSLGenamxModule(nb::module_ &m); +// extern void populateEUDSLGen_amxModule(nb::module_ &m); -extern void populateEUDSLGenarithModule(nb::module_ &m); +extern void populateEUDSLGen_arithModule(nb::module_ &m); -// extern void populateEUDSLGenarm_neonModule(nb::module_ &m); +// extern void populateEUDSLGen_arm_neonModule(nb::module_ &m); // too big -// extern void populateEUDSLGenarm_smeModule(nb::module_ &m); +// extern void populateEUDSLGen_arm_smeModule(nb::module_ &m); -// extern void populateEUDSLGenarm_sveModule(nb::module_ &m); +// extern void populateEUDSLGen_arm_sveModule(nb::module_ &m); -extern void populateEUDSLGenasyncModule(nb::module_ &m); +extern void populateEUDSLGen_asyncModule(nb::module_ &m); -extern void populateEUDSLGenbufferizationModule(nb::module_ &m); +extern void populateEUDSLGen_bufferizationModule(nb::module_ &m); -extern void populateEUDSLGencfModule(nb::module_ &m); +extern void populateEUDSLGen_cfModule(nb::module_ &m); -extern void populateEUDSLGencomplexModule(nb::module_ &m); +extern void populateEUDSLGen_complexModule(nb::module_ &m); -// extern void populateEUDSLGenDLTIDialectModule(nb::module_ &m); +// extern void populateEUDSLGen_DLTIDialectModule(nb::module_ &m); -extern void populateEUDSLGenemitcModule(nb::module_ &m); +extern void populateEUDSLGen_emitcModule(nb::module_ &m); -extern void populateEUDSLGenfuncModule(nb::module_ &m); +extern void populateEUDSLGen_funcModule(nb::module_ &m); -extern void populateEUDSLGengpuModule(nb::module_ &m); +extern void populateEUDSLGen_gpuModule(nb::module_ &m); -extern void populateEUDSLGenindexModule(nb::module_ &m); +extern void populateEUDSLGen_indexModule(nb::module_ &m); // error: use of class template 'ArrayRef' requires template arguments; argument // deduction not allowed in conversion function type void -// populateEUDSLGenirdlModule(nb::module_ &m) { +// populateEUDSLGen_irdlModule(nb::module_ &m) { // using namespace llvm; -// #include "EUDSLGenirdl.cpp.inc" +// #include "EUDSLGen_irdl.cpp.inc" // } -extern void populateEUDSLGenlinalgModule(nb::module_ &m); +extern void populateEUDSLGen_linalgModule(nb::module_ &m); -extern void populateEUDSLGenLLVMModule(nb::module_ &m); +extern void populateEUDSLGen_LLVMModule(nb::module_ &m); -extern void populateEUDSLGenmathModule(nb::module_ &m); +extern void populateEUDSLGen_mathModule(nb::module_ &m); -extern void populateEUDSLGenmemrefModule(nb::module_ &m); +extern void populateEUDSLGen_memrefModule(nb::module_ &m); -// extern void populateEUDSLGenmeshModule(nb::module_ &m); +// extern void populateEUDSLGen_meshModule(nb::module_ &m); -// extern void populateEUDSLGenml_programModule(nb::module_ &m); +// extern void populateEUDSLGen_ml_programModule(nb::module_ &m); -// extern void populateEUDSLGenmpiModule(nb::module_ &m); +// extern void populateEUDSLGen_mpiModule(nb::module_ &m); -extern void populateEUDSLGennvgpuModule(nb::module_ &m); +extern void populateEUDSLGen_nvgpuModule(nb::module_ &m); -extern void populateEUDSLGenNVVMModule(nb::module_ &m); +extern void populateEUDSLGen_NVVMModule(nb::module_ &m); // mlir::omp::TaskloopOp::getEffects(llvm::SmallVectorImpl>&) -// void populateEUDSLGenompModule(nb::module_ &m) { -// #include "EUDSLGenomp.cpp.inc" +// void populateEUDSLGen_ompModule(nb::module_ &m) { +// #include "EUDSLGen_omp.cpp.inc" // } -extern void populateEUDSLGenpdlModule(nb::module_ &m); +extern void populateEUDSLGen_pdlModule(nb::module_ &m); -extern void populateEUDSLGenpdl_interpModule(nb::module_ &m); +extern void populateEUDSLGen_pdl_interpModule(nb::module_ &m); -extern void populateEUDSLGenpolynomialModule(nb::module_ &m); +extern void populateEUDSLGen_polynomialModule(nb::module_ &m); -// extern void populateEUDSLGenptrModule(nb::module_ &m); +// extern void populateEUDSLGen_ptrModule(nb::module_ &m); -extern void populateEUDSLGenquantModule(nb::module_ &m); +extern void populateEUDSLGen_quantModule(nb::module_ &m); -extern void populateEUDSLGenROCDLModule(nb::module_ &m); +extern void populateEUDSLGen_ROCDLModule(nb::module_ &m); -extern void populateEUDSLGenscfModule(nb::module_ &m); +extern void populateEUDSLGen_scfModule(nb::module_ &m); -extern void populateEUDSLGenshapeModule(nb::module_ &m); +extern void populateEUDSLGen_shapeModule(nb::module_ &m); -// extern void populateEUDSLGensparse_tensorModule(nb::module_ &m); +// extern void populateEUDSLGen_sparse_tensorModule(nb::module_ &m); // nb::detail::nb_func_new("get_vce_triple_attr_name"): mismatched // static/instance method flags in function overloads! extern void -// populateEUDSLGenspirvModule(nb::module_ &m); +// populateEUDSLGen_spirvModule(nb::module_ &m); -extern void populateEUDSLGentensorModule(nb::module_ &m); +extern void populateEUDSLGen_tensorModule(nb::module_ &m); -extern void populateEUDSLGentosaModule(nb::module_ &m); +extern void populateEUDSLGen_tosaModule(nb::module_ &m); -extern void populateEUDSLGentransformModule(nb::module_ &m); +extern void populateEUDSLGen_transformModule(nb::module_ &m); -extern void populateEUDSLGenubModule(nb::module_ &m); +extern void populateEUDSLGen_ubModule(nb::module_ &m); // can't cast std::pair -// void populateEUDSLGenvectorModule(nb::module_ &m) { -// #include "EUDSLGenvector.cpp.inc" +// void populateEUDSLGen_vectorModule(nb::module_ &m) { +// #include "EUDSLGen_vector.cpp.inc" // } -extern void populateEUDSLGenx86vectorModule(nb::module_ &m); +extern void populateEUDSLGen_x86vectorModule(nb::module_ &m); -// extern void populateEUDSLGenxegpuModule(nb::module_ &m); +// extern void populateEUDSLGen_xegpuModule(nb::module_ &m); NB_MODULE(eudslpy_ext, m) { bind_array_ref_smallvector(m); @@ -589,137 +590,137 @@ NB_MODULE(eudslpy_ext, m) { auto dialectsModule = m.def_submodule("dialects"); // auto accModule = dialectsModule.def_submodule("acc"); - // populateEUDSLGenaccModule(accModule); + // populateEUDSLGen_accModule(accModule); auto affineModule = dialectsModule.def_submodule("affine"); - populateEUDSLGenaffineModule(affineModule); + populateEUDSLGen_affineModule(affineModule); auto amdgpuModule = dialectsModule.def_submodule("amdgpu"); - populateEUDSLGenamdgpuModule(amdgpuModule); + populateEUDSLGen_amdgpuModule(amdgpuModule); // auto amxModule = dialectsModule.def_submodule("amx"); - // populateEUDSLGenamxModule(amxModule); + // populateEUDSLGen_amxModule(amxModule); auto arithModule = dialectsModule.def_submodule("arith"); - populateEUDSLGenarithModule(arithModule); + populateEUDSLGen_arithModule(arithModule); // auto arm_neonModule = dialectsModule.def_submodule("arm_neon"); - // populateEUDSLGenarm_neonModule(arm_neonModule); + // populateEUDSLGen_arm_neonModule(arm_neonModule); // auto arm_smeModule = dialectsModule.def_submodule("arm_sme"); - // populateEUDSLGenarm_smeModule(arm_smeModule); + // populateEUDSLGen_arm_smeModule(arm_smeModule); // auto arm_sveModule = dialectsModule.def_submodule("arm_sve"); - // populateEUDSLGenarm_sveModule(arm_sveModule); + // populateEUDSLGen_arm_sveModule(arm_sveModule); auto asyncModule = dialectsModule.def_submodule("async"); - populateEUDSLGenasyncModule(asyncModule); + populateEUDSLGen_asyncModule(asyncModule); auto bufferizationModule = dialectsModule.def_submodule("bufferization"); - populateEUDSLGenbufferizationModule(bufferizationModule); + populateEUDSLGen_bufferizationModule(bufferizationModule); auto cfModule = dialectsModule.def_submodule("cf"); - populateEUDSLGencfModule(cfModule); + populateEUDSLGen_cfModule(cfModule); auto complexModule = dialectsModule.def_submodule("complex"); - populateEUDSLGencomplexModule(complexModule); + populateEUDSLGen_complexModule(complexModule); // auto DLTIDialectModule = dialectsModule.def_submodule("DLTIDialect"); - // populateEUDSLGenDLTIDialectModule(DLTIDialectModule); + // populateEUDSLGen_DLTIDialectModule(DLTIDialectModule); auto emitcModule = dialectsModule.def_submodule("emitc"); - populateEUDSLGenemitcModule(emitcModule); + populateEUDSLGen_emitcModule(emitcModule); auto funcModule = dialectsModule.def_submodule("func"); - populateEUDSLGenfuncModule(funcModule); + populateEUDSLGen_funcModule(funcModule); auto gpuModule = dialectsModule.def_submodule("gpu"); - populateEUDSLGengpuModule(gpuModule); + populateEUDSLGen_gpuModule(gpuModule); auto indexModule = dialectsModule.def_submodule("index"); - populateEUDSLGenindexModule(indexModule); + populateEUDSLGen_indexModule(indexModule); // auto irdlModule = dialectsModule.def_submodule("irdl"); - // populateEUDSLGenirdlModule(irdlModule); + // populateEUDSLGen_irdlModule(irdlModule); auto linalgModule = dialectsModule.def_submodule("linalg"); - populateEUDSLGenlinalgModule(linalgModule); + populateEUDSLGen_linalgModule(linalgModule); auto LLVMModule = dialectsModule.def_submodule("LLVM"); - populateEUDSLGenLLVMModule(LLVMModule); + populateEUDSLGen_LLVMModule(LLVMModule); auto mathModule = dialectsModule.def_submodule("math"); - populateEUDSLGenmathModule(mathModule); + populateEUDSLGen_mathModule(mathModule); auto memrefModule = dialectsModule.def_submodule("memref"); - populateEUDSLGenmemrefModule(memrefModule); + populateEUDSLGen_memrefModule(memrefModule); // auto meshModule = dialectsModule.def_submodule("mesh"); - // populateEUDSLGenmeshModule(meshModule); + // populateEUDSLGen_meshModule(meshModule); // auto ml_programModule = dialectsModule.def_submodule("ml_program"); - // populateEUDSLGenml_programModule(ml_programModule); + // populateEUDSLGen_ml_programModule(ml_programModule); // auto mpiModule = dialectsModule.def_submodule("mpi"); - // populateEUDSLGenmpiModule(mpiModule); + // populateEUDSLGen_mpiModule(mpiModule); auto nvgpuModule = dialectsModule.def_submodule("nvgpu"); - populateEUDSLGennvgpuModule(nvgpuModule); + populateEUDSLGen_nvgpuModule(nvgpuModule); auto NVVMModule = dialectsModule.def_submodule("NVVM"); - populateEUDSLGenNVVMModule(NVVMModule); + populateEUDSLGen_NVVMModule(NVVMModule); // auto ompModule = dialectsModule.def_submodule("omp"); - // populateEUDSLGenompModule(ompModule); + // populateEUDSLGen_ompModule(ompModule); auto pdlModule = dialectsModule.def_submodule("pdl"); - populateEUDSLGenpdlModule(pdlModule); + populateEUDSLGen_pdlModule(pdlModule); auto pdl_interpModule = dialectsModule.def_submodule("pdl_interp"); - populateEUDSLGenpdl_interpModule(pdl_interpModule); + populateEUDSLGen_pdl_interpModule(pdl_interpModule); auto polynomialModule = dialectsModule.def_submodule("polynomial"); - populateEUDSLGenpolynomialModule(polynomialModule); + populateEUDSLGen_polynomialModule(polynomialModule); // auto ptrModule = dialectsModule.def_submodule("ptr"); - // populateEUDSLGenptrModule(ptrModule); + // populateEUDSLGen_ptrModule(ptrModule); // auto quantModule = dialectsModule.def_submodule("quant"); - // populateEUDSLGenquantModule(quantModule); + // populateEUDSLGen_quantModule(quantModule); auto ROCDLModule = dialectsModule.def_submodule("ROCDL"); - populateEUDSLGenROCDLModule(ROCDLModule); + populateEUDSLGen_ROCDLModule(ROCDLModule); auto scfModule = dialectsModule.def_submodule("scf"); - populateEUDSLGenscfModule(scfModule); + populateEUDSLGen_scfModule(scfModule); auto shapeModule = dialectsModule.def_submodule("shape"); - populateEUDSLGenshapeModule(shapeModule); + populateEUDSLGen_shapeModule(shapeModule); // auto sparse_tensorModule = dialectsModule.def_submodule("sparse_tensor"); - // populateEUDSLGensparse_tensorModule(sparse_tensorModule); + // populateEUDSLGen_sparse_tensorModule(sparse_tensorModule); // auto spirvModule = dialectsModule.def_submodule("spirv"); - // populateEUDSLGenspirvModule(spirvModule); + // populateEUDSLGen_spirvModule(spirvModule); auto tensorModule = dialectsModule.def_submodule("tensor"); - populateEUDSLGentensorModule(tensorModule); + populateEUDSLGen_tensorModule(tensorModule); auto tosaModule = dialectsModule.def_submodule("tosa"); - populateEUDSLGentosaModule(tosaModule); + populateEUDSLGen_tosaModule(tosaModule); // auto transformModule = dialectsModule.def_submodule("transform"); - // populateEUDSLGentransformModule(transformModule); + // populateEUDSLGen_transformModule(transformModule); // auto ubModule = dialectsModule.def_submodule("ub"); - // populateEUDSLGenubModule(ubModule); + // populateEUDSLGen_ubModule(ubModule); // auto vectorModule = dialectsModule.def_submodule("vector"); - // populateEUDSLGenvectorModule(vectorModule); + // populateEUDSLGen_vectorModule(vectorModule); // auto x86vectorModule = dialectsModule.def_submodule("x86vector"); - // populateEUDSLGenx86vectorModule(x86vectorModule); + // populateEUDSLGen_x86vectorModule(x86vectorModule); // auto xegpuModule = dialectsModule.def_submodule("xegpu"); - // populateEUDSLGenxegpuModule(xegpuModule); + // populateEUDSLGen_xegpuModule(xegpuModule); } From dc21b375188fe52ce11e74f60b0fda04817d2928 Mon Sep 17 00:00:00 2001 From: Maksim Levental Date: Thu, 2 Jan 2025 18:30:46 -0500 Subject: [PATCH 4/6] [eudsl][nbgen] eudsl-nbgen into eudsl-py working with cibuildwheel --- .../workflows/build_test_release_eudsl.yml | 46 +++--- projects/eudsl-nbgen/CMakeLists.txt | 1 - .../cmake/eudsl_nbgen-config.cmake | 27 ++-- .../cmake/make_generated_shards.py | 81 ---------- projects/eudsl-nbgen/pyproject.toml | 52 ++++++ projects/eudsl-nbgen/src/eudsl-nbgen.cpp | 151 +++++++++++++++++- projects/eudsl-py/src/eudslpy_ext.cpp | 3 +- projects/eudsl-py/tests/test_bindings.py | 2 +- 8 files changed, 243 insertions(+), 120 deletions(-) delete mode 100644 projects/eudsl-nbgen/cmake/make_generated_shards.py diff --git a/.github/workflows/build_test_release_eudsl.yml b/.github/workflows/build_test_release_eudsl.yml index e631e829..46d97d05 100644 --- a/.github/workflows/build_test_release_eudsl.yml +++ b/.github/workflows/build_test_release_eudsl.yml @@ -104,6 +104,19 @@ jobs: curl -sLO $RELEASE_URL tar xf $RELEASE_PREFIX*.tar.gz + if [[ "${{ matrix.os }}" == "ubuntu" ]]; then + echo "LLVM_DIR=/host/$PWD/llvm-install/lib/cmake/llvm" >> $GITHUB_ENV + echo "MLIR_DIR=/host/$PWD/llvm-install/lib/cmake/mlir" >> $GITHUB_ENV + echo "Clang_DIR=/host/$PWD/llvm-install/lib/cmake/clang" >> $GITHUB_ENV + echo "CCACHE_DIR=/host/$CCACHE_DIR" >> $GITHUB_ENV + echo "PIP_FIND_LINKS=/host/$PWD/wheelhouse" >> $GITHUB_ENV + else + echo "LLVM_DIR=$PWD/llvm-install/lib/cmake/llvm" >> $GITHUB_ENV + echo "MLIR_DIR=$PWD/llvm-install/lib/cmake/mlir" >> $GITHUB_ENV + echo "Clang_DIR=$PWD/llvm-install/lib/cmake/clang" >> $GITHUB_ENV + echo "PIP_FIND_LINKS=$PWD/wheelhouse" >> $GITHUB_ENV + fi + # since linux builds execute in the cibuildwheel almalinux container if [[ "${{ matrix.os }}" == "ubuntu" ]]; then echo CC=clang >> $GITHUB_ENV @@ -123,35 +136,18 @@ jobs: - name: "Build eudsl-tblgen" run: | - if [[ "${{ matrix.os }}" == "ubuntu" ]]; then - export CMAKE_PREFIX_PATH="/host/$PWD/llvm-install" - export CCACHE_DIR="/host/$CCACHE_DIR" - else - export CMAKE_PREFIX_PATH="$PWD/llvm-install" - fi - $python3_command -m cibuildwheel "$PWD/projects/eudsl-tblgen" --output-dir wheelhouse - name: "Build eudsl-nbgen" + if: ${{ ! startsWith(matrix.os, 'windows') }} run: | - # this is not run in cibuildwheel container so its different - export CMAKE_PREFIX_PATH="$PWD/llvm-install" - $python3_command -m pip wheel "$PWD/projects/eudsl-nbgen" -w wheelhouse -v + $python3_command -m cibuildwheel "$PWD/projects/eudsl-nbgen" --output-dir wheelhouse - name: "Build eudsl-py" if: ${{ ! startsWith(matrix.os, 'windows') }} run: | - if [[ "${{ matrix.os }}" == "ubuntu" ]]; then - export CMAKE_PREFIX_PATH="/host/$PWD/llvm-install" - export CCACHE_DIR="/host/$CCACHE_DIR" - export PIP_FIND_LINKS="/host/$PWD/wheelhouse" - else - export CMAKE_PREFIX_PATH="$PWD/llvm-install" - export PIP_FIND_LINKS="$PWD/wheelhouse" - fi - # prevent OOM on free GHA export DISABLE_COMPILE_OPT="${{ inputs.debug_with_tmate }}" $python3_command -m cibuildwheel "$PWD/projects/eudsl-py" --output-dir wheelhouse @@ -299,10 +295,22 @@ jobs: id-token: write contents: write + strategy: + fail-fast: false + matrix: + runs-on: ["ubuntu-22.04"] + include: [ + {runs-on: "ubuntu-22.04", name: "ubuntu_x86_64"}, + {runs-on: "ubuntu-22.04", name: "macos_arm64"}, + {runs-on: "ubuntu-22.04", name: "macos_x86_64"}, + {runs-on: "ubuntu-22.04", name: "windows_x86_64"}, + ] + steps: - uses: actions/download-artifact@v4 with: + name: eudsl_${{ matrix.name }}_artifact path: wheelhouse - name: Release current commit diff --git a/projects/eudsl-nbgen/CMakeLists.txt b/projects/eudsl-nbgen/CMakeLists.txt index b1e92720..b5faa3f3 100644 --- a/projects/eudsl-nbgen/CMakeLists.txt +++ b/projects/eudsl-nbgen/CMakeLists.txt @@ -87,6 +87,5 @@ install( install( FILES cmake/eudsl_nbgen-config.cmake - cmake/make_generated_shards.py DESTINATION "${EUDSL_NBGEN_INSTALL_DATADIR}/cmake" ) diff --git a/projects/eudsl-nbgen/cmake/eudsl_nbgen-config.cmake b/projects/eudsl-nbgen/cmake/eudsl_nbgen-config.cmake index bfdc4968..235e8706 100644 --- a/projects/eudsl-nbgen/cmake/eudsl_nbgen-config.cmake +++ b/projects/eudsl-nbgen/cmake/eudsl_nbgen-config.cmake @@ -5,23 +5,25 @@ # copy-pasta from AddMLIR.cmake/AddLLVM.cmake/TableGen.cmake -function(eudslpygen target inputFile) - set(EUDSLPYGEN_TARGET_DEFINITIONS ${inputFile}) +function(eudslpygen target input_file) + set(EUDSLPYGEN_TARGET_DEFINITIONS ${input_file}) cmake_parse_arguments(ARG "" "" "DEPENDS;EXTRA_INCLUDES;NAMESPACES" ${ARGN}) if (IS_ABSOLUTE ${EUDSLPYGEN_TARGET_DEFINITIONS}) - set(EUDSLPYGEN_TARGET_DEFINITIONS_ABSOLUTE ${inputFile}) + set(EUDSLPYGEN_TARGET_DEFINITIONS_ABSOLUTE ${input_file}) else() - set(EUDSLPYGEN_TARGET_DEFINITIONS_ABSOLUTE ${CMAKE_CURRENT_SOURCE_DIR}/${inputFile}) + set(EUDSLPYGEN_TARGET_DEFINITIONS_ABSOLUTE ${CMAKE_CURRENT_SOURCE_DIR}/${input_file}) endif() if(NOT EXISTS "${EUDSLPYGEN_TARGET_DEFINITIONS_ABSOLUTE}") - message(FATAL_ERROR "${inputFile} does not exist") + message(FATAL_ERROR "${input_file} does not exist") endif() get_directory_property(eudslpygen_includes INCLUDE_DIRECTORIES) - list(APPEND eudslpygen_includes ${CMAKE_CXX_IMPLICIT_INCLUDE_DIRECTORIES} ${ARG_EXTRA_INCLUDES}) + list(TRANSFORM ARG_EXTRA_INCLUDES PREPEND -I) + list(APPEND eudslpygen_includes ${CMAKE_CXX_IMPLICIT_INCLUDE_DIRECTORIES}) list(REMOVE_ITEM eudslpygen_includes "") list(TRANSFORM eudslpygen_includes PREPEND -I) + list(APPEND eudslpygen_includes ${ARG_EXTRA_INCLUDES}) set(_gen_target_dir "${CMAKE_CURRENT_BINARY_DIR}/generated/${target}") file(MAKE_DIRECTORY ${_gen_target_dir}) @@ -41,7 +43,7 @@ function(eudslpygen target inputFile) -o ${_depfile} ) execute_process(COMMAND ${clang_command} RESULT_VARIABLE _had_error_depfile - COMMAND_ECHO STDOUT + # COMMAND_ECHO STDOUT ) if (IS_ABSOLUTE ${EUDSLPYGEN_TARGET_DEFINITIONS}) @@ -64,7 +66,7 @@ function(eudslpygen target inputFile) -o "${_full_gen_file}" WORKING_DIRECTORY ${CMAKE_BINARY_DIR} RESULT_VARIABLE _had_error_gen_cpp - COMMAND_ECHO STDOUT + # COMMAND_ECHO STDOUT ) if((_had_error_gen_cpp AND NOT _had_error_gen_cpp EQUAL 0) OR NOT EXISTS "${_full_gen_file}") message(FATAL_ERROR "failed to create ${_full_gen_file}: ${_had_error_gen_cpp}") @@ -72,11 +74,12 @@ function(eudslpygen target inputFile) # this is the specific thing connected the dependencies... set_property(DIRECTORY APPEND PROPERTY CMAKE_CONFIGURE_DEPENDS ${_full_gen_file}) execute_process( - COMMAND ${Python_EXECUTABLE} ${CMAKE_CURRENT_FUNCTION_LIST_DIR}/make_generated_shards.py - ${_full_gen_file} -t ${target} -I ${ARG_EXTRA_INCLUDES} ${EUDSLPYGEN_TARGET_DEFINITIONS_ABSOLUTE} + COMMAND ${EUDSLPY_EUDSLPYGEN_EXE} -shardify ${_full_gen_file} + # ARG_EXTRA_INCLUDES has already had -I prepended + -shard-target ${target} ${ARG_EXTRA_INCLUDES} -I ${EUDSLPYGEN_TARGET_DEFINITIONS_ABSOLUTE} WORKING_DIRECTORY ${CMAKE_BINARY_DIR} RESULT_VARIABLE _had_error_gen_sharded - COMMAND_ECHO STDOUT + # COMMAND_ECHO STDOUT ) if((_had_error_gen_sharded AND NOT _had_error_gen_sharded EQUAL 0) OR NOT EXISTS "${_full_gen_file}.sharded.cpp") message(FATAL_ERROR "failed to create ${_full_gen_file}.sharded.cpp: ${_had_error_gen_sharded}") @@ -148,4 +151,4 @@ function(patch_mlir_llvm_rpath target) endif() set_target_properties(${target} PROPERTIES INSTALL_RPATH "${_origin_prefix}") endif() -endfunction() \ No newline at end of file +endfunction() diff --git a/projects/eudsl-nbgen/cmake/make_generated_shards.py b/projects/eudsl-nbgen/cmake/make_generated_shards.py deleted file mode 100644 index 207fe4f4..00000000 --- a/projects/eudsl-nbgen/cmake/make_generated_shards.py +++ /dev/null @@ -1,81 +0,0 @@ -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# Copyright (c) 2024. - -import argparse -import re -from pathlib import Path -from textwrap import dedent - - -def make_source_shards(filename: Path, target, extra_includes, max_num_shards): - assert filename.name.endswith("cpp.gen"), "expected .cpp.gen file" - with open(filename) as f: - source = f.read() - shards = re.split(r"// eudslpy-gen-shard \d+", source) - if target is None: - target = filename.stem.split(".")[0] - for i, shar in enumerate(shards): - with open(f"{filename}.shard.{i}.cpp", "w") as f: - if extra_includes is not None: - for inc in extra_includes: - print(f'#include "{inc}"', file=f) - print( - dedent( - f"""\ - #include - namespace nb = nanobind; - using namespace nb::literals; - using namespace mlir; - using namespace llvm; - #include "type_casters.h" - - void populate{target}{i}Module(nb::module_ &m) {{""" - ), - file=f, - ) - print(shar, file=f) - print("}", file=f) - - if max_num_shards is not None: - if len(shards) > max_num_shards: - raise RuntimeError(f"expected less than {max_num_shards} shards") - for i in range(len(shards), max_num_shards): - with open(f"{filename}.shard.{i}.cpp", "w") as f: - print(f"// dummy shard {i}", file=f) - - with open(f"{filename}.sharded.cpp", "w") as f: - print( - dedent( - f"""\ - #include - namespace nb = nanobind; - using namespace nb::literals; - void populate{target}Module(nb::module_ &m) {{""" - ), - file=f, - ) - for i in range(len(shards)): - print( - dedent(f"extern void populate{target}{i}Module(nb::module_ &m);"), - file=f, - ) - for i in range(len(shards)): - print(dedent(f"populate{target}{i}Module(m);"), file=f) - - print("}", file=f) - - -if __name__ == "__main__": - parser = argparse.ArgumentParser() - parser.add_argument("filename") - parser.add_argument("-t", "--target") - parser.add_argument("-I", "--extra_includes", nargs="*") - parser.add_argument("-m", "--max-num-shards", type=int) - args = parser.parse_args() - if args.max_num_shards: - args.max_num_shards += 1 - make_source_shards( - Path(args.filename), args.target, args.extra_includes, args.max_num_shards - ) diff --git a/projects/eudsl-nbgen/pyproject.toml b/projects/eudsl-nbgen/pyproject.toml index f1e58aae..d6617a6b 100644 --- a/projects/eudsl-nbgen/pyproject.toml +++ b/projects/eudsl-nbgen/pyproject.toml @@ -33,3 +33,55 @@ CMAKE_C_COMPILER_LAUNCHER = { env = "CMAKE_C_COMPILER_LAUNCHER", default = "" } CMAKE_CXX_COMPILER_LAUNCHER = { env = "CMAKE_CXX_COMPILER_LAUNCHER", default = "" } CMAKE_CXX_VISIBILITY_PRESET = "hidden" CMAKE_VERBOSE_MAKEFILE = "ON" + +[tool.cibuildwheel] +build-verbosity = 1 +# build only one version +build = ["cp39*"] +skip = ["*-manylinux_i686", "*-musllinux*", "pp*", "*-win32"] +archs = ["auto64"] +manylinux-x86_64-image = "manylinux_2_28" +environment-pass = [ + "LLVM_DIR", + "CMAKE_GENERATOR", + "CMAKE_PREFIX_PATH", + "CC", + "CXX", + # ccache + "CCACHE_DIR", + "CCACHE_MAXSIZE=700M", + "CCACHE_SLOPPINESS", + "CCACHE_CPP2", + "CCACHE_UMASK", + "CMAKE_C_COMPILER_LAUNCHER", + "CMAKE_CXX_COMPILER_LAUNCHER" +] +before-build = [ + "export CCACHE_DIR=$CCACHE_DIR/$(python -c 'import platform; print(platform.python_version())')", + "mkdir -p $CCACHE_DIR", + "ccache -z" +] +# uncomment to make sure ccache is working inside containers +test-command = "ccache -s" +repair-wheel-command = [] + +[tool.cibuildwheel.linux] +before-all = [ + "yum install -y clang", + # ccache + "curl -sLO https://github.com/ccache/ccache/releases/download/v4.10.2/ccache-4.10.2-linux-x86_64.tar.xz", + "tar -xf ccache-4.10.2-linux-x86_64.tar.xz", + "pushd ccache-4.10.2-linux-x86_64 && make install && popd" +] +# synchronize TZ with host so ccache files have correct timestamp +container-engine = { name = "docker", create-args = ["-v", "/etc/timezone:/etc/timezone:ro", "-v", "/etc/localtime:/etc/localtime:ro"] } + +[tool.cibuildwheel.macos] +before-build = [ + "ccache -z" +] + +[tool.cibuildwheel.windows] +before-build = [ + "ccache -z" +] diff --git a/projects/eudsl-nbgen/src/eudsl-nbgen.cpp b/projects/eudsl-nbgen/src/eudsl-nbgen.cpp index 0b65ab64..6db162f1 100644 --- a/projects/eudsl-nbgen/src/eudsl-nbgen.cpp +++ b/projects/eudsl-nbgen/src/eudsl-nbgen.cpp @@ -15,7 +15,9 @@ #include "llvm/Support/FormatVariadic.h" #include "llvm/Support/ToolOutputFile.h" +#include #include +#include static llvm::cl::OptionCategory EUDSLPYGenCat("Options for eudslpy-gen"); static llvm::cl::opt InputFilename(llvm::cl::Positional, @@ -30,7 +32,7 @@ static llvm::cl::list static llvm::cl::opt OutputFilename("o", llvm::cl::desc("Output filename"), - llvm::cl::value_desc("filename"), llvm::cl::Required, + llvm::cl::value_desc("filename"), llvm::cl::cat(EUDSLPYGenCat)); static llvm::cl::list @@ -47,6 +49,22 @@ static llvm::cl::opt ShardSize("shard-size", llvm::cl::desc("Shard size"), llvm::cl::cat(EUDSLPYGenCat), llvm::cl::init(10)); +static llvm::cl::opt Shardify("shardify", + llvm::cl::desc("Split into shards"), + llvm::cl::value_desc("shardify"), + llvm::cl::cat(EUDSLPYGenCat), + llvm::cl::init(false)); + +static llvm::cl::opt + MaxNumShards("max-number-shards", + llvm::cl::desc("Maximum number of shards to split into"), + llvm::cl::value_desc("max number shards"), + llvm::cl::cat(EUDSLPYGenCat), llvm::cl::init(-1)); + +static llvm::cl::opt ShardTarget( + "shard-target", llvm::cl::desc("CMake target that is being sharded"), + llvm::cl::value_desc("sharding target"), llvm::cl::cat(EUDSLPYGenCat)); + static bool filterInNamespace(const std::string &s) { if (Namespaces.empty()) return true; @@ -648,9 +666,7 @@ struct GenerateBindingsAction : clang::ASTFrontendAction { std::shared_ptr outputFile; }; -int main(int argc, char **argv) { - llvm::cl::ParseCommandLineOptions(argc, argv); - +static int generate() { llvm::ErrorOr> fileOrErr = llvm::MemoryBuffer::getFileOrSTDIN(InputFilename); if (std::error_code ec = fileOrErr.getError()) { @@ -720,3 +736,130 @@ int main(int argc, char **argv) { return 0; } + +static int makeSourceShards() { + // Ensure the filename ends with ".cpp.gen" + if (InputFilename.substr(InputFilename.size() - 8) != ".cpp.gen") { + llvm::errs() << "expected .cpp.gen file\n"; + return -1; + } + + std::ifstream file(InputFilename); + if (!file.is_open()) { + llvm::errs() << "Failed to open file\n"; + return -1; + } + + std::stringstream buffer; + buffer << file.rdbuf(); + std::string source = buffer.str(); + + // Split source by regex pattern + std::regex re("// eudslpy-gen-shard \\d+"); + std::sregex_token_iterator iter(source.begin(), source.end(), re, -1); + std::sregex_token_iterator end; + + std::vector shards; + while (iter != end) + shards.push_back(*iter++); + + std::string finalTarget = + ShardTarget.empty() + ? InputFilename.substr(0, InputFilename.find_last_of(".")) + : ShardTarget; + + // Generate shard files + + for (size_t i = 0; i < shards.size(); ++i) { + std::ofstream shardFile(InputFilename.getValue() + ".shard." + + std::to_string(i) + ".cpp"); + if (!shardFile.is_open()) { + llvm::errs() << "Failed to create shard file"; + return -1; + } + + if (!IncludeDirs.empty()) { + for (const std::string &inc : IncludeDirs) + shardFile << "#include \"" << inc << "\"" << std::endl; + } + + // clang-format off + shardFile << R"( +#include +namespace nb = nanobind; +using namespace nb::literals; +using namespace mlir; +using namespace llvm; +#include "type_casters.h" + +void populate)" << finalTarget << i << R"(Module(nb::module_ &m) { +)"; + // clang-format on + shardFile << shards[i] << std::endl; + shardFile << "}" << std::endl; + shardFile.flush(); + shardFile.close(); + } + + // Handle max number of shards + if (MaxNumShards != -1 && shards.size() > static_cast(MaxNumShards)) { + llvm::errs() << "expected less than " + std::to_string(MaxNumShards) + + " shards"; + return -1; + } + if (MaxNumShards == -1) + MaxNumShards = shards.size(); + + for (size_t i = shards.size(); i < static_cast(MaxNumShards); ++i) { + std::ofstream dummyShardFile(InputFilename.getValue() + ".shard." + + std::to_string(i) + ".cpp"); + if (!dummyShardFile.is_open()) { + llvm::errs() << "Failed to create dummy shard file"; + return -1; + } + dummyShardFile << "// dummy shard " << i << std::endl; + dummyShardFile.flush(); + dummyShardFile.close(); + } + + // Generate the final sharded cpp file + std::ofstream finalShardedFile(InputFilename.getValue() + ".sharded.cpp"); + if (!finalShardedFile.is_open()) { + llvm::errs() << "Failed to create final sharded file"; + return -1; + } + + // clang-format off + finalShardedFile << R"( +#include +namespace nb = nanobind; +using namespace nb::literals; + +void populate)" << finalTarget << R"(Module(nb::module_ &m) { +)"; + // clang-format on + + for (size_t i = 0; i < shards.size(); ++i) + finalShardedFile << "extern void populate" << finalTarget << i + << "Module(nb::module_ &m);" << std::endl; + + for (size_t i = 0; i < shards.size(); ++i) + finalShardedFile << "populate" << finalTarget << i << "Module(m);" + << std::endl; + + finalShardedFile << "}" << std::endl; + finalShardedFile.flush(); + finalShardedFile.close(); + + return 0; +} + +int main(int argc, char **argv) { + llvm::cl::ParseCommandLineOptions(argc, argv); + + if (Shardify) + return makeSourceShards(); + return generate(); + + return 0; +} diff --git a/projects/eudsl-py/src/eudslpy_ext.cpp b/projects/eudsl-py/src/eudslpy_ext.cpp index b527527a..0b3fc7dc 100644 --- a/projects/eudsl-py/src/eudslpy_ext.cpp +++ b/projects/eudsl-py/src/eudslpy_ext.cpp @@ -1,8 +1,7 @@ -// Copyright (c) 2025. -// // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// Copyright (c) 2024. #include #include diff --git a/projects/eudsl-py/tests/test_bindings.py b/projects/eudsl-py/tests/test_bindings.py index f0feb9df..a48afad5 100644 --- a/projects/eudsl-py/tests/test_bindings.py +++ b/projects/eudsl-py/tests/test_bindings.py @@ -117,4 +117,4 @@ def test_types(): if __name__ == "__main__": test_array_ref() test_arith_dialect() - test_types() \ No newline at end of file + test_types() From 21b03c4aa9c2f931053c42e106ef32d85eaf65bc Mon Sep 17 00:00:00 2001 From: max Date: Thu, 2 Jan 2025 21:39:40 -0500 Subject: [PATCH 5/6] comment out some targets to reduce compile time --- .../workflows/build_test_release_eudsl.yml | 2 +- projects/eudsl-py/CMakeLists.txt | 129 +++++++++--------- projects/eudsl-py/pyproject.toml | 1 + 3 files changed, 67 insertions(+), 65 deletions(-) diff --git a/.github/workflows/build_test_release_eudsl.yml b/.github/workflows/build_test_release_eudsl.yml index 46d97d05..ce31af07 100644 --- a/.github/workflows/build_test_release_eudsl.yml +++ b/.github/workflows/build_test_release_eudsl.yml @@ -149,7 +149,7 @@ jobs: run: | # prevent OOM on free GHA - export DISABLE_COMPILE_OPT="${{ inputs.debug_with_tmate }}" + export EUDSLPY_DISABLE_COMPILE_OPT="${{ inputs.eudslpy_disable_compile_opt == 'true' && 'ON' || 'OFF' }}" $python3_command -m cibuildwheel "$PWD/projects/eudsl-py" --output-dir wheelhouse - name: "Save cache" diff --git a/projects/eudsl-py/CMakeLists.txt b/projects/eudsl-py/CMakeLists.txt index c212ea2b..7cabdb46 100644 --- a/projects/eudsl-py/CMakeLists.txt +++ b/projects/eudsl-py/CMakeLists.txt @@ -128,20 +128,20 @@ eudslpygen(EUDSLGen_amdgpu NAMESPACES mlir::amdgpu mlir::amdgpu::detail ) -eudslpygen(EUDSLGen_amx - ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h - NAMESPACES mlir::amx mlir::amx::detail -) +#eudslpygen(EUDSLGen_amx +# ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h +# NAMESPACES mlir::amx mlir::amx::detail +#) eudslpygen(EUDSLGen_arith ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::arith mlir::arith::detail ) -eudslpygen(EUDSLGen_arm_neon - ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h - NAMESPACES mlir::arm_neon mlir::arm_neon::detail -) +#eudslpygen(EUDSLGen_arm_neon +# ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h +# NAMESPACES mlir::arm_neon mlir::arm_neon::detail +#) # too big # eudslpygen(EUDSLGen_arm_sme @@ -149,10 +149,10 @@ eudslpygen(EUDSLGen_arm_neon # NAMESPACES mlir::arm_sme mlir::arm_sme::detail # ) -eudslpygen(EUDSLGen_arm_sve - ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h - NAMESPACES mlir::arm_sve mlir::arm_sve::detail -) +#eudslpygen(EUDSLGen_arm_sve +# ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h +# NAMESPACES mlir::arm_sve mlir::arm_sve::detail +#) eudslpygen(EUDSLGen_async ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h @@ -179,10 +179,10 @@ eudslpygen(EUDSLGen_complex NAMESPACES mlir::complex mlir::complex::detail ) -eudslpygen(EUDSLGen_DLTIDialect - ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h - NAMESPACES mlir::DLTIDialect mlir::DLTIDialect::detail -) +#eudslpygen(EUDSLGen_DLTIDialect +# ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h +# NAMESPACES mlir::DLTIDialect mlir::DLTIDialect::detail +#) eudslpygen(EUDSLGen_emitc ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h @@ -235,22 +235,22 @@ eudslpygen(EUDSLGen_memref NAMESPACES mlir::memref mlir::memref::detail ) -eudslpygen(EUDSLGen_mesh - ${MLIR_INCLUDE_DIR}/mlir/Dialect/Mesh/IR/MeshOps.h - NAMESPACES mlir::mesh mlir::mesh::detail - EXTRA_INCLUDES - mlir/Dialect/Mesh/IR/MeshOps.h -) +#eudslpygen(EUDSLGen_mesh +# ${MLIR_INCLUDE_DIR}/mlir/Dialect/Mesh/IR/MeshOps.h +# NAMESPACES mlir::mesh mlir::mesh::detail +# EXTRA_INCLUDES +# mlir/Dialect/Mesh/IR/MeshOps.h +#) -eudslpygen(EUDSLGen_ml_program - ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h - NAMESPACES mlir::ml_program mlir::ml_program::detail -) +#eudslpygen(EUDSLGen_ml_program +# ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h +# NAMESPACES mlir::ml_program mlir::ml_program::detail +#) -eudslpygen(EUDSLGen_mpi - ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h - NAMESPACES mlir::mpi mlir::mpi::detail -) +#eudslpygen(EUDSLGen_mpi +# ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h +# NAMESPACES mlir::mpi mlir::mpi::detail +#) eudslpygen(EUDSLGen_nvgpu ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h @@ -287,20 +287,20 @@ eudslpygen(EUDSLGen_polynomial mlir/IR/PatternMatch.h ) -eudslpygen(EUDSLGen_ptr - ${MLIR_INCLUDE_DIR}/mlir/Dialect/Ptr/IR/PtrOps.h - NAMESPACES mlir::ptr mlir::ptr::detail - EXTRA_INCLUDES - mlir/Dialect/Ptr/IR/PtrOps.h - mlir/IR/DialectImplementation.h -) +#eudslpygen(EUDSLGen_ptr +# ${MLIR_INCLUDE_DIR}/mlir/Dialect/Ptr/IR/PtrOps.h +# NAMESPACES mlir::ptr mlir::ptr::detail +# EXTRA_INCLUDES +# mlir/Dialect/Ptr/IR/PtrOps.h +# mlir/IR/DialectImplementation.h +#) -eudslpygen(EUDSLGen_quant - ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h - NAMESPACES mlir::quant mlir::quant::detail - EXTRA_INCLUDES - mlir/Dialect/Quant/IR/QuantTypes.h -) +#eudslpygen(EUDSLGen_quant +# ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h +# NAMESPACES mlir::quant mlir::quant::detail +# EXTRA_INCLUDES +# mlir/Dialect/Quant/IR/QuantTypes.h +#) eudslpygen(EUDSLGen_ROCDL ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h @@ -317,10 +317,10 @@ eudslpygen(EUDSLGen_shape NAMESPACES mlir::shape mlir::shape::detail ) -eudslpygen(EUDSLGen_sparse_tensor - ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h - NAMESPACES mlir::sparse_tensor mlir::sparse_tensor::detail -) +#eudslpygen(EUDSLGen_sparse_tensor +# ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h +# NAMESPACES mlir::sparse_tensor mlir::sparse_tensor::detail +#) # nb::detail::nb_func_new("get_vce_triple_attr_name"): mismatched static/instance method flags in function overloads! # eudslpygen(EUDSLGen_spirv @@ -341,30 +341,30 @@ eudslpygen(EUDSLGen_tosa NAMESPACES mlir::tosa mlir::tosa::detail ) -eudslpygen(EUDSLGen_transform - ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h - NAMESPACES mlir::transform mlir::transform::detail -) +#eudslpygen(EUDSLGen_transform +# ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h +# NAMESPACES mlir::transform mlir::transform::detail +#) -eudslpygen(EUDSLGen_ub - ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h - NAMESPACES mlir::ub mlir::ub::detail -) +#eudslpygen(EUDSLGen_ub +# ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h +# NAMESPACES mlir::ub mlir::ub::detail +#) #eudslpygen(EUDSLGen_vector # ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h # NAMESPACES mlir::vector mlir::vector::detail #) -eudslpygen(EUDSLGen_x86vector - ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h - NAMESPACES mlir::x86vector mlir::x86vector::detail -) +#eudslpygen(EUDSLGen_x86vector +# ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h +# NAMESPACES mlir::x86vector mlir::x86vector::detail +#) -eudslpygen(EUDSLGen_xegpu - ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h - NAMESPACES mlir::xegpu mlir::xegpu::detail -) +#eudslpygen(EUDSLGen_xegpu +# ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h +# NAMESPACES mlir::xegpu mlir::xegpu::detail +#) nanobind_add_module(eudslpy_ext NB_STATIC @@ -427,7 +427,8 @@ set_target_properties(eudslpy_ext ) target_link_libraries(eudslpy_ext PRIVATE MLIR) # hack - on GHA, linux, the build OOMs -if ("$ENV{DISABLE_COMPILE_OPT}" MATCHES "true") +set(EUDSLPY_DISABLE_COMPILE_OPT ON CACHE BOOL "") +if (EUDSLPY_DISABLE_COMPILE_OPT) target_compile_options(eudslpy_ext PRIVATE -O0) endif() diff --git a/projects/eudsl-py/pyproject.toml b/projects/eudsl-py/pyproject.toml index 60fb1786..ad121a33 100644 --- a/projects/eudsl-py/pyproject.toml +++ b/projects/eudsl-py/pyproject.toml @@ -37,6 +37,7 @@ Clang_DIR = { env = "Clang_DIR", default = "EMPTY" } CMAKE_PREFIX_PATH = { env = "CMAKE_PREFIX_PATH", default = "" } CMAKE_C_COMPILER_LAUNCHER = { env = "CMAKE_C_COMPILER_LAUNCHER", default = "" } CMAKE_CXX_COMPILER_LAUNCHER = { env = "CMAKE_CXX_COMPILER_LAUNCHER", default = "" } +EUDSLPY_DISABLE_COMPILE_OPT = { env = "EUDSLPY_DISABLE_COMPILE_OPT", default = "ON" } CMAKE_CXX_VISIBILITY_PRESET = "hidden" CMAKE_VERBOSE_MAKEFILE = "ON" From 732cc312cbe24d8ae400944e27a9429f3f325592 Mon Sep 17 00:00:00 2001 From: max Date: Fri, 3 Jan 2025 00:43:39 -0500 Subject: [PATCH 6/6] root CMakeLists.txt --- .../workflows/build_test_release_eudsl.yml | 29 ++- CMakeLists.txt | 75 +++++++ README.md | 64 +++++- projects/CMakeLists.txt | 10 + projects/eudsl-nbgen/CMakeLists.txt | 42 ++-- .../cmake/eudsl_nbgen-config.cmake | 190 ++++++++++++------ projects/eudsl-nbgen/src/eudsl-nbgen.cpp | 22 +- projects/eudsl-py/CMakeLists.txt | 120 +++++------ projects/eudsl-tblgen/CMakeLists.txt | 48 +++-- projects/eudsl-tblgen/requirements.txt | 3 - .../requirements.txt => requirements.txt | 1 - 11 files changed, 420 insertions(+), 184 deletions(-) create mode 100644 CMakeLists.txt create mode 100644 projects/CMakeLists.txt delete mode 100644 projects/eudsl-tblgen/requirements.txt rename projects/eudsl-py/requirements.txt => requirements.txt (67%) diff --git a/.github/workflows/build_test_release_eudsl.yml b/.github/workflows/build_test_release_eudsl.yml index ce31af07..02753e91 100644 --- a/.github/workflows/build_test_release_eudsl.yml +++ b/.github/workflows/build_test_release_eudsl.yml @@ -108,7 +108,6 @@ jobs: echo "LLVM_DIR=/host/$PWD/llvm-install/lib/cmake/llvm" >> $GITHUB_ENV echo "MLIR_DIR=/host/$PWD/llvm-install/lib/cmake/mlir" >> $GITHUB_ENV echo "Clang_DIR=/host/$PWD/llvm-install/lib/cmake/clang" >> $GITHUB_ENV - echo "CCACHE_DIR=/host/$CCACHE_DIR" >> $GITHUB_ENV echo "PIP_FIND_LINKS=/host/$PWD/wheelhouse" >> $GITHUB_ENV else echo "LLVM_DIR=$PWD/llvm-install/lib/cmake/llvm" >> $GITHUB_ENV @@ -131,27 +130,47 @@ jobs: echo MACOSX_DEPLOYMENT_TARGET=13.7 >> $GITHUB_ENV fi + # prevent OOM on free GHA + echo "EUDSLPY_DISABLE_COMPILE_OPT=${{ inputs.eudslpy_disable_compile_opt == 'true' && 'ON' || 'OFF' }}" >> $GITHUB_ENV + pip install cibuildwheel - name: "Build eudsl-tblgen" run: | + if [[ "${{ matrix.os }}" == "ubuntu" ]]; then + export CCACHE_DIR=/host/$CCACHE_DIR + fi $python3_command -m cibuildwheel "$PWD/projects/eudsl-tblgen" --output-dir wheelhouse - name: "Build eudsl-nbgen" - if: ${{ ! startsWith(matrix.os, 'windows') }} run: | + if [[ "${{ matrix.os }}" == "ubuntu" ]]; then + export CCACHE_DIR=/host/$CCACHE_DIR + fi $python3_command -m cibuildwheel "$PWD/projects/eudsl-nbgen" --output-dir wheelhouse - name: "Build eudsl-py" if: ${{ ! startsWith(matrix.os, 'windows') }} run: | - # prevent OOM on free GHA - export EUDSLPY_DISABLE_COMPILE_OPT="${{ inputs.eudslpy_disable_compile_opt == 'true' && 'ON' || 'OFF' }}" + if [[ "${{ matrix.os }}" == "ubuntu" ]]; then + export CCACHE_DIR=/host/$CCACHE_DIR + fi $python3_command -m cibuildwheel "$PWD/projects/eudsl-py" --output-dir wheelhouse + # just to/make sure total build continues to work + - name: "Build all of eudsl" + run: | + + pip install -r requirements.txt + cmake -B $PWD/eudsl-build -S $PWD \ + -DCMAKE_PREFIX_PATH=$PWD/llvm-install \ + -DCMAKE_BUILD_TYPE=Release \ + -DCMAKE_INSTALL_PREFIX=$PWD/eudsl-install + cmake --build "$PWD/eudsl-build" --target install + - name: "Save cache" uses: actions/cache/save@v3 if: ${{ !cancelled() && github.event_name == 'push' && github.ref_name == 'main' }} @@ -191,7 +210,7 @@ jobs: runs-on: ${{ matrix.runs-on }} - name: "Test eudsl ${{ matrix.name }}" + name: "Test eudsl-tblgen ${{ matrix.name }}" defaults: run: diff --git a/CMakeLists.txt b/CMakeLists.txt new file mode 100644 index 00000000..4f423581 --- /dev/null +++ b/CMakeLists.txt @@ -0,0 +1,75 @@ +# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +# See https://llvm.org/LICENSE.txt for license information. +# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +# Copyright (c) 2024. + +cmake_minimum_required(VERSION 3.29) +set(CMAKE_CXX_STANDARD 17) +set(CMAKE_EXPORT_COMPILE_COMMANDS ON) + +set(LLVM_SUBPROJECT_TITLE "EUDSL") +set(EUDSL_BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR}) + +if (CMAKE_SOURCE_DIR STREQUAL CMAKE_CURRENT_SOURCE_DIR) + message("Building ${LLVM_SUBPROJECT_TITLE} as a standalone project.") + project(${LLVM_SUBPROJECT_TITLE} CXX C) + set(EUDSL_STANDALONE_BUILD ON) +else() + enable_language(CXX C) + set(EUDSL_STANDALONE_BUILD OFF) +endif() + +if(EUDSL_STANDALONE_BUILD) + find_package(LLVM REQUIRED CONFIG) + find_package(MLIR REQUIRED CONFIG) + find_package(Clang REQUIRED CONFIG PATHS "${LLVM_BINARY_DIR}" NO_DEFAULT_PATH NO_CMAKE_FIND_ROOT_PATH) + + message(STATUS "Using LLVMConfig.cmake in: ${LLVM_DIR}") + message(STATUS "Using MLIRConfig.cmake in: ${MLIR_DIR}") + message(STATUS "Using ClangConfig.cmake in: ${Clang_DIR}") + + set(LLVM_RUNTIME_OUTPUT_INTDIR ${CMAKE_BINARY_DIR}/bin) + set(LLVM_LIBRARY_OUTPUT_INTDIR ${CMAKE_BINARY_DIR}/lib) + set(MLIR_BINARY_DIR ${CMAKE_BINARY_DIR}) + + list(APPEND CMAKE_MODULE_PATH "${MLIR_CMAKE_DIR}") + list(APPEND CMAKE_MODULE_PATH "${LLVM_CMAKE_DIR}") + list(APPEND CMAKE_MODULE_PATH "${CLANG_CMAKE_DIR}") + + include(TableGen) + include(AddLLVM) + include(AddMLIR) + include(AddClang) + + set(MLIR_INCLUDE_DIR ${MLIR_INCLUDE_DIRS}) +else() + # turning LLVM -DLLVM_OPTIMIZED_TABLEGEN=ON builds some stuff in the NATIVE dir + # but not everything so LLVM_BINARY_DIR isn't correct + string(REPLACE "NATIVE" "" LLVM_BINARY_DIR "${LLVM_BINARY_DIR}") + # Build via external projects mechanism + set(LLVM_INCLUDE_DIR ${LLVM_MAIN_SRC_DIR}/include) + set(LLVM_GENERATED_INCLUDE_DIR ${LLVM_BINARY_DIR}/include) + set(LLVM_INCLUDE_DIRS "${LLVM_INCLUDE_DIR};${LLVM_GENERATED_INCLUDE_DIR}") + + set(MLIR_MAIN_SRC_DIR ${LLVM_MAIN_SRC_DIR}/../mlir) + set(MLIR_INCLUDE_DIR ${MLIR_MAIN_SRC_DIR}/include) + set(MLIR_GENERATED_INCLUDE_DIR ${LLVM_BINARY_DIR}/tools/mlir/include) + set(MLIR_INCLUDE_DIRS "${MLIR_INCLUDE_DIR};${MLIR_GENERATED_INCLUDE_DIR}") + + set(CLANG_MAIN_SRC_DIR ${LLVM_MAIN_SRC_DIR}/../clang) + set(CLANG_INCLUDE_DIR ${CLANG_MAIN_SRC_DIR}/include) + set(CLANG_GENERATED_INCLUDE_DIR ${LLVM_BINARY_DIR}/tools/clang/include) + set(CLANG_INCLUDE_DIRS "${CLANG_INCLUDE_DIR};${CLANG_GENERATED_INCLUDE_DIR}") +endif() + +include_directories(${LLVM_INCLUDE_DIRS}) +include_directories(${MLIR_INCLUDE_DIRS}) +include_directories(${CLANG_INCLUDE_DIRS}) +link_directories(${LLVM_BUILD_LIBRARY_DIR}) +add_definitions(${LLVM_DEFINITIONS}) + +if(NOT TARGET LLVMSupport) + message(FATAL_ERROR "LLVMSupport not found") +endif() + +add_subdirectory(projects) \ No newline at end of file diff --git a/README.md b/README.md index 1208cabe..433a611b 100644 --- a/README.md +++ b/README.md @@ -7,6 +7,68 @@ This repository contains the source code for `EUDSL`, a toolkit for the construction of embedded DSLs, in arbitrary languages, for targeting [MLIR](https://mlir.llvm.org). +FYI: this project is currently "alpha" quality. + +Currently, there are three components: + +1. [eudsl-tblgen](./projects/eudsl-tblgen): Python bindings to [LLVM's TableGen library](https://github.com/llvm/llvm-project/tree/659192b1843c4af180700783caca4cdc7afa3eab/llvm/lib/TableGen); +2. [eudsl-nbgen](./projects/eudsl-nbgen): A source-to-source translator that translates MLIR headers[^1] into direct `nanobind` bindings; +3. [eudsl-py](./projects/eudsl-py): Direct Python bindings to MLIR, generated using `eudsl-nbgen`; + * Currently only TableGen outputs (various `*.h.inc` files generated by `mlir-tblgen`) are automatically translated but the `eudsl-nbgen` tool can be manually run to translate any MLIR header (to varying degrees of success); + * See [projects/eudsl-py/tests](./projects/eudsl-py/tests) for a few examples of what these bindings look like. + ## Getting started -Coming soon... +Python wheels of all the tools are available at the [`latest` release page](https://github.com/llvm/eudsl/releases/tag/latest). +They are also `pip install`-able with .e.g + +```shell +$ cd +$ pip install eudsl-py -f https://github.com/llvm/eudsl/releases/expanded_assets/latest +``` + +`eudsl-py` has a slowly growing set of tests @ [projects/eudsl-py/tests](./projects/eudsl-py/tests). + +If you don't want to install locally, here is a [colab notebook minimal working example](https://colab.research.google.com/drive/1l-6rVnsUM3ypn7rKcaF_V6XVdopEM4Df?usp=sharing). + +## Developing/building + +**Strong recommendation**: check the CI scripts @ [.github/workflows](.github/workflows) - they do a fresh checkout and build on every commit and are written to be read by a non-CI expert. + +Firstly, you need a distribution of LLVM. You can build LLVM from source using our submodule by doing (on Mac/Linux or mingw): + +```shell +$ cd +$ ./build_tools/build_llvm.sh +``` + +Alternatively you can download a distribution for your platform from the [`llvm` release page](https://github.com/llvm/eudsl/releases/tag/llvm). + +Then each of the above tools can both be built as a conventional, standalone, CMake project and as a Python wheel. +The wheel build looks something like: + +```shell +$ cd +$ export CMAKE_PREFIX_PATH=$PWD/llvm-install +$ pip wheel projects/eudsl-nbgen -w wheelhouse -v +$ pip wheel projects/eudsl-py -w wheelhouse -v --find-links $PWD/wheelhouse +``` + +Note, the trailing `--find-links $PWD/wheelhouse` on `pip wheel projects/eudsl-py` is because `eudsl-nbgen` is a dependency of `eudsl-py` (that can be satisfied using the `eudsl-nbgen` wheel). + +If you want to build an individual tool via CMake you can do something like: + +```shell +$ cd +$ pip install -r requirements.txt +$ export CMAKE_PREFIX_PATH=$PWD/llvm-install +$ cmake -B $PWD/eudsl-nbgen-build -S $PWD/projects/eudsl-nbgen -DCMAKE_BUILD_TYPE=Release -DCMAKE_INSTALL_PREFIX=$PWD/eudsl-nbgen-install +$ cmake --build "$PWD/eudsl-build" --target install +``` + +If you want to build all the tools at once using CMake you can use the root [`CMakeLists.txt`](./CMakeLists.txt). +Note, in this case, `eudsl-nbgen` will automatically be built prior to `eudsl-py`. + +## Footnotes + +[^1]: Yes C++ headers... \ No newline at end of file diff --git a/projects/CMakeLists.txt b/projects/CMakeLists.txt new file mode 100644 index 00000000..e3b16175 --- /dev/null +++ b/projects/CMakeLists.txt @@ -0,0 +1,10 @@ +# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +# See https://llvm.org/LICENSE.txt for license information. +# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +# Copyright (c) 2024. + +if(NOT WIN32) + add_subdirectory(eudsl-py) +endif() +add_subdirectory(eudsl-nbgen) +add_subdirectory(eudsl-tblgen) diff --git a/projects/eudsl-nbgen/CMakeLists.txt b/projects/eudsl-nbgen/CMakeLists.txt index b5faa3f3..a63b9287 100644 --- a/projects/eudsl-nbgen/CMakeLists.txt +++ b/projects/eudsl-nbgen/CMakeLists.txt @@ -30,25 +30,7 @@ if(EUDSL_NBGEN_STANDALONE_BUILD) include(TableGen) include(AddLLVM) include(AddClang) - # TODO(max): probably don't need this anymore after landing the nanobind fix? - # technically we need this on windows too but our LLVM is compiled without exception handling - # and that breaks windows - if(NOT WIN32) - include(HandleLLVMOptions) - endif() -else() - # turning LLVM -DLLVM_OPTIMIZED_TABLEGEN=ON builds some stuff in the NATIVE dir - # but not everything so LLVM_BINARY_DIR isn't correct - string(REPLACE "NATIVE" "" LLVM_BINARY_DIR "${LLVM_BINARY_DIR}") - # Build via external projects mechanism - set(LLVM_INCLUDE_DIR ${LLVM_MAIN_SRC_DIR}/include) - set(LLVM_GENERATED_INCLUDE_DIR ${LLVM_BINARY_DIR}/include) - set(LLVM_INCLUDE_DIRS "${LLVM_INCLUDE_DIR};${LLVM_GENERATED_INCLUDE_DIR}") - - set(CLANG_MAIN_SRC_DIR ${LLVM_MAIN_SRC_DIR}/../clang) - set(CLANG_INCLUDE_DIR ${CLANG_MAIN_SRC_DIR}/include) - set(CLANG_GENERATED_INCLUDE_DIR ${LLVM_BINARY_DIR}/tools/clang/include) - set(CLANG_INCLUDE_DIRS "${CLANG_INCLUDE_DIR};${CLANG_GENERATED_INCLUDE_DIR}") + include(HandleLLVMOptions) endif() include_directories(${LLVM_INCLUDE_DIRS}) @@ -72,13 +54,25 @@ target_link_libraries(eudsl-nbgen ) string(TOLOWER ${LLVM_SUBPROJECT_TITLE} EUDSL_NBGEN_INSTALL_DATADIR) -# https://github.com/scikit-build/scikit-build-core/blob/a887a9b6c057b4ce9d3cfd53ae24e73caf1395a2/docs/build.md?plain=1#L139-L148 -# actually installs to venv/bin -install(TARGETS eudsl-nbgen RUNTIME DESTINATION "${EUDSL_NBGEN_INSTALL_DATADIR}.data/scripts") -# this actually installs into venv/lib -install(IMPORTED_RUNTIME_ARTIFACTS LLVM LIBRARY DESTINATION "${EUDSL_NBGEN_INSTALL_DATADIR}.data/data/lib") +if (NOT "$ENV{PIP_BUILD_TRACKER}" STREQUAL "") + # pip install + # actually installs to venv/bin + # https://github.com/scikit-build/scikit-build-core/blob/a887a9b6c057b4ce9d3cfd53ae24e73caf1395a2/docs/build.md?plain=1#L139-L148 + install(TARGETS eudsl-nbgen RUNTIME DESTINATION "${EUDSL_NBGEN_INSTALL_DATADIR}.data/scripts") + if (NOT WIN32) + # this actually installs into venv/lib + install(IMPORTED_RUNTIME_ARTIFACTS LLVM LIBRARY DESTINATION "${EUDSL_NBGEN_INSTALL_DATADIR}.data/data/lib") + endif() +else() + # pip cmake install + install(TARGETS eudsl-nbgen RUNTIME DESTINATION "${EUDSL_NBGEN_INSTALL_DATADIR}/bin") + if (NOT WIN32) + install(IMPORTED_RUNTIME_ARTIFACTS LLVM LIBRARY DESTINATION "${EUDSL_NBGEN_INSTALL_DATADIR}/lib") + endif() +endif() install( + # the slash is load-bearing... DIRECTORY src/ DESTINATION "${EUDSL_NBGEN_INSTALL_DATADIR}" FILES_MATCHING PATTERN "*\.py" diff --git a/projects/eudsl-nbgen/cmake/eudsl_nbgen-config.cmake b/projects/eudsl-nbgen/cmake/eudsl_nbgen-config.cmake index 235e8706..2548f4e5 100644 --- a/projects/eudsl-nbgen/cmake/eudsl_nbgen-config.cmake +++ b/projects/eudsl-nbgen/cmake/eudsl_nbgen-config.cmake @@ -5,25 +5,25 @@ # copy-pasta from AddMLIR.cmake/AddLLVM.cmake/TableGen.cmake -function(eudslpygen target input_file) - set(EUDSLPYGEN_TARGET_DEFINITIONS ${input_file}) - cmake_parse_arguments(ARG "" "" "DEPENDS;EXTRA_INCLUDES;NAMESPACES" ${ARGN}) - if (IS_ABSOLUTE ${EUDSLPYGEN_TARGET_DEFINITIONS}) - set(EUDSLPYGEN_TARGET_DEFINITIONS_ABSOLUTE ${input_file}) +function(eudsl_nbgen target input_file) + set(EUDSL_NBGEN_TARGET_DEFINITIONS ${input_file}) + cmake_parse_arguments(ARG "" "" "LINK_LIBS;EXTRA_INCLUDES;NAMESPACES" ${ARGN}) + if (IS_ABSOLUTE ${EUDSL_NBGEN_TARGET_DEFINITIONS}) + set(EUDSL_NBGEN_TARGET_DEFINITIONS_ABSOLUTE ${input_file}) else() - set(EUDSLPYGEN_TARGET_DEFINITIONS_ABSOLUTE ${CMAKE_CURRENT_SOURCE_DIR}/${input_file}) + set(EUDSL_NBGEN_TARGET_DEFINITIONS_ABSOLUTE ${CMAKE_CURRENT_SOURCE_DIR}/${input_file}) endif() - if(NOT EXISTS "${EUDSLPYGEN_TARGET_DEFINITIONS_ABSOLUTE}") + if (NOT EXISTS "${EUDSL_NBGEN_TARGET_DEFINITIONS_ABSOLUTE}") message(FATAL_ERROR "${input_file} does not exist") endif() - get_directory_property(eudslpygen_includes INCLUDE_DIRECTORIES) + get_directory_property(eudsl_nbgen_includes INCLUDE_DIRECTORIES) list(TRANSFORM ARG_EXTRA_INCLUDES PREPEND -I) - list(APPEND eudslpygen_includes ${CMAKE_CXX_IMPLICIT_INCLUDE_DIRECTORIES}) - list(REMOVE_ITEM eudslpygen_includes "") - list(TRANSFORM eudslpygen_includes PREPEND -I) - list(APPEND eudslpygen_includes ${ARG_EXTRA_INCLUDES}) + list(APPEND eudsl_nbgen_includes ${CMAKE_CXX_IMPLICIT_INCLUDE_DIRECTORIES}) + list(REMOVE_ITEM eudsl_nbgen_includes "") + list(TRANSFORM eudsl_nbgen_includes PREPEND -I) + list(APPEND eudsl_nbgen_includes ${ARG_EXTRA_INCLUDES}) set(_gen_target_dir "${CMAKE_CURRENT_BINARY_DIR}/generated/${target}") file(MAKE_DIRECTORY ${_gen_target_dir}) @@ -33,61 +33,107 @@ function(eudslpygen target input_file) # hack but most of the time we're loading headers that are downstream of tds anyway # this could be smarter by asking people to list td targets or something but that's too onerous file(GLOB_RECURSE global_tds "${MLIR_INCLUDE_DIR}/mlir/*.td") - # use cc -MM to collect all transitive headers - set(clang_command ${CMAKE_CXX_COMPILER} - # -v - -xc++ "-std=c++${CMAKE_CXX_STANDARD}" - -MM ${EUDSLPYGEN_TARGET_DEFINITIONS_ABSOLUTE} - -MT ${_full_gen_file} - ${eudslpygen_includes} - -o ${_depfile} - ) - execute_process(COMMAND ${clang_command} RESULT_VARIABLE _had_error_depfile - # COMMAND_ECHO STDOUT - ) + if (NOT EXISTS ${_depfile}) + # use cc -MM to collect all transitive headers + set(clang_command ${CMAKE_CXX_COMPILER} + # -v + -xc++ "-std=c++${CMAKE_CXX_STANDARD}" + -MM ${EUDSL_NBGEN_TARGET_DEFINITIONS_ABSOLUTE} + -MT ${_full_gen_file} + ${eudsl_nbgen_includes} + -o ${_depfile} + ) + execute_process(COMMAND ${clang_command} RESULT_VARIABLE _had_error_depfile + # COMMAND_ECHO STDOUT + ) + endif() - if (IS_ABSOLUTE ${EUDSLPYGEN_TARGET_DEFINITIONS}) - set(EUDSLPYGEN_TARGET_DEFINITIONS_ABSOLUTE ${EUDSLPYGEN_TARGET_DEFINITIONS}) + if (IS_ABSOLUTE ${EUDSL_NBGEN_TARGET_DEFINITIONS}) + set(EUDSL_NBGEN_TARGET_DEFINITIONS_ABSOLUTE ${EUDSL_NBGEN_TARGET_DEFINITIONS}) else() - set(EUDSLPYGEN_TARGET_DEFINITIONS_ABSOLUTE - ${CMAKE_CURRENT_SOURCE_DIR}/${EUDSLPYGEN_TARGET_DEFINITIONS}) + set(EUDSL_NBGEN_TARGET_DEFINITIONS_ABSOLUTE + ${CMAKE_CURRENT_SOURCE_DIR}/${EUDSL_NBGEN_TARGET_DEFINITIONS}) endif() - string(REPLACE " " ";" eudslpygen_defines "${LLVM_DEFINITIONS}") + string(REPLACE " " ";" eudsl_nbgen_defines "${LLVM_DEFINITIONS}") list(JOIN ARG_NAMESPACES "," namespaces) - find_program(EUDSLPY_EUDSLPYGEN_EXE "eudsl-nbgen" REQUIRED) - execute_process( - COMMAND ${EUDSLPY_EUDSLPYGEN_EXE} ${EUDSLPYGEN_TARGET_DEFINITIONS_ABSOLUTE} - -I${CMAKE_CURRENT_SOURCE_DIR} - -namespaces=${namespaces} - ${eudslpygen_includes} - ${eudslpygen_defines} - -o "${_full_gen_file}" - WORKING_DIRECTORY ${CMAKE_BINARY_DIR} - RESULT_VARIABLE _had_error_gen_cpp - # COMMAND_ECHO STDOUT + set(eudsl_nbgen_generate_cmd + ${EUDSL_NBGEN_TARGET_DEFINITIONS_ABSOLUTE} + -I${CMAKE_CURRENT_LIST_DIR} -namespaces=${namespaces} + ${eudsl_nbgen_includes} ${eudsl_nbgen_defines} + -o "${_full_gen_file}" ) - if((_had_error_gen_cpp AND NOT _had_error_gen_cpp EQUAL 0) OR NOT EXISTS "${_full_gen_file}") - message(FATAL_ERROR "failed to create ${_full_gen_file}: ${_had_error_gen_cpp}") - endif() - # this is the specific thing connected the dependencies... - set_property(DIRECTORY APPEND PROPERTY CMAKE_CONFIGURE_DEPENDS ${_full_gen_file}) - execute_process( - COMMAND ${EUDSLPY_EUDSLPYGEN_EXE} -shardify ${_full_gen_file} - # ARG_EXTRA_INCLUDES has already had -I prepended - -shard-target ${target} ${ARG_EXTRA_INCLUDES} -I ${EUDSLPYGEN_TARGET_DEFINITIONS_ABSOLUTE} - WORKING_DIRECTORY ${CMAKE_BINARY_DIR} - RESULT_VARIABLE _had_error_gen_sharded - # COMMAND_ECHO STDOUT + set(eudsl_nbgen_shardify_cmd + -shardify ${_full_gen_file} + # ARG_EXTRA_INCLUDES has already had -I prepended + -shard-target ${target} ${ARG_EXTRA_INCLUDES} -I ${EUDSL_NBGEN_TARGET_DEFINITIONS_ABSOLUTE} ) - if((_had_error_gen_sharded AND NOT _had_error_gen_sharded EQUAL 0) OR NOT EXISTS "${_full_gen_file}.sharded.cpp") - message(FATAL_ERROR "failed to create ${_full_gen_file}.sharded.cpp: ${_had_error_gen_sharded}") - endif() - set_property(DIRECTORY APPEND PROPERTY CMAKE_CONFIGURE_DEPENDS "${_full_gen_file}.sharded.cpp") - file(GLOB _shards CONFIGURE_DEPENDS "${_gen_target_dir}/*shard*cpp") - if(NOT _shards) - message(FATAL_ERROR "no shards created") + + find_program(EUDSL_NBGEN_EXE "eudsl-nbgen") + if (EUDSL_NBGEN_EXE STREQUAL "EUDSL_NBGEN_EXE-NOTFOUND") + ################################## + # not standalone build + ################################## + if (WIN32) + set(EUDSL_NBGEN_EXE "eudsl-nbgen.exe") + else() + set(EUDSL_NBGEN_EXE "eudsl-nbgen") + endif() + + string(REPLACE " " ";" eudsl_nbgen_defines "${LLVM_DEFINITIONS}") + list(JOIN ARG_NAMESPACES "," namespaces) + + add_custom_command(OUTPUT ${_full_gen_file} + COMMAND ${EUDSL_NBGEN_EXE} ${eudsl_nbgen_generate_cmd} + WORKING_DIRECTORY ${CMAKE_BINARY_DIR} + DEPENDS ${EUDSL_NBGEN_EXE} ${global_tds} + DEPFILE ${_depfile} + COMMENT "eudsl-nbgen: Generating ${_full_gen_file}..." + ) + # epic hack to specify all shards that will be generated even though we don't know them before hand + set(_shards) + # lol spirv has 260 ops + set(_max_num_shards 30) + # note this count [0, 30] <- inclusive + foreach(i RANGE ${_max_num_shards}) + list(APPEND _shards "${_full_gen_file}.shard.${i}.cpp") + endforeach() + + add_custom_command(OUTPUT "${_full_gen_file}.sharded.cpp" + COMMAND ${EUDSL_NBGEN_EXE} ${eudsl_nbgen_shardify_cmd} + -max-number-shards ${_max_num_shards} + BYPRODUCTS ${_shards} + WORKING_DIRECTORY ${CMAKE_BINARY_DIR} + DEPENDS ${_full_gen_file} ${EUDSL_NBGEN_EXE} + COMMENT "eudsl-nbgen: Generating ${_full_gen_file}.sharded.cpp..." + ) + else() + ################################## + # standalone build + ################################## + execute_process( + COMMAND ${EUDSL_NBGEN_EXE} ${eudsl_nbgen_generate_cmd} + WORKING_DIRECTORY ${CMAKE_BINARY_DIR} + RESULT_VARIABLE _had_error_gen_cpp + # COMMAND_ECHO STDOUT + ) + if ((_had_error_gen_cpp AND NOT _had_error_gen_cpp EQUAL 0) OR NOT EXISTS "${_full_gen_file}") + message(FATAL_ERROR "failed to create ${_full_gen_file}: ${_had_error_gen_cpp}") + endif() + execute_process( + COMMAND ${EUDSL_NBGEN_EXE} ${eudsl_nbgen_shardify_cmd} + WORKING_DIRECTORY ${CMAKE_BINARY_DIR} + RESULT_VARIABLE _had_error_gen_sharded + # COMMAND_ECHO STDOUT + ) + if ((_had_error_gen_sharded AND NOT _had_error_gen_sharded EQUAL 0) OR NOT EXISTS "${_full_gen_file}.sharded.cpp") + message(FATAL_ERROR "failed to create ${_full_gen_file}.sharded.cpp: ${_had_error_gen_sharded}") + endif() + file(GLOB _shards CONFIGURE_DEPENDS "${_gen_target_dir}/*shard*cpp") + if (NOT _shards) + message(FATAL_ERROR "no shards created") + endif() endif() add_library(${target} STATIC "${_full_gen_file}.sharded.cpp" ${_shards}) @@ -96,32 +142,42 @@ function(eudslpygen target input_file) OUTPUT_STRIP_TRAILING_WHITESPACE OUTPUT_VARIABLE nanobind_include_dir RESULT_VARIABLE _has_err_find_nanobind ) - if((_has_err_find_nanobind AND NOT _has_err_find_nanobind EQUAL 0) OR NOT EXISTS "${nanobind_include_dir}") + if ((_has_err_find_nanobind AND NOT _has_err_find_nanobind EQUAL 0) OR NOT EXISTS "${nanobind_include_dir}") message(FATAL_ERROR "couldn't find nanobind include dir: ${_has_err_find_nanobind}") endif() target_include_directories(${target} PRIVATE - ${eudslpygen_includes} + ${eudsl_nbgen_includes} ${Python_INCLUDE_DIRS} ${nanobind_include_dir} ) - target_compile_options(${target} PRIVATE -Wno-cast-qual) + # not sure why unix needs this buy not apple (and really only in root-cmake build...) + if(UNIX AND NOT APPLE) + set_property(TARGET ${target} PROPERTY POSITION_INDEPENDENT_CODE ON) + endif() + set(_link_libs ${ARG_LINK_LIBS}) + if (NOT ARG_LINK_LIBS) + set(_link_libs MLIR LLVM) + endif() + target_link_libraries(${target} PUBLIC ${_link_libs}) + target_compile_options(${target} PUBLIC -Wno-cast-qual) # `make clean' must remove all those generated files: # TODO(max): clean up dep files set_property(DIRECTORY ${CMAKE_CURRENT_LIST_DIR} - APPEND PROPERTY ADDITIONAL_MAKE_CLEAN_FILES ${_shards}) + APPEND PROPERTY ADDITIONAL_MAKE_CLEAN_FILES ${_shards} ${_depfile}) set_source_files_properties(${_shards} PROPERTIES GENERATED 1) endfunction() function(patch_mlir_llvm_rpath target) + cmake_parse_arguments(ARG "STANDALONE" "" "" ${ARGN}) # hack so we can move libMLIR and libLLVM into the wheel # see AddLLVM.cmake#llvm_setup_rpath - if(APPLE OR UNIX) + if (APPLE OR UNIX) set(_origin_prefix "\$ORIGIN") - if(APPLE) + if (APPLE) set(_origin_prefix "@loader_path") endif() - if (EUDSLPY_STANDALONE_BUILD) + if (STANDALONE) get_target_property(_mlir_loc MLIR LOCATION) get_target_property(_llvm_loc LLVM LOCATION) else() @@ -129,7 +185,7 @@ function(patch_mlir_llvm_rpath target) set(_llvm_loc "$") endif() set(_old_rpath "${_origin_prefix}/../lib${LLVM_LIBDIR_SUFFIX}") - if(APPLE) + if (APPLE) if (EXISTS ${_mlir_loc}) execute_process(COMMAND install_name_tool -rpath "${_old_rpath}" ${_origin_prefix} "${_mlir_loc}" ERROR_VARIABLE rpath_err) endif() diff --git a/projects/eudsl-nbgen/src/eudsl-nbgen.cpp b/projects/eudsl-nbgen/src/eudsl-nbgen.cpp index 6db162f1..de053aac 100644 --- a/projects/eudsl-nbgen/src/eudsl-nbgen.cpp +++ b/projects/eudsl-nbgen/src/eudsl-nbgen.cpp @@ -16,8 +16,10 @@ #include "llvm/Support/ToolOutputFile.h" #include +#include #include #include +#include static llvm::cl::OptionCategory EUDSLPYGenCat("Options for eudslpy-gen"); static llvm::cl::opt InputFilename(llvm::cl::Positional, @@ -737,6 +739,14 @@ static int generate() { return 0; } +// cannot fathom why but std::to_string is return ascii characters +template +static std::string numberToString(T number) { + std::ostringstream ss; + ss << number; + return ss.str(); +} + static int makeSourceShards() { // Ensure the filename ends with ".cpp.gen" if (InputFilename.substr(InputFilename.size() - 8) != ".cpp.gen") { @@ -769,10 +779,9 @@ static int makeSourceShards() { : ShardTarget; // Generate shard files - for (size_t i = 0; i < shards.size(); ++i) { std::ofstream shardFile(InputFilename.getValue() + ".shard." + - std::to_string(i) + ".cpp"); + numberToString(i) + ".cpp"); if (!shardFile.is_open()) { llvm::errs() << "Failed to create shard file"; return -1; @@ -803,16 +812,19 @@ void populate)" << finalTarget << i << R"(Module(nb::module_ &m) { // Handle max number of shards if (MaxNumShards != -1 && shards.size() > static_cast(MaxNumShards)) { - llvm::errs() << "expected less than " + std::to_string(MaxNumShards) + - " shards"; + llvm::errs() << "expected less than " + + numberToString(MaxNumShards.getValue()) + " shards"; return -1; } if (MaxNumShards == -1) MaxNumShards = shards.size(); + else + // _max_num_shards in cmake counts +1 + MaxNumShards.getValue() += 1; for (size_t i = shards.size(); i < static_cast(MaxNumShards); ++i) { std::ofstream dummyShardFile(InputFilename.getValue() + ".shard." + - std::to_string(i) + ".cpp"); + numberToString(i) + ".cpp"); if (!dummyShardFile.is_open()) { llvm::errs() << "Failed to create dummy shard file"; return -1; diff --git a/projects/eudsl-py/CMakeLists.txt b/projects/eudsl-py/CMakeLists.txt index 7cabdb46..bad1d4cf 100644 --- a/projects/eudsl-py/CMakeLists.txt +++ b/projects/eudsl-py/CMakeLists.txt @@ -53,24 +53,6 @@ if(EUDSLPY_STANDALONE_BUILD) # for out-of-tree MLIR_INCLUDE_DIR points to the build dir by default # and MLIR_INCLUDE_DIRS points to the correct place set(MLIR_INCLUDE_DIR ${MLIR_INCLUDE_DIRS}) -else() - # turning LLVM -DLLVM_OPTIMIZED_TABLEGEN=ON builds some stuff in the NATIVE dir - # but not everything so LLVM_BINARY_DIR isn't correct - string(REPLACE "NATIVE" "" LLVM_BINARY_DIR "${LLVM_BINARY_DIR}") - # Build via external projects mechanism - set(LLVM_INCLUDE_DIR ${LLVM_MAIN_SRC_DIR}/include) - set(LLVM_GENERATED_INCLUDE_DIR ${LLVM_BINARY_DIR}/include) - set(LLVM_INCLUDE_DIRS "${LLVM_INCLUDE_DIR};${LLVM_GENERATED_INCLUDE_DIR}") - - set(MLIR_MAIN_SRC_DIR ${LLVM_MAIN_SRC_DIR}/../mlir) - set(MLIR_INCLUDE_DIR ${MLIR_MAIN_SRC_DIR}/include) - set(MLIR_GENERATED_INCLUDE_DIR ${LLVM_BINARY_DIR}/tools/mlir/include) - set(MLIR_INCLUDE_DIRS "${MLIR_INCLUDE_DIR};${MLIR_GENERATED_INCLUDE_DIR}") - - set(CLANG_MAIN_SRC_DIR ${LLVM_MAIN_SRC_DIR}/../clang) - set(CLANG_INCLUDE_DIR ${CLANG_MAIN_SRC_DIR}/include) - set(CLANG_GENERATED_INCLUDE_DIR ${LLVM_BINARY_DIR}/tools/clang/include) - set(CLANG_INCLUDE_DIRS "${CLANG_INCLUDE_DIR};${CLANG_GENERATED_INCLUDE_DIR}") endif() include_directories(${LLVM_INCLUDE_DIRS}) @@ -107,15 +89,18 @@ include_directories(${EUDSLPY_SRC_DIR}) execute_process( COMMAND "${Python_EXECUTABLE}" -m eudsl_nbgen --cmake_dir OUTPUT_STRIP_TRAILING_WHITESPACE OUTPUT_VARIABLE eudsl_nbgen_DIR) -find_package(eudsl_nbgen CONFIG REQUIRED) +find_package(eudsl_nbgen CONFIG) +if(NOT eudsl_nbgen__FOUND) + include("${CMAKE_CURRENT_LIST_DIR}/../eudsl-nbgen/cmake/eudsl_nbgen-config.cmake") +endif() # too big -# eudslpygen(EUDSLGen_acc +# eudsl_nbgen(EUDSLGen_acc # ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h # NAMESPACES mlir::acc mlir::acc::detail # ) -eudslpygen(EUDSLGen_affine +eudsl_nbgen(EUDSLGen_affine ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::affine mlir::affine::detail EXTRA_INCLUDES @@ -123,50 +108,50 @@ eudslpygen(EUDSLGen_affine mlir/Dialect/Affine/IR/AffineValueMap.h ) -eudslpygen(EUDSLGen_amdgpu +eudsl_nbgen(EUDSLGen_amdgpu ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::amdgpu mlir::amdgpu::detail ) -#eudslpygen(EUDSLGen_amx +#eudsl_nbgen(EUDSLGen_amx # ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h # NAMESPACES mlir::amx mlir::amx::detail #) -eudslpygen(EUDSLGen_arith +eudsl_nbgen(EUDSLGen_arith ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::arith mlir::arith::detail ) -#eudslpygen(EUDSLGen_arm_neon +#eudsl_nbgen(EUDSLGen_arm_neon # ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h # NAMESPACES mlir::arm_neon mlir::arm_neon::detail #) # too big -# eudslpygen(EUDSLGen_arm_sme +# eudsl_nbgen(EUDSLGen_arm_sme # ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h # NAMESPACES mlir::arm_sme mlir::arm_sme::detail # ) -#eudslpygen(EUDSLGen_arm_sve +#eudsl_nbgen(EUDSLGen_arm_sve # ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h # NAMESPACES mlir::arm_sve mlir::arm_sve::detail #) -eudslpygen(EUDSLGen_async +eudsl_nbgen(EUDSLGen_async ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::async mlir::async::detail ) -eudslpygen(EUDSLGen_bufferization +eudsl_nbgen(EUDSLGen_bufferization ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::bufferization mlir::bufferization::detail EXTRA_INCLUDES mlir/Dialect/Bufferization/Transforms/Bufferize.h ) -eudslpygen(EUDSLGen_cf +eudsl_nbgen(EUDSLGen_cf ${MLIR_INCLUDE_DIR}/mlir/Dialect/ControlFlow/IR/ControlFlowOps.h NAMESPACES mlir::cf mlir::cf::detail EXTRA_INCLUDES @@ -174,27 +159,27 @@ eudslpygen(EUDSLGen_cf mlir/IR/PatternMatch.h ) -eudslpygen(EUDSLGen_complex +eudsl_nbgen(EUDSLGen_complex ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::complex mlir::complex::detail ) -#eudslpygen(EUDSLGen_DLTIDialect +#eudsl_nbgen(EUDSLGen_DLTIDialect # ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h # NAMESPACES mlir::DLTIDialect mlir::DLTIDialect::detail #) -eudslpygen(EUDSLGen_emitc +eudsl_nbgen(EUDSLGen_emitc ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::emitc mlir::emitc::detail ) -eudslpygen(EUDSLGen_func +eudsl_nbgen(EUDSLGen_func ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::func mlir::func::detail ) -eudslpygen(EUDSLGen_gpu +eudsl_nbgen(EUDSLGen_gpu ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::gpu mlir::gpu::detail EXTRA_INCLUDES @@ -202,7 +187,7 @@ eudslpygen(EUDSLGen_gpu llvm/IR/IRBuilder.h ) -eudslpygen(EUDSLGen_index +eudsl_nbgen(EUDSLGen_index ${MLIR_INCLUDE_DIR}/mlir/Dialect/Index/IR/IndexOps.h NAMESPACES mlir::index mlir::index::detail EXTRA_INCLUDES @@ -210,76 +195,76 @@ eudslpygen(EUDSLGen_index mlir/IR/PatternMatch.h ) -#eudslpygen(EUDSLGen_irdl +#eudsl_nbgen(EUDSLGen_irdl # ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h # NAMESPACES mlir::irdl mlir::irdl::detail #) -eudslpygen(EUDSLGen_linalg +eudsl_nbgen(EUDSLGen_linalg ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::linalg mlir::linalg::detail ) -eudslpygen(EUDSLGen_LLVM +eudsl_nbgen(EUDSLGen_LLVM ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::LLVM mlir::LLVM::detail ) -eudslpygen(EUDSLGen_math +eudsl_nbgen(EUDSLGen_math ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::math mlir::math::detail ) -eudslpygen(EUDSLGen_memref +eudsl_nbgen(EUDSLGen_memref ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::memref mlir::memref::detail ) -#eudslpygen(EUDSLGen_mesh +#eudsl_nbgen(EUDSLGen_mesh # ${MLIR_INCLUDE_DIR}/mlir/Dialect/Mesh/IR/MeshOps.h # NAMESPACES mlir::mesh mlir::mesh::detail # EXTRA_INCLUDES # mlir/Dialect/Mesh/IR/MeshOps.h #) -#eudslpygen(EUDSLGen_ml_program +#eudsl_nbgen(EUDSLGen_ml_program # ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h # NAMESPACES mlir::ml_program mlir::ml_program::detail #) -#eudslpygen(EUDSLGen_mpi +#eudsl_nbgen(EUDSLGen_mpi # ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h # NAMESPACES mlir::mpi mlir::mpi::detail #) -eudslpygen(EUDSLGen_nvgpu +eudsl_nbgen(EUDSLGen_nvgpu ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::nvgpu mlir::nvgpu::detail ) -eudslpygen(EUDSLGen_NVVM +eudsl_nbgen(EUDSLGen_NVVM ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::NVVM mlir::NVVM::detail ) -#eudslpygen(EUDSLGen_omp +#eudsl_nbgen(EUDSLGen_omp # ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h # NAMESPACES mlir::omp mlir::omp::detail #) -eudslpygen(EUDSLGen_pdl +eudsl_nbgen(EUDSLGen_pdl ${MLIR_INCLUDE_DIR}/mlir/Dialect/PDL/IR/PDLOps.h NAMESPACES mlir::pdl mlir::pdl::detail EXTRA_INCLUDES mlir/Dialect/PDL/IR/PDLOps.h ) -eudslpygen(EUDSLGen_pdl_interp +eudsl_nbgen(EUDSLGen_pdl_interp ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::pdl_interp mlir::pdl_interp::detail ) -eudslpygen(EUDSLGen_polynomial +eudsl_nbgen(EUDSLGen_polynomial ${MLIR_INCLUDE_DIR}/mlir/Dialect/Polynomial/IR/PolynomialOps.h NAMESPACES mlir::polynomial mlir::polynomial::detail EXTRA_INCLUDES @@ -287,7 +272,7 @@ eudslpygen(EUDSLGen_polynomial mlir/IR/PatternMatch.h ) -#eudslpygen(EUDSLGen_ptr +#eudsl_nbgen(EUDSLGen_ptr # ${MLIR_INCLUDE_DIR}/mlir/Dialect/Ptr/IR/PtrOps.h # NAMESPACES mlir::ptr mlir::ptr::detail # EXTRA_INCLUDES @@ -295,35 +280,35 @@ eudslpygen(EUDSLGen_polynomial # mlir/IR/DialectImplementation.h #) -#eudslpygen(EUDSLGen_quant +#eudsl_nbgen(EUDSLGen_quant # ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h # NAMESPACES mlir::quant mlir::quant::detail # EXTRA_INCLUDES # mlir/Dialect/Quant/IR/QuantTypes.h #) -eudslpygen(EUDSLGen_ROCDL +eudsl_nbgen(EUDSLGen_ROCDL ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::ROCDL mlir::ROCDL::detail ) -eudslpygen(EUDSLGen_scf +eudsl_nbgen(EUDSLGen_scf ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::scf mlir::scf::detail ) -eudslpygen(EUDSLGen_shape +eudsl_nbgen(EUDSLGen_shape ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::shape mlir::shape::detail ) -#eudslpygen(EUDSLGen_sparse_tensor +#eudsl_nbgen(EUDSLGen_sparse_tensor # ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h # NAMESPACES mlir::sparse_tensor mlir::sparse_tensor::detail #) # nb::detail::nb_func_new("get_vce_triple_attr_name"): mismatched static/instance method flags in function overloads! -# eudslpygen(EUDSLGen_spirv +# eudsl_nbgen(EUDSLGen_spirv # ${MLIR_INCLUDE_DIR}/mlir/Dialect/SPIRV/IR/SPIRVOps.h # NAMESPACES mlir::spirv mlir::spirv::detail # EXTRA_INCLUDES @@ -331,37 +316,37 @@ eudslpygen(EUDSLGen_shape # mlir/IR/PatternMatch.h # ) -eudslpygen(EUDSLGen_tensor +eudsl_nbgen(EUDSLGen_tensor ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::tensor mlir::tensor::detail ) -eudslpygen(EUDSLGen_tosa +eudsl_nbgen(EUDSLGen_tosa ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h NAMESPACES mlir::tosa mlir::tosa::detail ) -#eudslpygen(EUDSLGen_transform +#eudsl_nbgen(EUDSLGen_transform # ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h # NAMESPACES mlir::transform mlir::transform::detail #) -#eudslpygen(EUDSLGen_ub +#eudsl_nbgen(EUDSLGen_ub # ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h # NAMESPACES mlir::ub mlir::ub::detail #) -#eudslpygen(EUDSLGen_vector +#eudsl_nbgen(EUDSLGen_vector # ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h # NAMESPACES mlir::vector mlir::vector::detail #) -#eudslpygen(EUDSLGen_x86vector +#eudsl_nbgen(EUDSLGen_x86vector # ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h # NAMESPACES mlir::x86vector mlir::x86vector::detail #) -#eudslpygen(EUDSLGen_xegpu +#eudsl_nbgen(EUDSLGen_xegpu # ${MLIR_INCLUDE_DIR}/mlir/InitAllDialects.h # NAMESPACES mlir::xegpu mlir::xegpu::detail #) @@ -425,14 +410,17 @@ set_target_properties(eudslpy_ext PROPERTIES LIBRARY_OUTPUT_DIRECTORY "${EUDSLPY_SRC_DIR}/eudsl" ) -target_link_libraries(eudslpy_ext PRIVATE MLIR) # hack - on GHA, linux, the build OOMs set(EUDSLPY_DISABLE_COMPILE_OPT ON CACHE BOOL "") if (EUDSLPY_DISABLE_COMPILE_OPT) target_compile_options(eudslpy_ext PRIVATE -O0) endif() -patch_mlir_llvm_rpath(eudslpy_ext) +set(STANDALONE) +if(EUDSLPY_STANDALONE_BUILD OR EUDSL_STANDALONE_BUILD) + set(STANDALONE STANDALONE) +endif() +patch_mlir_llvm_rpath(eudslpy_ext ${STANDALONE}) # copy libMLIR into the ext dir for wheels install(IMPORTED_RUNTIME_ARTIFACTS MLIR LLVM LIBRARY DESTINATION eudsl) diff --git a/projects/eudsl-tblgen/CMakeLists.txt b/projects/eudsl-tblgen/CMakeLists.txt index d8ff2530..0396b6ea 100644 --- a/projects/eudsl-tblgen/CMakeLists.txt +++ b/projects/eudsl-tblgen/CMakeLists.txt @@ -4,17 +4,40 @@ # Copyright (c) 2024. cmake_minimum_required(VERSION 3.29) -project(eudsl_tblgen CXX C) set(CMAKE_CXX_STANDARD 17) +set(CMAKE_EXPORT_COMPILE_COMMANDS ON) + +set(LLVM_SUBPROJECT_TITLE "EUDSL_TBLGEN") +set(EUDSL_TBLGEN_BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR}) + +if (CMAKE_SOURCE_DIR STREQUAL CMAKE_CURRENT_LIST_DIR) + message("Building ${LLVM_SUBPROJECT_TITLE} as a standalone project.") + project(${LLVM_SUBPROJECT_TITLE} CXX C) + find_package(LLVM REQUIRED CONFIG) + + message(STATUS "Using LLVMConfig.cmake in: ${LLVM_DIR}") + + set(LLVM_RUNTIME_OUTPUT_INTDIR ${CMAKE_BINARY_DIR}/bin) + set(LLVM_LIBRARY_OUTPUT_INTDIR ${CMAKE_BINARY_DIR}/lib) + + list(APPEND CMAKE_MODULE_PATH "${LLVM_CMAKE_DIR}") + + include(TableGen) + include(AddLLVM) + # TODO(max): probably don't need this anymore after landing the nanobind fix? + # technically we need this on windows too but our LLVM is compiled without exception handling + # and that breaks windows + if(NOT WIN32) + include(HandleLLVMOptions) + endif() +endif() -find_package(LLVM REQUIRED CONFIG) -message(STATUS "Using LLVMConfig.cmake in: ${LLVM_DIR}") include_directories(${LLVM_INCLUDE_DIRS}) +link_directories(${LLVM_BUILD_LIBRARY_DIR}) +add_definitions(${LLVM_DEFINITIONS}) -# technically we need this but our LLVM is compiled without exception handling -# and that breaks windows -if(NOT WIN32) - include(HandleLLVMOptions) +if(NOT TARGET LLVMSupport) + message(FATAL_ERROR "LLVMSupport not found") endif() find_package(Python 3.8 @@ -26,7 +49,7 @@ execute_process( OUTPUT_STRIP_TRAILING_WHITESPACE OUTPUT_VARIABLE nanobind_DIR) find_package(nanobind CONFIG REQUIRED) -set(CMAKE_LIBRARY_OUTPUT_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/eudsl_tblgen) +set(CMAKE_LIBRARY_OUTPUT_DIRECTORY ${CMAKE_CURRENT_LIST_DIR}/eudsl_tblgen) set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/bin) nanobind_add_module(eudsl_tblgen_ext NB_STATIC STABLE_ABI @@ -37,6 +60,7 @@ nanobind_add_module(eudsl_tblgen_ext NB_STATIC STABLE_ABI target_link_libraries(eudsl_tblgen_ext PRIVATE LLVMTableGenCommon LLVMTableGen) target_compile_options(eudsl_tblgen_ext PUBLIC + -Wno-cast-qual $<$:-fexceptions -frtti> $<$:-fexceptions -frtti> $<$:/EHsc /GR>) @@ -44,21 +68,21 @@ target_compile_options(eudsl_tblgen_ext nanobind_add_stub( eudsl_tblgen_ext_stub MODULE eudsl_tblgen_ext - OUTPUT ${CMAKE_CURRENT_SOURCE_DIR}/eudsl_tblgen/eudsl_tblgen_ext.pyi + OUTPUT ${CMAKE_CURRENT_LIST_DIR}/eudsl_tblgen/eudsl_tblgen_ext.pyi PYTHON_PATH $ DEPENDS eudsl_tblgen_ext ) nanobind_add_stub( eudsl_tblgen_stub MODULE eudsl_tblgen - OUTPUT ${CMAKE_CURRENT_SOURCE_DIR}/eudsl_tblgen/__init__.pyi - PYTHON_PATH ${CMAKE_CURRENT_SOURCE_DIR} + OUTPUT ${CMAKE_CURRENT_LIST_DIR}/eudsl_tblgen/__init__.pyi + PYTHON_PATH ${CMAKE_CURRENT_LIST_DIR} DEPENDS eudsl_tblgen_ext ) install(TARGETS eudsl_tblgen_ext LIBRARY DESTINATION eudsl_tblgen) install( - DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/eudsl_tblgen + DIRECTORY ${CMAKE_CURRENT_LIST_DIR}/eudsl_tblgen DESTINATION ${CMAKE_INSTALL_PREFIX} PATTERN "*.so" EXCLUDE PATTERN "*.a" EXCLUDE diff --git a/projects/eudsl-tblgen/requirements.txt b/projects/eudsl-tblgen/requirements.txt deleted file mode 100644 index 1f942f7e..00000000 --- a/projects/eudsl-tblgen/requirements.txt +++ /dev/null @@ -1,3 +0,0 @@ -click==8.1.7 -pytest==8.3.3 -nanobind==2.2.0 diff --git a/projects/eudsl-py/requirements.txt b/requirements.txt similarity index 67% rename from projects/eudsl-py/requirements.txt rename to requirements.txt index 003b5eee..53635869 100644 --- a/projects/eudsl-py/requirements.txt +++ b/requirements.txt @@ -1,3 +1,2 @@ -pytest==8.3.3 nanobind==2.4.0 numpy==2.0.2