[DELETE ME] - leaving for reference bwl1289/feat/eugo-changes

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 ===

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 &registry) {
   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 &registry) {
   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::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::proton::ProtonDialect>();
+              // mlir::triton::amdgpu::TritonAMDGPUDialect, // @EUGO_CHANGE - no AMD support
+              // mlir::ROCDL::ROCDLDialect>(); // @EUGO_CHANGE - no ROCM support
 }

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"
+)

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

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

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"]

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"]