diff --git a/CMakeLists.txt b/CMakeLists.txt index b0fd7ceabaca..c22bf5510d50 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -8,7 +8,8 @@ endif() include(ExternalProject) -set(CMAKE_CXX_STANDARD 17) + +set(CMAKE_CXX_STANDARD 17) # @EUGO_CHANGE: # TODO+ Override their C++ standard to C++23 when migrating to new version set(CMAKE_INCLUDE_CURRENT_DIR ON) @@ -70,10 +71,6 @@ if(NOT CMAKE_BUILD_TYPE) set(CMAKE_BUILD_TYPE "Release") endif() -if(NOT WIN32) - find_library(TERMINFO_LIBRARY tinfo) -endif() - if(TRITON_BUILD_UT) # This is an aggregate target for all unit tests. add_custom_target(TritonUnitTests) @@ -84,7 +81,7 @@ endif() # Compiler flags include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include) if(NOT MSVC) - set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -D__STDC_FORMAT_MACROS -fPIC -std=gnu++17") + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -D__STDC_FORMAT_MACROS -fPIC -std=gnu++23") # @EUGO_CHANGE: Override their C++ standard to C++23 else() set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -D__STDC_FORMAT_MACROS") endif() @@ -142,7 +139,7 @@ endfunction() # Disable warnings that show up in external code (gtest;pybind11) if(NOT MSVC) - set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Werror -Wno-covered-switch-default -fvisibility=hidden") + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wno-covered-switch-default -fvisibility=hidden") else() set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /wd4244 /wd4624 /wd4715 /wd4530") endif() @@ -162,8 +159,8 @@ add_subdirectory(lib) # TODO: Figure out which target is sufficient to fix errors; triton is # apparently not enough. Currently set linking libstdc++fs for all targets # to support some old version GCC compilers like 8.3.0. -if (NOT WIN32 AND NOT APPLE AND NOT BSD) - link_libraries(stdc++fs) +if (NOT WIN32 AND NOT APPLE) + # link_libraries(stdc++fs) # @EUGO_CHANGE - we don't have GCC installed endif() @@ -214,7 +211,7 @@ if(TRITON_BUILD_PYTHON_MODULE) ${triton_plugins} # mlir - MLIRAMDGPUDialect + # MLIRAMDGPUDialect # @EUGO_CHANGE - no AMD support MLIRNVVMDialect MLIRNVVMToLLVMIRTranslation MLIRGPUToNVVMTransforms @@ -229,19 +226,19 @@ if(TRITON_BUILD_PYTHON_MODULE) MLIRSupport MLIRTargetLLVMIRExport MLIRMathToLLVM - MLIRROCDLToLLVMIRTranslation + # MLIRROCDLToLLVMIRTranslation # @EUGO_CHANGE - no ROCM support MLIRGPUDialect MLIRSCFToControlFlow MLIRIndexToLLVM - MLIRGPUToROCDLTransforms + # MLIRGPUToROCDLTransforms # @EUGO_CHANGE - no ROCM support MLIRUBToLLVM # LLVM LLVMPasses LLVMNVPTXCodeGen # LLVMNVPTXAsmPrinter - LLVMAMDGPUCodeGen - LLVMAMDGPUAsmParser + # LLVMAMDGPUCodeGen # @EUGO_CHANGE - no AMD support + # LLVMAMDGPUAsmParser # @EUGO_CHANGE - no AMD support Python3::Module pybind11::headers @@ -349,3 +346,9 @@ if(TRITON_BUILD_UT) USES_TERMINAL ) endif() + + +# === @begin: Eugo wrapper @EUGO_CHANGE === +# `SKBUILD_PLATLIB_DIR` resolves to `site-packages` +include(eugo_wrapper.cmake) +# === @end: Eugo wrapper === diff --git a/bin/RegisterTritonDialects.h b/bin/RegisterTritonDialects.h index f808fff35950..ece0f56a2cf6 100644 --- a/bin/RegisterTritonDialects.h +++ b/bin/RegisterTritonDialects.h @@ -1,6 +1,6 @@ #pragma once -#include "amd/include/Dialect/TritonAMDGPU/IR/Dialect.h" -#include "amd/include/TritonAMDGPUTransforms/Passes.h" +// #include "amd/include/Dialect/TritonAMDGPU/IR/Dialect.h" @EUGO_CHANGE - no AMD support +// #include "amd/include/TritonAMDGPUTransforms/Passes.h" @EUGO_CHANGE - no AMD support #include "third_party/nvidia/include/Dialect/NVGPU/IR/Dialect.h" #include "third_party/nvidia/include/Dialect/NVWS/IR/Dialect.h" #include "third_party/proton/dialect/include/Dialect/Proton/IR/Dialect.h" @@ -9,9 +9,9 @@ #include "triton/Dialect/TritonNvidiaGPU/IR/Dialect.h" // Below headers will allow registration to ROCm passes -#include "TritonAMDGPUToLLVM/Passes.h" -#include "TritonAMDGPUTransforms/Passes.h" -#include "TritonAMDGPUTransforms/TritonGPUConversion.h" +// #include "TritonAMDGPUToLLVM/Passes.h" @EUGO_CHANGE - no AMD support +// #include "TritonAMDGPUTransforms/Passes.h" @EUGO_CHANGE - no AMD support +// #include "TritonAMDGPUTransforms/TritonGPUConversion.h" @EUGO_CHANGE - no AMD support #include "triton/Dialect/Triton/Transforms/Passes.h" #include "triton/Dialect/TritonGPU/Transforms/Passes.h" @@ -26,33 +26,27 @@ #include "triton/Target/LLVMIR/Passes.h" #include "mlir/Dialect/LLVMIR/NVVMDialect.h" -#include "mlir/Dialect/LLVMIR/ROCDLDialect.h" +// #include "mlir/Dialect/LLVMIR/ROCDLDialect.h" @EUGO_CHANGE - no ROCM support #include "mlir/InitAllPasses.h" -namespace mlir { -namespace test { -void registerTestAliasPass(); -void registerTestAlignmentPass(); -void registerTestAllocationPass(); -void registerTestMembarPass(); -void registerTestAMDGPUMembarPass(); -void registerTestTritonAMDGPURangeAnalysis(); -void registerTestLoopPeelingPass(); -} // namespace test -} // namespace mlir +// namespace mlir { @EUGO_CHANGE - Don't build tests +// namespace test { +// void registerTestAliasPass(); +// void registerTestAlignmentPass(); +// void registerTestAllocationPass(); +// void registerTestMembarPass(); +// } // namespace test +// } // namespace mlir inline void registerTritonDialects(mlir::DialectRegistry ®istry) { mlir::registerAllPasses(); mlir::triton::registerTritonPasses(); mlir::triton::gpu::registerTritonGPUPasses(); - mlir::triton::nvidia_gpu::registerTritonNvidiaGPUPasses(); - mlir::test::registerTestAliasPass(); - mlir::test::registerTestAlignmentPass(); - mlir::test::registerTestAllocationPass(); - mlir::test::registerTestMembarPass(); - mlir::test::registerTestLoopPeelingPass(); - mlir::test::registerTestAMDGPUMembarPass(); - mlir::test::registerTestTritonAMDGPURangeAnalysis(); + mlir::registerTritonNvidiaGPUPasses(); + // mlir::test::registerTestAliasPass(); @EUGO_CHANGE - Don't build tests + // mlir::test::registerTestAlignmentPass(); // @EUGO_CHANGE - Don't build tests + // mlir::test::registerTestAllocationPass(); // @EUGO_CHANGE - Don't build tests + // mlir::test::registerTestMembarPass(); // @EUGO_CHANGE - Don't build tests mlir::triton::registerConvertTritonToTritonGPUPass(); mlir::triton::registerRelayoutTritonGPUPass(); mlir::triton::gpu::registerAllocateSharedMemoryPass(); @@ -64,39 +58,31 @@ inline void registerTritonDialects(mlir::DialectRegistry ®istry) { mlir::registerLLVMDIScope(); // TritonAMDGPUToLLVM passes - mlir::triton::registerConvertTritonAMDGPUToLLVM(); - mlir::triton::registerConvertBuiltinFuncToLLVM(); - mlir::triton::registerOptimizeAMDLDSUsage(); + // mlir::triton::registerConvertTritonAMDGPUToLLVM(); // @EUGO_CHANGE - no AMD support + // mlir::triton::registerConvertBuiltinFuncToLLVM(); // @EUGO_CHANGE - no AMD support (this comes from ./third_party/amd/include/TritonAMDGPUToLLVM/Passes.td and is defined like `ConvertBuiltinFuncToLLVM`) + // mlir::triton::registerDecomposeUnsupportedAMDConversions(); // @EUGO_CHANGE - no AMD support + // mlir::triton::registerOptimizeAMDLDSUsage(); // @EUGO_CHANGE - no AMD support // TritonAMDGPUTransforms passes - mlir::registerTritonAMDGPUAccelerateMatmul(); - mlir::registerTritonAMDGPUOptimizeEpilogue(); - mlir::registerTritonAMDGPUHoistLayoutConversions(); - mlir::registerTritonAMDGPUReorderInstructions(); - mlir::registerTritonAMDGPUBlockPingpong(); - mlir::registerTritonAMDGPUStreamPipeline(); - mlir::registerTritonAMDGPUCanonicalizePointers(); - mlir::registerTritonAMDGPUConvertToBufferOps(); - mlir::registerTritonAMDGPUInThreadTranspose(); - mlir::registerTritonAMDGPUCoalesceAsyncCopy(); - mlir::registerTritonAMDGPUUpdateAsyncWaitCount(); - mlir::triton::registerTritonAMDGPUInsertInstructionSchedHints(); - mlir::triton::registerTritonAMDGPULowerInstructionSchedHints(); - mlir::registerTritonAMDFoldTrueCmpI(); + // mlir::registerTritonAMDGPUAccelerateMatmul(); // @EUGO_CHANGE - no AMD support + // mlir::registerTritonAMDGPUOptimizeEpilogue(); // @EUGO_CHANGE - no AMD support + // mlir::registerTritonAMDGPUReorderInstructions(); // @EUGO_CHANGE - no AMD support + // mlir::registerTritonAMDGPUBlockPingpong(); // @EUGO_CHANGE - no AMD support + // mlir::registerTritonAMDGPUStreamPipeline(); // @EUGO_CHANGE - no AMD support + // mlir::registerTritonAMDGPUCanonicalizePointers(); // @EUGO_CHANGE - no AMD support + // mlir::registerTritonAMDGPUConvertToBufferOps(); // @EUGO_CHANGE - no AMD support + // mlir::triton::registerTritonAMDGPUInsertInstructionSchedHints(); // @EUGO_CHANGE - no AMD support + // mlir::triton::registerTritonAMDGPULowerInstructionSchedHints(); // @EUGO_CHANGE - no AMD support - // NVWS passes - mlir::triton::registerNVWSTransformsPasses(); - - // NVGPU transform passes - mlir::registerNVHopperTransformsPasses(); - - registry.insert< - mlir::triton::TritonDialect, mlir::cf::ControlFlowDialect, - mlir::triton::nvidia_gpu::TritonNvidiaGPUDialect, - mlir::triton::gpu::TritonGPUDialect, mlir::math::MathDialect, - mlir::arith::ArithDialect, mlir::scf::SCFDialect, mlir::gpu::GPUDialect, - mlir::LLVM::LLVMDialect, mlir::NVVM::NVVMDialect, - mlir::triton::nvgpu::NVGPUDialect, mlir::triton::nvws::NVWSDialect, - mlir::triton::amdgpu::TritonAMDGPUDialect, - mlir::triton::proton::ProtonDialect, mlir::ROCDL::ROCDLDialect>(); + // TODO: register Triton & TritonGPU passes + registry + .insert(); + // mlir::triton::amdgpu::TritonAMDGPUDialect, // @EUGO_CHANGE - no AMD support + // mlir::ROCDL::ROCDLDialect>(); // @EUGO_CHANGE - no ROCM support } diff --git a/eugo_wrapper.cmake b/eugo_wrapper.cmake new file mode 100644 index 000000000000..5cc5e49661ba --- /dev/null +++ b/eugo_wrapper.cmake @@ -0,0 +1,57 @@ +# @EUGO_CHANGE +install( + TARGETS triton proton LIBRARY + DESTINATION ${SKBUILD_PLATLIB_DIR}/triton/_C +) + +install( + TARGETS GPUInstrumentationTestLib LIBRARY + DESTINATION ${SKBUILD_PLATLIB_DIR}/triton/instrumentation +) + +# @HELP - after updating to new version, we will unlikely have to manually add files missing in our older commit due to .gitignore mess introduced by triton team, as these paths are covered by this function call. +# IMPORTANT: we need to exclude them twice. First when copying python/triton (because they are nested here), and second in `pyproject.toml` when they are copied together with library targets. +install( + DIRECTORY "python/triton/" + DESTINATION ${SKBUILD_PLATLIB_DIR}/triton + FILES_MATCHING + PATTERN "*.py" + PATTERN "_C/include" EXCLUDE + PATTERN "_C/include/*" EXCLUDE +) + +install( + DIRECTORY "third_party/nvidia/backend/" + DESTINATION ${SKBUILD_PLATLIB_DIR}/triton/backends/nvidia + FILES_MATCHING + PATTERN "*.py" + PATTERN "*.c" +) + +install( + FILES "third_party/nvidia/backend/lib/libdevice.10.bc" + DESTINATION ${SKBUILD_PLATLIB_DIR}/triton/backends/nvidia/lib +) + +install( + DIRECTORY "third_party/nvidia/language/cuda/" + DESTINATION ${SKBUILD_PLATLIB_DIR}/triton/language/extra/cuda + FILES_MATCHING + PATTERN "*.py" +) + +install( + DIRECTORY "third_party/proton/proton/" + DESTINATION ${SKBUILD_PLATLIB_DIR}/triton/profiler + FILES_MATCHING + PATTERN "*.py" + # PATTERN "context.py" EXCLUDE # they exclude in their latest version, but it's probably incorrect +) + +install( + DIRECTORY "third_party/nvidia/tools/cuda/" + DESTINATION ${SKBUILD_PLATLIB_DIR}/triton/tools + FILES_MATCHING + PATTERN "*.c" + PATTERN "*.h" +) \ No newline at end of file diff --git a/lib/Conversion/TritonGPUToLLVM/CMakeLists.txt b/lib/Conversion/TritonGPUToLLVM/CMakeLists.txt index 2e2412025f79..493bf6fe589f 100644 --- a/lib/Conversion/TritonGPUToLLVM/CMakeLists.txt +++ b/lib/Conversion/TritonGPUToLLVM/CMakeLists.txt @@ -29,7 +29,7 @@ add_triton_library(TritonGPUToLLVM MLIRPass MLIRGPUDialect MLIRGPUToNVVMTransforms - MLIRGPUToROCDLTransforms + # MLIRGPUToROCDLTransforms // @EUGO_CHANGE - no ROCM support MLIRGPUTransforms TritonAnalysis TritonIR diff --git a/lib/Target/LLVMIR/CMakeLists.txt b/lib/Target/LLVMIR/CMakeLists.txt index f2f9adf8f493..6dfcd5ab22f4 100644 --- a/lib/Target/LLVMIR/CMakeLists.txt +++ b/lib/Target/LLVMIR/CMakeLists.txt @@ -15,7 +15,7 @@ add_triton_library(TritonLLVMIR MLIRLLVMDialect MLIRLLVMToLLVMIRTranslation MLIRNVVMToLLVMIRTranslation - MLIRROCDLToLLVMIRTranslation + # MLIRROCDLToLLVMIRTranslation // @EUGO_CHANGE - no ROCM support MLIRSCFToControlFlow MLIRSupport MLIRTargetLLVMIRExport diff --git a/pyproject.toml b/pyproject.toml index 414f4420cba1..c749c0399bf6 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -1,6 +1,12 @@ +# @EUGO_CHANGE [build-system] -requires = ["setuptools>=40.8.0", "cmake>=3.20,<4.0", "ninja>=1.11.1", "pybind11>=2.13.1"] -build-backend = "setuptools.build_meta" + +requires = [ + "scikit_build_core", + "pybind11" +] +build-backend = "scikit_build_core.build" + [tool.mypy] mypy_path = "$MYPY_CONFIG_FILE_DIR/python" @@ -16,6 +22,26 @@ files = [ exclude = ["/build/"] follow_imports = "silent" + +# @EUGO_CHANGE +[project] +name = "triton" +version = "3.4.0" # @TODO+: Keep in sync with upstream version + +# @EUGO_CHANGE +[tool.scikit-build] +sdist.exclude = [] +install.strip = false # Default value is `true`, but we perform stripping ourselves + + +# @EUGO_CHANGE +# IMPORTANT: we need to exclude them twice. First in `CMakeLists.txt` when copying python/triton (because they are nested there), and second here in `pyproject.toml` when they are copied together with library targets. +[tool.scikit-build.wheel] +exclude = [ + "triton/_C/include/**" +] + + [tool.yapf] based_on_style = "pep8" column_limit = 120 @@ -30,8 +56,20 @@ aggressive = 1 ignore = "E501,E701,E731,W690,W503" max_line_length = 88 +[tool.yapfignore] +ignore_patterns = [ + # This exclusion is also specified in .pre-commit-config.yaml. + # - We put it here because if you run yapf directly, we want it to skip the + # file. + # - We also put it in .pre-commit-config because yapf raises an error if + # pre-commit runs it but all of the files it might touch are ignored! + "python/test/unit/language/test_line_info.py" +] + + [tool.ruff] line-length = 120 + [tool.ruff.lint] ignore = ["E501", "E701", "E731", "E741"] diff --git a/python/pyproject.toml b/python/pyproject.toml new file mode 100644 index 000000000000..d96af50a543e --- /dev/null +++ b/python/pyproject.toml @@ -0,0 +1,15 @@ + +[build-system] +requires = ["setuptools>=40.8.0", "wheel", "cmake>=3.18", "ninja>=1.11.1", "pybind11>=2.13.1"] + +# We're incrementally switching from autopep8 to ruff. +[tool.autopep8] +aggressive = 1 +ignore = "E501,E701,E731,W690,W503" +max_line_length = 88 + +[tool.ruff] +line-length = 120 + +[tool.ruff.lint] +ignore = ["E501", "E701", "E731", "E741"] diff --git a/third_party/amd/CMakeLists.txt b/third_party/amd/CMakeLists.txt index b030dbbd1a65..dc96e300c26b 100644 --- a/third_party/amd/CMakeLists.txt +++ b/third_party/amd/CMakeLists.txt @@ -1,12 +1,12 @@ -include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include) -include_directories(${CMAKE_CURRENT_BINARY_DIR}/include) -add_subdirectory(include) -add_subdirectory(lib) -if(TRITON_BUILD_PYTHON_MODULE) - add_triton_plugin(TritonAMD ${CMAKE_CURRENT_SOURCE_DIR}/python/triton_amd.cc LINK_LIBS TritonAMDGPUToLLVM TritonAMDGPUTransforms TritonAMDGPUDialectToLLVM) - target_link_libraries(TritonAMD PRIVATE Python3::Module pybind11::headers) -endif() -if(TRITON_BUILD_UT) - add_subdirectory(unittest) -endif() -add_subdirectory(test) +# @EUGO_CHANGE - no AMD support +# include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include) +# include_directories(${CMAKE_CURRENT_BINARY_DIR}/include) +# add_subdirectory(include) +# add_subdirectory(lib) +# if(TRITON_BUILD_PYTHON_MODULE) +# add_triton_plugin(TritonAMD ${CMAKE_CURRENT_SOURCE_DIR}/python/triton_amd.cc LINK_LIBS TritonAMDGPUToLLVM TritonAMDGPUTransforms TritonAMDGPUDialectToLLVM) +# target_link_libraries(TritonAMD PRIVATE Python3::Module pybind11::headers) +# endif() +# if(TRITON_BUILD_UT) +# add_subdirectory(unittest) +# endif() diff --git a/third_party/amd/include/CMakeLists.txt b/third_party/amd/include/CMakeLists.txt index 08707d601b23..dfe50d0c7ef0 100644 --- a/third_party/amd/include/CMakeLists.txt +++ b/third_party/amd/include/CMakeLists.txt @@ -1,3 +1,4 @@ -add_subdirectory(Dialect) -add_subdirectory(TritonAMDGPUToLLVM) -add_subdirectory(TritonAMDGPUTransforms) +# @EUGO_CHANGE - no AMD support +# add_subdirectory(Dialect) +# add_subdirectory(TritonAMDGPUToLLVM) +# add_subdirectory(TritonAMDGPUTransforms) diff --git a/third_party/amd/include/Dialect/CMakeLists.txt b/third_party/amd/include/Dialect/CMakeLists.txt index 4f9163bdf016..f922f9c2ae52 100644 --- a/third_party/amd/include/Dialect/CMakeLists.txt +++ b/third_party/amd/include/Dialect/CMakeLists.txt @@ -1 +1,2 @@ -add_subdirectory(TritonAMDGPU) +# @EUGO_CHANGE - no AMD support +# add_subdirectory(TritonAMDGPU) diff --git a/third_party/amd/include/Dialect/TritonAMDGPU/CMakeLists.txt b/third_party/amd/include/Dialect/TritonAMDGPU/CMakeLists.txt index f33061b2d87c..d53a4b14b2bc 100644 --- a/third_party/amd/include/Dialect/TritonAMDGPU/CMakeLists.txt +++ b/third_party/amd/include/Dialect/TritonAMDGPU/CMakeLists.txt @@ -1 +1,2 @@ -add_subdirectory(IR) +# @EUGO_CHANGE - no AMD support +# add_subdirectory(IR) diff --git a/third_party/amd/include/Dialect/TritonAMDGPU/IR/CMakeLists.txt b/third_party/amd/include/Dialect/TritonAMDGPU/IR/CMakeLists.txt index 094ecfc7d4e9..b15d92b2a5a1 100644 --- a/third_party/amd/include/Dialect/TritonAMDGPU/IR/CMakeLists.txt +++ b/third_party/amd/include/Dialect/TritonAMDGPU/IR/CMakeLists.txt @@ -1,4 +1,5 @@ -set(MLIR_BINARY_DIR ${CMAKE_BINARY_DIR}) +# @EUGO_CHANGE - no AMD support +# set(MLIR_BINARY_DIR ${CMAKE_BINARY_DIR}) set(LLVM_TARGET_DEFINITIONS TritonAMDGPUOps.td) mlir_tablegen(Dialect.h.inc -gen-dialect-decls -dialect=amdgpu) @@ -10,9 +11,9 @@ add_mlir_doc(TritonAMDGPUDialect TritonAMDGPUDialect dialects/ -gen-dialect-doc) add_mlir_doc(TritonAMDGPUOps TritonAMDGPUOps dialects/ -gen-op-doc) add_public_tablegen_target(TritonAMDGPUTableGen) -set(LLVM_TARGET_DEFINITIONS TritonAMDGPUAttrDefs.td) -mlir_tablegen(TritonAMDGPUEnums.h.inc -gen-enum-decls) -mlir_tablegen(TritonAMDGPUEnums.cpp.inc -gen-enum-defs) -mlir_tablegen(TritonAMDGPUAttrDefs.h.inc -gen-attrdef-decls) -mlir_tablegen(TritonAMDGPUAttrDefs.cpp.inc -gen-attrdef-defs) -add_public_tablegen_target(TritonAMDGPUAttrDefsIncGen) +# set(LLVM_TARGET_DEFINITIONS TritonAMDGPUAttrDefs.td) +# mlir_tablegen(TritonAMDGPUEnums.h.inc -gen-enum-decls) +# mlir_tablegen(TritonAMDGPUEnums.cpp.inc -gen-enum-defs) +# mlir_tablegen(TritonAMDGPUAttrDefs.h.inc -gen-attrdef-decls) +# mlir_tablegen(TritonAMDGPUAttrDefs.cpp.inc -gen-attrdef-defs) +# add_public_tablegen_target(TritonAMDGPUAttrDefsIncGen) diff --git a/third_party/amd/include/Dialect/TritonAMDGPU/IR/Dialect.h b/third_party/amd/include/Dialect/TritonAMDGPU/IR/Dialect.h index 9d91da924a13..427aa4c5ba20 100644 --- a/third_party/amd/include/Dialect/TritonAMDGPU/IR/Dialect.h +++ b/third_party/amd/include/Dialect/TritonAMDGPU/IR/Dialect.h @@ -34,14 +34,14 @@ #include "triton/Dialect/TritonGPU/IR/Dialect.h" // clang-format off -#include "amd/include/Dialect/TritonAMDGPU/IR/Dialect.h.inc" -#include "amd/include/Dialect/TritonAMDGPU/IR/TritonAMDGPUEnums.h.inc" +// #include "amd/include/Dialect/TritonAMDGPU/IR/Dialect.h.inc" // @EUGO_CHANGE - no AMD support +// #include "amd/include/Dialect/TritonAMDGPU/IR/TritonAMDGPUEnums.h.inc" // @EUGO_CHANGE - no AMD support // clang-format on #define GET_ATTRDEF_CLASSES -#include "amd/include/Dialect/TritonAMDGPU/IR/TritonAMDGPUAttrDefs.h.inc" +// #include "amd/include/Dialect/TritonAMDGPU/IR/TritonAMDGPUAttrDefs.h.inc" // @EUGO_CHANGE - no AMD support #define GET_OP_CLASSES -#include "amd/include/Dialect/TritonAMDGPU/IR/Ops.h.inc" +// #include "amd/include/Dialect/TritonAMDGPU/IR/Ops.h.inc" // @EUGO_CHANGE - no AMD support #endif // TRITON_THIRD_PARTY_AMD_INCLUDE_DIALECT_TRITONAMDGPU_IR_DIALECT_H_ diff --git a/third_party/amd/include/TritonAMDGPUTransforms/CMakeLists.txt b/third_party/amd/include/TritonAMDGPUTransforms/CMakeLists.txt index b8c9325f4394..1999abab3a96 100644 --- a/third_party/amd/include/TritonAMDGPUTransforms/CMakeLists.txt +++ b/third_party/amd/include/TritonAMDGPUTransforms/CMakeLists.txt @@ -1,3 +1,4 @@ -set(LLVM_TARGET_DEFINITIONS Passes.td) -mlir_tablegen(Passes.h.inc -gen-pass-decls -name TritonAMDGPU) -add_public_tablegen_target(TritonAMDGPUTransformsIncGen) +# @EUGO_CHANGE - no AMD support +# set(LLVM_TARGET_DEFINITIONS Passes.td) +# mlir_tablegen(Passes.h.inc -gen-pass-decls -name TritonAMDGPU) +# add_public_tablegen_target(TritonAMDGPUTransformsIncGen) diff --git a/third_party/amd/include/TritonAMDGPUTransforms/Passes.h b/third_party/amd/include/TritonAMDGPUTransforms/Passes.h index ca48e3863a42..c8be0335ee91 100644 --- a/third_party/amd/include/TritonAMDGPUTransforms/Passes.h +++ b/third_party/amd/include/TritonAMDGPUTransforms/Passes.h @@ -1,9 +1,11 @@ -#ifndef TRITON_THIRD_PARTY_AMD_INCLUDE_TRITONAMDGPUTRANSFORMS_PASSES_H_ -#define TRITON_THIRD_PARTY_AMD_INCLUDE_TRITONAMDGPUTRANSFORMS_PASSES_H_ +// @EUGO_CHANGE - no AMD support +// #ifndef TRITON_THIRD_PARTY_AMD_INCLUDE_TRITONAMDGPUTRANSFORMS_PASSES_H_ +// #define TRITON_THIRD_PARTY_AMD_INCLUDE_TRITONAMDGPUTRANSFORMS_PASSES_H_ -#include "mlir/Dialect/LLVMIR/ROCDLDialect.h" -#include "mlir/Pass/Pass.h" -#include "third_party/amd/include/Dialect/TritonAMDGPU/IR/Dialect.h" +// #include "mlir/Dialect/LLVMIR/ROCDLDialect.h" +// #include "mlir/Pass/Pass.h" +// #include "third_party/amd/include/Dialect/TritonAMDGPU/IR/Dialect.h" +// #include "triton/Dialect/TritonNvidiaGPU/IR/Dialect.h" namespace mlir { diff --git a/third_party/proton/CMakeLists.txt b/third_party/proton/CMakeLists.txt index 7dfbe35f1ffd..de2bcd929e27 100644 --- a/third_party/proton/CMakeLists.txt +++ b/third_party/proton/CMakeLists.txt @@ -63,11 +63,8 @@ if(NOT _proton_obj_libs) message(WARNING "No object libraries were defined in 'PROTON_LIBS'!") endif() -set(_proton_obj_sources "") -foreach(_lib IN LISTS _proton_obj_libs) - list(APPEND _proton_obj_sources $) - message(STATUS "Collecting object files from ${_lib}") -endforeach() +# include_directories(SYSTEM ${ROCTRACER_INCLUDE_DIR}) # @EUGO_CHANGE - no ROCM support +# target_compile_definitions(proton PRIVATE __HIP_PLATFORM_AMD__) # @EUGO_CHANGE - no HIP/AMD support add_library(proton SHARED ${_proton_obj_sources}) diff --git a/third_party/proton/csrc/lib/Driver/Device.cpp b/third_party/proton/csrc/lib/Driver/Device.cpp index 1fb1f2361e90..73e157470018 100644 --- a/third_party/proton/csrc/lib/Driver/Device.cpp +++ b/third_party/proton/csrc/lib/Driver/Device.cpp @@ -1,6 +1,6 @@ #include "Driver/Device.h" #include "Driver/GPU/CudaApi.h" -#include "Driver/GPU/HipApi.h" +// #include "Driver/GPU/HipApi.h" // @EUGO_CHANGE - no HIP support #include "Utility/Errors.h" @@ -10,18 +10,18 @@ Device getDevice(DeviceType type, uint64_t index) { if (type == DeviceType::CUDA) { return cuda::getDevice(index); } - if (type == DeviceType::HIP) { - return hip::getDevice(index); - } + // if (type == DeviceType::HIP) { // @EUGO_CHANGE - no HIP support + // return hip::getDevice(index); + // } throw std::runtime_error("DeviceType not supported"); } const std::string getDeviceTypeString(DeviceType type) { if (type == DeviceType::CUDA) { return DeviceTraits::name; - } else if (type == DeviceType::HIP) { - return DeviceTraits::name; - } + } // else if (type == DeviceType::HIP) { // @EUGO_CHANGE - no HIP support + // return DeviceTraits::name; // @EUGO_CHANGE - no HIP support + // } // @EUGO_CHANGE - no HIP support throw std::runtime_error("DeviceType not supported"); } diff --git a/third_party/proton/csrc/lib/Driver/GPU/HipApi.cpp b/third_party/proton/csrc/lib/Driver/GPU/HipApi.cpp index 9e8ef8d225fc..adf02778fea3 100644 --- a/third_party/proton/csrc/lib/Driver/GPU/HipApi.cpp +++ b/third_party/proton/csrc/lib/Driver/GPU/HipApi.cpp @@ -1,7 +1,8 @@ -#include "Driver/GPU/HipApi.h" -#include "Driver/Dispatch.h" -#include "hip/hip_runtime_api.h" -#include +// @EUGO_CHANGE - no HIP support +// #include "Driver/GPU/HipApi.h" +// #include "Driver/Dispatch.h" +// #include "hip/hip_runtime_api.h" +// #include namespace proton { diff --git a/third_party/proton/csrc/lib/Driver/GPU/HsaApi.cpp b/third_party/proton/csrc/lib/Driver/GPU/HsaApi.cpp index 7c607b4b99ec..1956a7b8c09e 100644 --- a/third_party/proton/csrc/lib/Driver/GPU/HsaApi.cpp +++ b/third_party/proton/csrc/lib/Driver/GPU/HsaApi.cpp @@ -1,5 +1,6 @@ -#include "Driver/GPU/HsaApi.h" -#include "Driver/Dispatch.h" +// @EUGO_CHANGE - no HIP support +// #include "Driver/GPU/HsaApi.h" +// #include "Driver/Dispatch.h" namespace proton { diff --git a/third_party/proton/csrc/lib/Driver/GPU/RoctracerApi.cpp b/third_party/proton/csrc/lib/Driver/GPU/RoctracerApi.cpp index a6dcdcf346ea..cc09a768f31d 100644 --- a/third_party/proton/csrc/lib/Driver/GPU/RoctracerApi.cpp +++ b/third_party/proton/csrc/lib/Driver/GPU/RoctracerApi.cpp @@ -1,5 +1,6 @@ -#include "Driver/GPU/RoctracerApi.h" -#include "Driver/Dispatch.h" +// @EUGO_CHANGE - no ROCM support +// #include "Driver/GPU/RoctracerApi.h" +// #include "Driver/Dispatch.h" namespace proton { diff --git a/third_party/proton/csrc/lib/Profiler/RocTracer/RoctracerProfiler.cpp b/third_party/proton/csrc/lib/Profiler/RocTracer/RoctracerProfiler.cpp index 2439e887ff3b..82a96e36f149 100644 --- a/third_party/proton/csrc/lib/Profiler/RocTracer/RoctracerProfiler.cpp +++ b/third_party/proton/csrc/lib/Profiler/RocTracer/RoctracerProfiler.cpp @@ -1,13 +1,14 @@ -#include "Profiler/Roctracer/RoctracerProfiler.h" -#include "Context/Context.h" -#include "Data/Metric.h" -#include "Driver/GPU/HipApi.h" -#include "Driver/GPU/HsaApi.h" -#include "Driver/GPU/RoctracerApi.h" - -#include "hip/amd_detail/hip_runtime_prof.h" -#include "roctracer/roctracer_ext.h" -#include "roctracer/roctracer_hip.h" +// @EUGO_CHANGE - no ROCM/HIP/HSA support +// #include "Profiler/Roctracer/RoctracerProfiler.h" +// #include "Context/Context.h" +// #include "Data/Metric.h" +// #include "Driver/GPU/HipApi.h" +// #include "Driver/GPU/HsaApi.h" +// #include "Driver/GPU/RoctracerApi.h" + +// #include "hip/amd_detail/hip_runtime_prof.h" +// #include "roctracer/roctracer_ext.h" +// #include "roctracer/roctracer_hip.h" #include #include diff --git a/third_party/proton/csrc/lib/Session/Session.cpp b/third_party/proton/csrc/lib/Session/Session.cpp index 8a476355f033..c2fc6fefe910 100644 --- a/third_party/proton/csrc/lib/Session/Session.cpp +++ b/third_party/proton/csrc/lib/Session/Session.cpp @@ -3,7 +3,7 @@ #include "Context/Shadow.h" #include "Data/TreeData.h" #include "Profiler/Cupti/CuptiProfiler.h" -#include "Profiler/Roctracer/RoctracerProfiler.h" +// #include "Profiler/Roctracer/RoctracerProfiler.h" // @EUGO_CHANGE - no ROCM support #include "Utility/String.h" namespace proton { @@ -16,10 +16,10 @@ Profiler *getProfiler(const std::string &name, const std::string &path) { if (proton::toLower(name) == "cupti_pcsampling") { return &CuptiProfiler::instance().setLibPath(path).enablePCSampling(); } - if (proton::toLower(name) == "roctracer") { - return &RoctracerProfiler::instance(); - } - throw std::runtime_error("Unknown profiler: " + name); + // if (proton::toLower(profilerName) == "roctracer") { // @EUGO_CHANGE - no ROCM support + // return &RoctracerProfiler::instance(); // @EUGO_CHANGE - no ROCM support + // } // @EUGO_CHANGE - no ROCM support + throw std::runtime_error("Unknown profiler: " + profilerName); } std::unique_ptr makeData(const std::string &dataName,