[CI] Add manylinux auditwheel repair

[executor] Introduce a pass to expand unsupproted Math operations

Adds path `executor-expand-math-ops` to expand unsupported Math
operations into compositions of supported operations. This pass
utilizes upstream Math dialect rewrite pattern sets (e.g.
`math-expand-ops`, `math-polynomial-approx` etc).

GitOrigin-RevId: 5d1736f1a5a398ee37a003461ee2d32cd0a5c06e

[mlir-tensorrt] Add Stablehlo patch for various Stablehlo upstream pass issues

Add a Stablehlo patch for the following issues:

1. No Atan2 scalar float support in StablehloToScalarOp template
2. Fix crash when convert `stablehlo.bitcast_convert` with complex
   type to arith::BitcastOp
3. Fix StablehloRefineShapes can accidentally erase functions with
   side effects
4. Fix StablehloAggressiveFolder can crash when folding
   `stablehlo.compare` if the result type has static dimensions
   erased.

[compiler][emitc] Add support for embedding and emitting runtime files

This change enables the compiler to emit the required Standalone runtime
sources and headers as artifacts when translating to EmitC, eliminating the
need for users to manually locate and copy runtime files from the source tree.

The StandaloneCPP runtime files (*.cpp, *.h) are now embedded directly into
the compiler binary at build time using a new CMake script
(GenerateEmbeddedStandaloneCPP.cmake) that generates a translation unit
containing the file contents as raw string literals.

A new pass `EmitCppSupportFilesPass` analyzes the EmitC module to determine
which runtime components are required (Core, CUDA, TensorRT) and emits them
as `executor.file_artifact` operations. The pass can also optionally generate
an example CMakeLists.txt and a test driver source file.

New compiler options control emission:
- `--emitc-emit-support-files`: Emit all support files (runtime, CMake, driver)
- `--emitc-emit-runtime-files`: Emit only the required runtime subset
- `--emitc-emit-cmake-file`: Emit an example CMake file
- `--emitc-emit-test-driver`: Emit a C++ test driver source file

The `-o` output path option is now scoped to global CLI parsing only
to avoid conflicts when parsing option strings programmatically.

GitOrigin-RevId: 330f1f17e78584131b1b4a482e7600b39d3dfb27

[compiler] NFC: Add missing `memref-to-cuda` test cases

[integrations/PJRT] Fix CMake configuration for PJRT library

- Previous change to PJRT CMake config broke the build
  when BUILD_SHARED_LIBS is set to ON.
- The fix is simply to undo the change to the location
  of the CMAKE_CXX_VISIBILITY_PRESET setting.
- In addition, the visiblility fix using the linker option
  is improved using CMake `LINKER:` prefix.
- Finally, the linker option fixed the issue that was causing
  us to build an additional PJRT static library for unit tests,
  so we can eliminate the extra library and just build the one
  monolithic shared library for testing and deployment. This
  additionally seems to act as a check against LLVM/MLIR
  symbol visibility.

GitOrigin-RevId: 24c3090fca08602319466e2eefeb9f6fb0f68677

NFC: Consolidate CUDA integration tests and simplify test commands

GitOrigin-RevId: 94c9f43ca8c51c6f02edcc8f26725268a84210a1

[mlir-tensorrt] Integrate internal changes

--- [compiler] Add cuda.get_program_device op

Introduce `cuda.get_program_device` as a pure/speculatable way to map
a program logical device id (i32) to a CUDA device ordinal (i32).

GitOrigin-RevId: 00512cc5a9e9c61023e1d9de734b2383da369bcf

--- [compiler] Refactor device management and stream creation utilities

This commit introduces a new device management model to support
multi-device SPMD and MPMD programs and refactors stream creation to
use reusable utility functions. The primary motivation is to enable
more flexible device assignment where programs can be assigned to
specific CUDA ordinals via logical device IDs, laying the groundwork
for better multi-device support.

GitOrigin-RevId: 447b72743e64f394671f866fcdfdb0d6f0f3d579

---[compiler|executor] Refactor plugin call stream handling

This change refactors how CUDA streams are handled for plugin calls in
the executor dialect. Previously, when no stream was provided to a
CallPluginOp, the lowering would create and use a global CUDA stream
(stream0). This approach had several issues:

1. It tightly coupled the executor dialect to CUDA-specific stream
   creation
2. It required maintaining global stream state across compilation
3. It made the stream handling implicit and harder to reason about

The new approach uses null streams (nullptr) when no explicit stream
is provided. This is the standard CUDA convention where a null stream
represents the default stream. The changes include:

- Modified `executor.call_plugin` op to accept any type for the stream
  operand (not just `!executor.ptr<host>`), allowing frontend dialects
  to pass their own stream representations (e.g. `!cuda.stream`)
- Updated the assembly format to print the stream type for clarity
- Removed `getGlobalCudaStream` helper method from
  ConvertToExecutorPattern
- Changed CallPluginConversionPattern to create a null pointer
  (inttoptr 0) when no stream is provided instead of creating a global
  stream
- Updated StablehloToPlan conversion to use
  `cuda::getOrCreateDefaultStream0` to explicitly create CUDA streams
  when converting TVM FFI custom calls
- Added CUDADialect dependency to StablehloToPlan pass and CMakeLists

This makes stream handling more explicit and flexible, allowing
different frontend dialects to manage their own stream creation while
falling back to null streams (CUDA default stream) when appropriate.

GitOrigin-RevId: 764238bc58308d5d284f8e32da91c7e5f90fdf0c

Author: lanluo-nvidia

.github/workflows/mlir-tensorrt-build-test.yml

@@ -4,7 +4,7 @@ on:
   workflow_call:
     inputs:
       channel:
-        description: 'Channel, valid values are "nightly", "test", or "release"'
+        description: 'Channel, valid values are "nightly", "test", "release" or "pypi-release"'
         default: "test"
         type: string
       build-matrix:
@@ -26,7 +26,7 @@ jobs:
     env:
       # eg. TENSORRT_VERSION: 10.12 or 10.13
       MLIR_TRT_DOWNLOAD_TENSORRT_VERSION: ${{ matrix.trt }}
-      # eg. CHANNEL: nightly, test or release
+      # eg. CHANNEL: nightly, test, release or pypi-release
       CHANNEL: ${{ inputs.channel }}
       ARCH: ${{ matrix.arch }}
       CMAKE_PRESET: ${{ matrix.cmake_preset }}

.github/workflows/mlir-tensorrt-pypi.yml

@@ -0,0 +1,194 @@
+name: MLIR-TensorRT PyPI Release CI
+
+on:
+  workflow_dispatch:
+    inputs:
+      confirm_publish_to_pypi:
+        description: 'I confirm that I have verified version to be published to pypi.org'
+        required: true
+        type: boolean
+        default: false
+
+defaults:
+  run:
+    shell: bash
+
+jobs:
+  generate-matrix:
+    name: Generate Build Matrix (pypi-release)
+    runs-on: ubuntu-latest
+    outputs:
+      matrix: ${{ steps.generate-matrix.outputs.matrix }}
+    steps:
+      - uses: actions/checkout@v6
+        with:
+          fetch-depth: 5
+      - name: Generate Build Matrix
+        id: generate-matrix
+        run: |
+          set -euo pipefail
+          set -x
+          MATRIX_BLOB="$(python3 ./.github/workflows/mlir-tensorrt/generate-matrix.py --channel pypi-release)"
+          echo "${MATRIX_BLOB}"
+          echo "matrix=${MATRIX_BLOB}" >> "${GITHUB_OUTPUT}"
+
+  pypi-release-wheels-build:
+    name: ${{ matrix.arch}} - Build PyPI Release Wheels
+    needs:
+      - generate-matrix
+    permissions:
+      id-token: write
+      packages: write
+      contents: read
+    strategy:
+      matrix: ${{ fromJSON(needs.generate-matrix.outputs.matrix) }}
+    runs-on: ${{ matrix.github_runner }}
+    env:
+      MLIR_TRT_DOWNLOAD_TENSORRT_VERSION: ${{ matrix.trt }}
+      CMAKE_PRESET: ${{ matrix.cmake_preset }}
+      CCACHE_RESTORE_KEY: mlir-tensorrt-ccache-v1-${{ matrix.arch }}-${{ matrix.cmake_preset }}
+      CPM_RESTORE_KEY: mlir-tensorrt-cpm-v1
+    timeout-minutes: 120
+    container:
+      # pypi audit wheel repair requires rockylinux8
+      image: ${{ matrix.docker_image }}
+      options: >-
+        --gpus all
+        --shm-size=1g
+    steps:
+      - name: Checkout TensorRT-Incubator
+        uses: actions/checkout@v6
+        with:
+          fetch-depth: 5
+
+      - name: Create Cache Folders
+        run: |
+          set -euo pipefail
+          set -x
+          export CPM_SOURCE_CACHE=${GITHUB_WORKSPACE}/mlir-tensorrt/.cache.cpm
+          export CCACHE_DIR=${GITHUB_WORKSPACE}/mlir-tensorrt/ccache
+
+          echo "CPM_SOURCE_CACHE=$CPM_SOURCE_CACHE" >> "$GITHUB_ENV"
+          echo "CCACHE_DIR=$CCACHE_DIR" >> "$GITHUB_ENV"
+
+          mkdir -p ${CCACHE_DIR}
+          mkdir -p ${CPM_SOURCE_CACHE}
+
+      - name: Compute CCache Key
+        id: ccache-key
+        run: |
+          hash=$( (find mlir-tensorrt/compiler \
+                        mlir-tensorrt/common \
+                        mlir-tensorrt/kernel \
+                        mlir-tensorrt/tensorrt \
+                        mlir-tensorrt/integrations \
+                        mlir-tensorrt/executor \
+                   -type f \( -name '*.cpp' -o -name '*.h' \) \
+                   -exec sha256sum {} \; ; \
+                  sha256sum mlir-tensorrt/DependencyProvider.cmake \
+                            mlir-tensorrt/CMakeLists.txt) \
+                 | sort | sha256sum | cut -d' ' -f1)
+          echo "key=${{ env.CCACHE_RESTORE_KEY }}-${hash}" >> $GITHUB_OUTPUT
+
+      - name: Compute CPM Key
+        id: cpm-key
+        run: |
+          hash=$(sha256sum mlir-tensorrt/DependencyProvider.cmake | cut -d' ' -f1)
+          echo "key=${{ env.CPM_RESTORE_KEY }}-${hash}" >> $GITHUB_OUTPUT
+
+      - name: Restore CCache
+        id: restore-ccache
+        uses: actions/cache/restore@v4
+        with:
+          key: ${{ steps.ccache-key.outputs.key }}
+          restore-keys: |
+            ${{ env.CCACHE_RESTORE_KEY }}
+          path: |
+            ${{ env.CCACHE_DIR }}
+
+      - name: Restore CPM cache
+        id: restore-cpm
+        uses: actions/cache/restore@v4
+        with:
+          key: ${{ steps.cpm-key.outputs.key }}
+          enableCrossOsArchive: true
+          restore-keys: |
+            ${{ env.CPM_RESTORE_KEY }}
+          path: |
+            mlir-tensorrt/.cache.cpm/*
+            !mlir-tensorrt/.cache.cpm/tensorrt
+            !mlir-tensorrt/.cache.cpm/tensorrt/**
+
+      - name: Build Wheels With CUDA:${{ matrix.cuda }} + TensorRT:${{ matrix.trt }}
+        env:
+          MLIR_TRT_DOWNLOAD_TENSORRT_VERSION: ${{ matrix.trt }}
+          ARCH: ${{ matrix.arch }}
+          CMAKE_PRESET: distribution-wheels
+        run: |
+          set -euo pipefail
+          set -x
+          cd mlir-tensorrt
+          # Build only pjrt wheels for PyPI upload, with auditwheel repair
+          MLIR_TRT_PYPI=1 PACKAGES="pjrt" ./build_tools/scripts/cicd-build-wheels.sh
+
+      - name: Upload Wheels
+        uses: actions/upload-artifact@v4
+        with:
+          name: release-wheels-${{ matrix.arch }}-cu${{ matrix.cuda }}-trt${{ matrix.trt }}
+          path: mlir-tensorrt/dist
+          if-no-files-found: error
+
+  test-pypi-release-wheels-publish:
+    name: Publish to TestPyPI
+    needs: [pypi-release-wheels-build]
+    runs-on: ubuntu-latest
+    environment:
+      name: testpypi
+      url: https://test.pypi.org/project/mlir-tensorrt-jax/
+    permissions:
+      id-token: write
+      packages: write
+      contents: read
+    steps:
+      - name: Download built wheels
+        uses: actions/download-artifact@v4
+        with:
+          pattern: release-wheels-*
+          merge-multiple: true
+          path: dist
+
+      - name: Publish to TestPyPI
+        uses: pypa/gh-action-pypi-publish@release/v1
+        with:
+          skip-existing: true
+          verbose: true
+          repository-url: https://test.pypi.org/legacy/
+
+  pypi-release-wheels-publish:
+    name: Publish to PyPI
+    if: ${{ inputs.confirm_publish_to_pypi }}
+    needs: [pypi-release-wheels-build]
+    runs-on: ubuntu-latest
+    permissions:
+      id-token: write
+      packages: write
+      contents: read
+    steps:
+      - name: Download built wheels
+        uses: actions/download-artifact@v4
+        with:
+          pattern: release-wheels-*
+          merge-multiple: true
+          path: dist
+
+      - name: Publish to PyPI
+        uses: pypa/gh-action-pypi-publish@release/v1
+        with:
+          verbose: true
+          skip-existing: true
+          user: __token__
+          password: ${{ secrets.PYPI_API_TOKEN }}
+
+concurrency:
+  group: ${{ github.workflow }}-${{ github.event.pull_request.number || github.ref_name }}-mlir-tensorrt-pypi
+  cancel-in-progress: true

.github/workflows/mlir-tensorrt-release.yml

@@ -23,8 +23,8 @@ jobs:
       # eg. 10.12 or 10.13
       MLIR_TRT_DOWNLOAD_TENSORRT_VERSION: ${{ matrix.trt }}
       ARCH: ${{ matrix.arch }}
-      CMAKE_PRESET: distribution-wheels
-      CCACHE_RESTORE_KEY: mlir-tensorrt-ccache-v1-${{ matrix.arch }}-distribution-wheels
+      CMAKE_PRESET: ${{ matrix.cmake_preset }}
+      CCACHE_RESTORE_KEY: mlir-tensorrt-ccache-v1-${{ matrix.arch }}-${{ matrix.cmake_preset }}
       CPM_RESTORE_KEY: mlir-tensorrt-cpm-v1
     runs-on: ${{ matrix.github_runner }}
     timeout-minutes: 120

.github/workflows/mlir-tensorrt/generate-matrix.py

@@ -31,12 +31,19 @@
             "trt": "10.13",
         },
     ],
+    "pypi-release": [
+        {
+            "cuda": "13.0",
+            "trt": "10.13",
+        },
+    ],
 }
 
 ARCH_LIST_DICT = {
     "test": ["x86_64"],
     "release": ["x86_64", "aarch64"],
     "nightly": ["x86_64", "aarch64"],
+    "pypi-release": ["x86_64", "aarch64"],
 }
 
 GH_RUNNER_DICT = {
@@ -47,9 +54,8 @@
 CMAKE_PRESET_DICT = {
     "nightly": "github-cicd",
     "test": "github-cicd",
-    # release should use the release wheel build preset
-    # TODO: add the release wheel build preset
-    "release": "github-cicd",
+    "release": "distribution-wheels",
+    "pypi-release": "distribution-wheels",
 }
 
 DOCKER_IMAGE_DICT = {
@@ -71,6 +77,9 @@
             "13.0": "ghcr.io/nvidia/tensorrt-incubator/mlir-tensorrt:cuda13.0-rockylinux9-0.1",
         },
     },
+    "pypi-release": {
+        "13.0": "ghcr.io/nvidia/tensorrt-incubator/mlir-tensorrt:cuda13.0-rockylinux8-0.1",
+    },
 }
 
 
@@ -84,9 +93,9 @@ def main(args: list[str]) -> None:
     )
 
     options = parser.parse_args(args)
-    if options.channel not in ("nightly", "test", "release"):
+    if options.channel not in ("nightly", "test", "release", "pypi-release"):
         raise Exception(
-            "--channel is invalid, please choose from nightly, test or release"
+            "--channel is invalid, please choose from nightly, test, release or pypi-release"
         )
 
     channel = options.channel

mlir-tensorrt/.devcontainer/cuda12.9-rockylinux8-prebuilt/devcontainer.json

@@ -0,0 +1,88 @@
+{
+  "name": "cuda12.9-rockylinux8-prebuilt",
+  "image": "ghcr.io/nvidia/tensorrt-incubator/mlir-tensorrt:cuda12.9-rockylinux8-0.1",
+  "remoteUser": "nvidia",
+  "updateRemoteUserUID": true,
+  "runArgs": [
+    "--name",
+    "cuda12.9-rockylinux8-prebuilt-${localEnv:USER:nvidia}-${devcontainerId}",
+    "--cap-add=SYS_PTRACE",
+    "--security-opt",
+    "seccomp=unconfined",
+    "--shm-size=1g",
+    "--ulimit",
+    "memlock=-1",
+    "--network=host"
+  ],
+  "hostRequirements": {
+    "gpu": "optional"
+  },
+  "workspaceMount": "source=${localWorkspaceFolder}/..,target=/workspaces/TensorRT-Incubator,type=bind,consistency=cached",
+  "workspaceFolder": "/workspaces/TensorRT-Incubator/mlir-tensorrt",
+  "customizations": {
+    "vscode": {
+      "extensions": [
+        "llvm-vs-code-extensions.vscode-clangd",
+        "llvm-vs-code-extensions.vscode-mlir",
+        "eamodio.gitlens",
+        "ms-python.black-formatter",
+        "ms-python.python"
+      ],
+      "settings": {
+        "[python]": {
+          "editor.defaultFormatter": "ms-python.black-formatter"
+        },
+        "mlir.pdll_compilation_databases": [
+          "build/pdll_compile_commands.yml"
+        ],
+        "mlir.server_path": "build/bin/mlir-tensorrt-lsp-server",
+        "files.exclude": {
+          "**/.git": true,
+          "**/.cache": true,
+          "**/.venv*": true
+        },
+        "files.watcherExclude": {
+          "**/.git/objects/**": true,
+          "**/.git/subtree-cache/**": true,
+          "**/.private*": true,
+          "**/.venv*/**": true,
+          "**/build/**": true
+        },
+        "search.exclude": {
+          "**/.private*": true,
+          "**/.venv*": true,
+          "**/build": true
+        },
+        "python.analysis.include": [
+          "integrations/python",
+          "integrations/python/internal"
+        ],
+        "python.analysis.typeCheckingMode": "basic",
+        "python.analysis.extraPaths": [
+          "build/python_packages/mlir_tensorrt_compiler",
+          "build/python_packages/mlir_tensorrt_runtime",
+          "build/python_packages/tools"
+        ],
+        "python.analysis.exclude": [
+          "**/build/**",
+          "**/.cache.cpm/**",
+          "**/*bazel*/**",
+          "**/build_tools/**",
+          "third_party"
+        ]
+      }
+    }
+  },
+  "features": {
+    "ghcr.io/devcontainers/features/common-utils:2": {
+      "installZsh": true,
+      "installOhMyZsh": true,
+      "configureZshAsDefaultShell": false,
+      "upgradePackages": false,
+      "username": "nvidia",
+      "userUid": "automatic",
+      "userGid": "automatic"
+    },
+    "ghcr.io/devcontainers/features/git:1": {}
+  }
+}