Merge remote-tracking branch 'origin' into amyachev/bmg-skiplist2

Author: anmyachev

.github/WINDOWS.md

@@ -56,7 +56,7 @@ If you do not have a system Python installed at this step, you can install one w
 For example:
 
 ```
-choco install python --version=3.9.13
+choco install python --version=3.10.11
 ```
 
 ### Git

.github/workflows/build-test-python.yml

@@ -58,9 +58,9 @@ jobs:
         id: matrix
         run: |
           if [[ -n "${{ inputs.runner_label }}" ]]; then
-            matrix='{"python": ["3.9", "3.10", "3.11", "3.12", "3.13"]}'
+            matrix='{"python": ["3.10", "3.11", "3.12", "3.13"]}'
           else
-            matrix='{"python": ["3.9", "3.10", "3.11", "3.12", "3.13"], "driver": ["rolling", "lts"]}'
+            matrix='{"python": ["3.10", "3.11", "3.12", "3.13"], "driver": ["rolling", "lts"]}'
           fi
           echo "matrix=$matrix" | tee -a $GITHUB_OUTPUT
 

.github/workflows/nightly-wheels.yml

@@ -32,7 +32,6 @@ jobs:
     strategy:
       matrix:
         python:
-          - "3.9"
           - "3.10"
           - "3.11"
           - "3.12"

.github/workflows/try-latest-pytorch.yml

@@ -68,7 +68,7 @@ jobs:
       - name: Matrix
         id: matrix
         run: |
-          integration_matrix='{"python": ["3.9", "3.10", "3.11", "3.12"], "driver": ["rolling", "lts"]}'
+          integration_matrix='{"python": ["3.10", "3.11", "3.12"], "driver": ["rolling", "lts"]}'
 
           echo "integration_matrix=$integration_matrix" | tee -a $GITHUB_OUTPUT
           e2e_matrix='{
@@ -97,7 +97,7 @@ jobs:
         inductor/test_max_autotune.py
         inductor/test_compile_subprocess.py
       runner_label: ${{ inputs.runner_label }}
-      python_version: "3.9"
+      python_version: "3.10"
 
   integration-tests:
     name: Integration tests

.github/workflows/wheels-pytorch.yml

@@ -21,7 +21,6 @@ jobs:
     strategy:
       matrix:
         python:
-          - "3.9"
           - "3.10"
           - "3.11"
           - "3.12"

.github/workflows/wheels-triton.yml

@@ -16,7 +16,6 @@ jobs:
     strategy:
       matrix:
         python:
-          - "3.9"
           - "3.10"
           - "3.11"
           - "3.12"

README.md

@@ -27,7 +27,7 @@ You can install the latest stable release of Triton from pip:
 pip install triton
 ```
 
-Binary wheels are available for CPython 3.9-3.13.
+Binary wheels are available for CPython 3.10-3.13.
 
 # Install from source
 

bin/RegisterTritonDialects.h

@@ -47,6 +47,7 @@
 
 #include "mlir/Dialect/LLVMIR/NVVMDialect.h"
 #include "mlir/Dialect/LLVMIR/ROCDLDialect.h"
+#include "mlir/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.h"
 #include "mlir/InitAllPasses.h"
 
 namespace mlir {
@@ -107,6 +108,8 @@ inline void registerTritonDialects(mlir::DialectRegistry &registry) {
   mlir::triton::registerConvertTritonGENToLLVM();
   mlir::triton::registerTritonGENToLLVMPasses();
   mlir::triton::registerTritonGENToSPIRVPasses();
+  mlir::LLVM::registerInlinerInterface(registry);
+  mlir::NVVM::registerInlinerInterface(registry);
 
   // TritonAMDGPUToLLVM passes
   mlir::triton::registerAllocateAMDGPUSharedMemory();

docs/getting-started/installation.rst

@@ -14,7 +14,7 @@ You can install the latest stable release of Triton from pip:
 
       pip install triton
 
-Binary wheels are available for CPython 3.9-3.13.
+Binary wheels are available for CPython 3.10-3.13.
 
 -----------
 From Source

include/triton/Dialect/TritonNvidiaGPU/IR/TritonNvidiaGPUAttrDefs.td

@@ -36,14 +36,16 @@ def TTG_TensorMemoryEncodingAttr : AttrDef<TritonNvidiaGPU_Dialect, "TensorMemor
   let attrName = "triton.gpu.tensor_memory_encoding";
   let description = [{
     An encoding to represent the different way the tensor memory is laid out.
-    `unpacked` attributes indicates whether types smaller than 32bits are unpacked (take full 32bits)
-    or are packed (N elements are stored within one 32bits row).
+    `colStride` describes the stride in elements along the column dimension,
+    that is, the stride between two elements in the same row.
+    When colStride is 1 the tensor memory is packed. When colStride > 1, the
+    tensor memory between elements is undefined.
   }];
   let parameters = (
     ins
     "unsigned":$blockM,
     "unsigned":$blockN,
-    "bool":$unpacked,
+    "unsigned":$colStride,
     DefaultValuedParameter<"unsigned", "1">:$CTASplitM,
     DefaultValuedParameter<"unsigned", "1">:$CTASplitN
   );