Drop inferred amdgpu-no-* attributes from instrumented kernel clones by zhanglx13 · Pull Request #44 · AMDResearch/omniprobe

zhanglx13 · 2026-09-26T06:08:43Z

Fixes #43.

Problem

When the LLVM pipeline has a target machine, the AMDGPU attributor runs at OptimizerLastEP before this plugin. It marks each kernel with amdgpu-no-* attributes for implicit inputs the kernel does not use, for example amdgpu-no-workgroup-id-y and amdgpu-no-implicitarg-ptr. cloneKernelWithExtraArg copies those attributes onto the instrumented clone, but the injected dh_comms calls read blockIdx.y/z, threadIdx.y/z and gridDim. The backend trusts the attributes and does not set those inputs up, so the reads are undefined (per LLVM's AMDGPU documentation for these attributes).

Triton currently hides this by dropping the target machine whenever LLVM_PASS_PLUGIN_PATH is set, so the attributor never runs. triton-lang/triton#10849 removes that special case.

Fix

Remove the amdgpu-no-* attributes from the clone after CloneFunctionInto (which copies the original's attributes again, so removing them earlier has no effect). The backend then provides all implicit inputs to instrumented kernels. Uninstrumented kernels keep their attributes.

Testing

Built libAMDGCNSubmitAddressMessages-triton.so from this branch against Triton's LLVM (b010a18d) and compiled a Triton vector-add on gfx950 with the plugin loaded, using current Triton main plus triton-lang/triton#10849 (target machine kept). Instrumented clone __amd_crk_vaddPv:

	`amdgpu-no-*` attributes	`.amdhsa_system_sgpr_workgroup_id_y`	`.amdhsa_kernarg_size`
Triton `main` (no target machine), before and after	0	1	312
Triton with #10849, before this fix	21	0	56
Triton with #10849, with this fix	0	1	312

Without the fix, the clone reserves no implicit arguments, so gridDim reads land outside the kernel-argument segment. A separate probe plugin that injects a non-inlined call reading workgroup_id.y and the hidden block count confirmed the runtime effect: 32737 instead of 2 and 0 instead of 4 with the target machine, correct values once the attributes are removed.

Not tested: the full Omniprobe runtime and the HIP (-fpass-plugin) path. The HIP device compile also runs the attributor before plugins, so the same fix should apply there.

For the Triton test I also had to relax validateAMDGPUTarget locally (not part of this PR), because Triton main now emits the new-style triple amdgpu9.50-amd-amdhsa instead of amdgcn-amd-amdhsa, and triton_install.sh's assert len(names) == 1 patch was needed as usual. Details in #43.

When the LLVM pipeline has a target machine, the AMDGPU attributor runs at OptimizerLastEP before this plugin and marks kernels with "amdgpu-no-*" attributes for implicit inputs they do not use, such as the workgroup id y or the implicit-argument pointer. cloneKernelWithExtraArg copied those onto the instrumented clone, whose injected dh_comms calls read blockIdx.y/z, threadIdx.y/z and gridDim. The backend then did not set those inputs up, so the reads were undefined. Remove the attributes from the clone after CloneFunctionInto (which copies the original's attributes again), so the backend provides all implicit inputs to instrumented kernels. Uninstrumented kernels are unchanged.

`optimize_module` drops the target machine whenever an LLVM-IR pass plugin is loaded (`LLVM_PASS_PLUGIN_PATH`). That guard was added in #4669 for instrumentation plugins that splice in externally (Clang-)compiled code. This PR creates the target machine whenever `arch` is known, whether or not a plugin is loaded. ## Why Loading a plugin currently changes how the kernel itself is optimized, not only what the plugin adds. Without a target machine, the O3 pipeline runs without the AMDGPU target passes (for example `AMDGPUAttributor`) and without target cost models. Measured on gfx950 with a vector-add kernel: | | no plugin | in-tree `GPUHello` plugin loaded | |---|---:|---:| | `amdgpu-no-*` attributes in LLIR, `main` | 21 | 0 | | `amdgpu-no-*` attributes in LLIR, this PR | 21 | 21 | `GPUHello` only prints, yet loading it removes every attribute the attributor infers for this kernel. How much that matters depends on the kernel; for a plain matmul the generated code is nearly identical, while scheduler-style plugins, such as the out-of-tree LLIR scheduler used by the gfx950 Gluon tutorials, lose performance without the target cost models. With this PR the optimized IR no longer depends on whether a plugin is loaded. ## What #4669 was working around With a target machine, the AMDGPU attributor runs at `OptimizerLastEP` before any plugin pass, because the target machine registers its callbacks in the `PassBuilder` constructor. It marks kernels `amdgpu-no-*` for implicit inputs they do not use. An instrumentation plugin that then injects calls reading those inputs (workgroup id y/z, the implicit-argument pointer, ...) leaves the kernel with attributes that no longer hold, and the backend does not set those inputs up. Dropping the target machine avoided this by never running the attributor, at the cost of disabling every target pass for all plugins. The correct fix is in the plugin, which should remove `amdgpu-no-*` from the functions it instruments. For Omniprobe, the instrumentation tool this guard was written for, that fix is AMDResearch/omniprobe#44 (reproduction and details in AMDResearch/omniprobe#43). I verified on gfx950 that with this PR and that fix, Omniprobe's instrumented kernels get the same implicit-input setup as today. ## Scope Today this only affects AMD. The NVIDIA backend calls `optimize_module` without `arch`, so it never creates a target machine, with or without a plugin. If #11156 lands and passes `arch` for NVIDIA, NVIDIA plugin users get the target machine as well. ## Testing - Merged current `main`; `python/test/unit/instrumentation/test_gpuhello.py` passes on gfx950 with `LLVM_PASS_PLUGIN_PATH` set, as in the AMD CI job. - The table above, before and after this change. - Omniprobe's Triton plugin with and without AMDResearch/omniprobe#44, as described above.

…en arch is specified (#10849)' (facebookexperimental#3911) Summary: Pull Request resolved: facebookexperimental#3911 This is a backport of an upstream PR: triton-lang/triton#10849 allow-large-files bypass-github-export-checks false positive from LLVM update Upstream commit message: ``` > [LLVM] Always set the target machine when arch is specified (#10849) > `optimize_module` drops the target machine whenever an LLVM-IR pass > plugin is loaded (`LLVM_PASS_PLUGIN_PATH`). That guard was added in > facebookexperimental#4669 for instrumentation plugins that splice in externally > (Clang-)compiled code. This PR creates the target machine whenever > `arch` is known, whether or not a plugin is loaded. > ## Why > Loading a plugin currently changes how the kernel itself is optimized, > not only what the plugin adds. Without a target machine, the O3 pipeline > runs without the AMDGPU target passes (for example `AMDGPUAttributor`) > and without target cost models. Measured on gfx950 with a vector-add > kernel: > | | no plugin | in-tree `GPUHello` plugin loaded | > |---|---:|---:| > | `amdgpu-no-*` attributes in LLIR, `main` | 21 | 0 | > | `amdgpu-no-*` attributes in LLIR, this PR | 21 | 21 | > `GPUHello` only prints, yet loading it removes every attribute the > attributor infers for this kernel. How much that matters depends on the > kernel; for a plain matmul the generated code is nearly identical, while > scheduler-style plugins, such as the out-of-tree LLIR scheduler used by > the gfx950 Gluon tutorials, lose performance without the target cost > models. With this PR the optimized IR no longer depends on whether a > plugin is loaded. > ## What facebookexperimental#4669 was working around > With a target machine, the AMDGPU attributor runs at `OptimizerLastEP` > before any plugin pass, because the target machine registers its > callbacks in the `PassBuilder` constructor. It marks kernels > `amdgpu-no-*` for implicit inputs they do not use. An instrumentation > plugin that then injects calls reading those inputs (workgroup id y/z, > the implicit-argument pointer, ...) leaves the kernel with attributes > that no longer hold, and the backend does not set those inputs up. > Dropping the target machine avoided this by never running the > attributor, at the cost of disabling every target pass for all plugins. > The correct fix is in the plugin, which should remove `amdgpu-no-*` from > the functions it instruments. For Omniprobe, the instrumentation tool > this guard was written for, that fix is AMDResearch/omniprobe#44 > (reproduction and details in AMDResearch/omniprobe#43). I verified on > gfx950 that with this PR and that fix, Omniprobe's instrumented kernels > get the same implicit-input setup as today. > ## Scope > Today this only affects AMD. The NVIDIA backend calls `optimize_module` > without `arch`, so it never creates a target machine, with or without a > plugin. If #11156 lands and passes `arch` for NVIDIA, NVIDIA plugin > users get the target machine as well. > ## Testing > - Merged current `main`; > `python/test/unit/instrumentation/test_gpuhello.py` passes on gfx950 > with `LLVM_PASS_PLUGIN_PATH` set, as in the AMD CI job. > - The table above, before and after this change. > - Omniprobe's Triton plugin with and without AMDResearch/omniprobe#44, > as described above. ``` ***Do not remove the following line from this commit*** Reactor Backport Revision: 4fd7cc5cfd0a84acfeb5dadb1578dffa409903fc --- This diff was generated by running: ``` buck run fbcode//triton/tools/reactor:reactor -- backport --pr 10849 ``` Reviewed By: agron911, stashuk-olek Differential Revision: D122891292

…en arch is specified (#10849)' Summary: This is a backport of an upstream PR: triton-lang/triton#10849 allow-large-files bypass-github-export-checks false positive from LLVM update Upstream commit message: ``` > [LLVM] Always set the target machine when arch is specified (#10849) > `optimize_module` drops the target machine whenever an LLVM-IR pass > plugin is loaded (`LLVM_PASS_PLUGIN_PATH`). That guard was added in > facebookexperimental#4669 for instrumentation plugins that splice in externally > (Clang-)compiled code. This PR creates the target machine whenever > `arch` is known, whether or not a plugin is loaded. > ## Why > Loading a plugin currently changes how the kernel itself is optimized, > not only what the plugin adds. Without a target machine, the O3 pipeline > runs without the AMDGPU target passes (for example `AMDGPUAttributor`) > and without target cost models. Measured on gfx950 with a vector-add > kernel: > | | no plugin | in-tree `GPUHello` plugin loaded | > |---|---:|---:| > | `amdgpu-no-*` attributes in LLIR, `main` | 21 | 0 | > | `amdgpu-no-*` attributes in LLIR, this PR | 21 | 21 | > `GPUHello` only prints, yet loading it removes every attribute the > attributor infers for this kernel. How much that matters depends on the > kernel; for a plain matmul the generated code is nearly identical, while > scheduler-style plugins, such as the out-of-tree LLIR scheduler used by > the gfx950 Gluon tutorials, lose performance without the target cost > models. With this PR the optimized IR no longer depends on whether a > plugin is loaded. > ## What facebookexperimental#4669 was working around > With a target machine, the AMDGPU attributor runs at `OptimizerLastEP` > before any plugin pass, because the target machine registers its > callbacks in the `PassBuilder` constructor. It marks kernels > `amdgpu-no-*` for implicit inputs they do not use. An instrumentation > plugin that then injects calls reading those inputs (workgroup id y/z, > the implicit-argument pointer, ...) leaves the kernel with attributes > that no longer hold, and the backend does not set those inputs up. > Dropping the target machine avoided this by never running the > attributor, at the cost of disabling every target pass for all plugins. > The correct fix is in the plugin, which should remove `amdgpu-no-*` from > the functions it instruments. For Omniprobe, the instrumentation tool > this guard was written for, that fix is AMDResearch/omniprobe#44 > (reproduction and details in AMDResearch/omniprobe#43). I verified on > gfx950 that with this PR and that fix, Omniprobe's instrumented kernels > get the same implicit-input setup as today. > ## Scope > Today this only affects AMD. The NVIDIA backend calls `optimize_module` > without `arch`, so it never creates a target machine, with or without a > plugin. If #11156 lands and passes `arch` for NVIDIA, NVIDIA plugin > users get the target machine as well. > ## Testing > - Merged current `main`; > `python/test/unit/instrumentation/test_gpuhello.py` passes on gfx950 > with `LLVM_PASS_PLUGIN_PATH` set, as in the AMD CI job. > - The table above, before and after this change. > - Omniprobe's Triton plugin with and without AMDResearch/omniprobe#44, > as described above. ``` ***Do not remove the following line from this commit*** Reactor Backport Revision: 4fd7cc5cfd0a84acfeb5dadb1578dffa409903fc --- This diff was generated by running: ``` buck run fbcode//triton/tools/reactor:reactor -- backport --pr 10849 ``` Reviewed By: agron911 Differential Revision: D122891292

…en arch is specified (#10849)' (facebookexperimental#3911) Summary: Pull Request resolved: facebookexperimental#3911 This is a backport of an upstream PR: triton-lang/triton#10849 allow-large-files bypass-github-export-checks false positive from LLVM update Upstream commit message: ``` > [LLVM] Always set the target machine when arch is specified (#10849) > `optimize_module` drops the target machine whenever an LLVM-IR pass > plugin is loaded (`LLVM_PASS_PLUGIN_PATH`). That guard was added in > facebookexperimental#4669 for instrumentation plugins that splice in externally > (Clang-)compiled code. This PR creates the target machine whenever > `arch` is known, whether or not a plugin is loaded. > ## Why > Loading a plugin currently changes how the kernel itself is optimized, > not only what the plugin adds. Without a target machine, the O3 pipeline > runs without the AMDGPU target passes (for example `AMDGPUAttributor`) > and without target cost models. Measured on gfx950 with a vector-add > kernel: > | | no plugin | in-tree `GPUHello` plugin loaded | > |---|---:|---:| > | `amdgpu-no-*` attributes in LLIR, `main` | 21 | 0 | > | `amdgpu-no-*` attributes in LLIR, this PR | 21 | 21 | > `GPUHello` only prints, yet loading it removes every attribute the > attributor infers for this kernel. How much that matters depends on the > kernel; for a plain matmul the generated code is nearly identical, while > scheduler-style plugins, such as the out-of-tree LLIR scheduler used by > the gfx950 Gluon tutorials, lose performance without the target cost > models. With this PR the optimized IR no longer depends on whether a > plugin is loaded. > ## What facebookexperimental#4669 was working around > With a target machine, the AMDGPU attributor runs at `OptimizerLastEP` > before any plugin pass, because the target machine registers its > callbacks in the `PassBuilder` constructor. It marks kernels > `amdgpu-no-*` for implicit inputs they do not use. An instrumentation > plugin that then injects calls reading those inputs (workgroup id y/z, > the implicit-argument pointer, ...) leaves the kernel with attributes > that no longer hold, and the backend does not set those inputs up. > Dropping the target machine avoided this by never running the > attributor, at the cost of disabling every target pass for all plugins. > The correct fix is in the plugin, which should remove `amdgpu-no-*` from > the functions it instruments. For Omniprobe, the instrumentation tool > this guard was written for, that fix is AMDResearch/omniprobe#44 > (reproduction and details in AMDResearch/omniprobe#43). I verified on > gfx950 that with this PR and that fix, Omniprobe's instrumented kernels > get the same implicit-input setup as today. > ## Scope > Today this only affects AMD. The NVIDIA backend calls `optimize_module` > without `arch`, so it never creates a target machine, with or without a > plugin. If #11156 lands and passes `arch` for NVIDIA, NVIDIA plugin > users get the target machine as well. > ## Testing > - Merged current `main`; > `python/test/unit/instrumentation/test_gpuhello.py` passes on gfx950 > with `LLVM_PASS_PLUGIN_PATH` set, as in the AMD CI job. > - The table above, before and after this change. > - Omniprobe's Triton plugin with and without AMDResearch/omniprobe#44, > as described above. ``` ***Do not remove the following line from this commit*** Reactor Backport Revision: 4fd7cc5cfd0a84acfeb5dadb1578dffa409903fc --- This diff was generated by running: ``` buck run fbcode//triton/tools/reactor:reactor -- backport --pr 10849 ``` Reviewed By: agron911, stashuk-olek Differential Revision: D122891292

…en arch is specified (#10849)' (facebookexperimental#3911) Summary: Pull Request resolved: facebookexperimental#3911 This is a backport of an upstream PR: triton-lang/triton#10849 allow-large-files bypass-github-export-checks false positive from LLVM update Upstream commit message: ``` > [LLVM] Always set the target machine when arch is specified (#10849) > `optimize_module` drops the target machine whenever an LLVM-IR pass > plugin is loaded (`LLVM_PASS_PLUGIN_PATH`). That guard was added in > facebookexperimental#4669 for instrumentation plugins that splice in externally > (Clang-)compiled code. This PR creates the target machine whenever > `arch` is known, whether or not a plugin is loaded. > ## Why > Loading a plugin currently changes how the kernel itself is optimized, > not only what the plugin adds. Without a target machine, the O3 pipeline > runs without the AMDGPU target passes (for example `AMDGPUAttributor`) > and without target cost models. Measured on gfx950 with a vector-add > kernel: > | | no plugin | in-tree `GPUHello` plugin loaded | > |---|---:|---:| > | `amdgpu-no-*` attributes in LLIR, `main` | 21 | 0 | > | `amdgpu-no-*` attributes in LLIR, this PR | 21 | 21 | > `GPUHello` only prints, yet loading it removes every attribute the > attributor infers for this kernel. How much that matters depends on the > kernel; for a plain matmul the generated code is nearly identical, while > scheduler-style plugins, such as the out-of-tree LLIR scheduler used by > the gfx950 Gluon tutorials, lose performance without the target cost > models. With this PR the optimized IR no longer depends on whether a > plugin is loaded. > ## What facebookexperimental#4669 was working around > With a target machine, the AMDGPU attributor runs at `OptimizerLastEP` > before any plugin pass, because the target machine registers its > callbacks in the `PassBuilder` constructor. It marks kernels > `amdgpu-no-*` for implicit inputs they do not use. An instrumentation > plugin that then injects calls reading those inputs (workgroup id y/z, > the implicit-argument pointer, ...) leaves the kernel with attributes > that no longer hold, and the backend does not set those inputs up. > Dropping the target machine avoided this by never running the > attributor, at the cost of disabling every target pass for all plugins. > The correct fix is in the plugin, which should remove `amdgpu-no-*` from > the functions it instruments. For Omniprobe, the instrumentation tool > this guard was written for, that fix is AMDResearch/omniprobe#44 > (reproduction and details in AMDResearch/omniprobe#43). I verified on > gfx950 that with this PR and that fix, Omniprobe's instrumented kernels > get the same implicit-input setup as today. > ## Scope > Today this only affects AMD. The NVIDIA backend calls `optimize_module` > without `arch`, so it never creates a target machine, with or without a > plugin. If #11156 lands and passes `arch` for NVIDIA, NVIDIA plugin > users get the target machine as well. > ## Testing > - Merged current `main`; > `python/test/unit/instrumentation/test_gpuhello.py` passes on gfx950 > with `LLVM_PASS_PLUGIN_PATH` set, as in the AMD CI job. > - The table above, before and after this change. > - Omniprobe's Triton plugin with and without AMDResearch/omniprobe#44, > as described above. ``` ***Do not remove the following line from this commit*** Reactor Backport Revision: 4fd7cc5cfd0a84acfeb5dadb1578dffa409903fc --- This diff was generated by running: ``` buck run fbcode//triton/tools/reactor:reactor -- backport --pr 10849 ``` Reviewed By: agron911, plotfi, stashuk-olek Differential Revision: D122891292

…en arch is specified (#10849)' (facebookexperimental#3911) Summary: Pull Request resolved: facebookexperimental#3911 This is a backport of an upstream PR: triton-lang/triton#10849 allow-large-files bypass-github-export-checks false positive from LLVM update Upstream commit message: ``` > [LLVM] Always set the target machine when arch is specified (#10849) > `optimize_module` drops the target machine whenever an LLVM-IR pass > plugin is loaded (`LLVM_PASS_PLUGIN_PATH`). That guard was added in > facebookexperimental#4669 for instrumentation plugins that splice in externally > (Clang-)compiled code. This PR creates the target machine whenever > `arch` is known, whether or not a plugin is loaded. > ## Why > Loading a plugin currently changes how the kernel itself is optimized, > not only what the plugin adds. Without a target machine, the O3 pipeline > runs without the AMDGPU target passes (for example `AMDGPUAttributor`) > and without target cost models. Measured on gfx950 with a vector-add > kernel: > | | no plugin | in-tree `GPUHello` plugin loaded | > |---|---:|---:| > | `amdgpu-no-*` attributes in LLIR, `main` | 21 | 0 | > | `amdgpu-no-*` attributes in LLIR, this PR | 21 | 21 | > `GPUHello` only prints, yet loading it removes every attribute the > attributor infers for this kernel. How much that matters depends on the > kernel; for a plain matmul the generated code is nearly identical, while > scheduler-style plugins, such as the out-of-tree LLIR scheduler used by > the gfx950 Gluon tutorials, lose performance without the target cost > models. With this PR the optimized IR no longer depends on whether a > plugin is loaded. > ## What facebookexperimental#4669 was working around > With a target machine, the AMDGPU attributor runs at `OptimizerLastEP` > before any plugin pass, because the target machine registers its > callbacks in the `PassBuilder` constructor. It marks kernels > `amdgpu-no-*` for implicit inputs they do not use. An instrumentation > plugin that then injects calls reading those inputs (workgroup id y/z, > the implicit-argument pointer, ...) leaves the kernel with attributes > that no longer hold, and the backend does not set those inputs up. > Dropping the target machine avoided this by never running the > attributor, at the cost of disabling every target pass for all plugins. > The correct fix is in the plugin, which should remove `amdgpu-no-*` from > the functions it instruments. For Omniprobe, the instrumentation tool > this guard was written for, that fix is AMDResearch/omniprobe#44 > (reproduction and details in AMDResearch/omniprobe#43). I verified on > gfx950 that with this PR and that fix, Omniprobe's instrumented kernels > get the same implicit-input setup as today. > ## Scope > Today this only affects AMD. The NVIDIA backend calls `optimize_module` > without `arch`, so it never creates a target machine, with or without a > plugin. If #11156 lands and passes `arch` for NVIDIA, NVIDIA plugin > users get the target machine as well. > ## Testing > - Merged current `main`; > `python/test/unit/instrumentation/test_gpuhello.py` passes on gfx950 > with `LLVM_PASS_PLUGIN_PATH` set, as in the AMD CI job. > - The table above, before and after this change. > - Omniprobe's Triton plugin with and without AMDResearch/omniprobe#44, > as described above. ``` ***Do not remove the following line from this commit*** Reactor Backport Revision: 4fd7cc5cfd0a84acfeb5dadb1578dffa409903fc --- This diff was generated by running: ``` buck run fbcode//triton/tools/reactor:reactor -- backport --pr 10849 ``` Reviewed By: agron911, stashuk-olek Differential Revision: D122891292

…en arch is specified (#10849)' (facebookexperimental#3911) Summary: Pull Request resolved: facebookexperimental#3911 This is a backport of an upstream PR: triton-lang/triton#10849 allow-large-files bypass-github-export-checks false positive from LLVM update Upstream commit message: ``` > [LLVM] Always set the target machine when arch is specified (#10849) > `optimize_module` drops the target machine whenever an LLVM-IR pass > plugin is loaded (`LLVM_PASS_PLUGIN_PATH`). That guard was added in > facebookexperimental#4669 for instrumentation plugins that splice in externally > (Clang-)compiled code. This PR creates the target machine whenever > `arch` is known, whether or not a plugin is loaded. > ## Why > Loading a plugin currently changes how the kernel itself is optimized, > not only what the plugin adds. Without a target machine, the O3 pipeline > runs without the AMDGPU target passes (for example `AMDGPUAttributor`) > and without target cost models. Measured on gfx950 with a vector-add > kernel: > | | no plugin | in-tree `GPUHello` plugin loaded | > |---|---:|---:| > | `amdgpu-no-*` attributes in LLIR, `main` | 21 | 0 | > | `amdgpu-no-*` attributes in LLIR, this PR | 21 | 21 | > `GPUHello` only prints, yet loading it removes every attribute the > attributor infers for this kernel. How much that matters depends on the > kernel; for a plain matmul the generated code is nearly identical, while > scheduler-style plugins, such as the out-of-tree LLIR scheduler used by > the gfx950 Gluon tutorials, lose performance without the target cost > models. With this PR the optimized IR no longer depends on whether a > plugin is loaded. > ## What facebookexperimental#4669 was working around > With a target machine, the AMDGPU attributor runs at `OptimizerLastEP` > before any plugin pass, because the target machine registers its > callbacks in the `PassBuilder` constructor. It marks kernels > `amdgpu-no-*` for implicit inputs they do not use. An instrumentation > plugin that then injects calls reading those inputs (workgroup id y/z, > the implicit-argument pointer, ...) leaves the kernel with attributes > that no longer hold, and the backend does not set those inputs up. > Dropping the target machine avoided this by never running the > attributor, at the cost of disabling every target pass for all plugins. > The correct fix is in the plugin, which should remove `amdgpu-no-*` from > the functions it instruments. For Omniprobe, the instrumentation tool > this guard was written for, that fix is AMDResearch/omniprobe#44 > (reproduction and details in AMDResearch/omniprobe#43). I verified on > gfx950 that with this PR and that fix, Omniprobe's instrumented kernels > get the same implicit-input setup as today. > ## Scope > Today this only affects AMD. The NVIDIA backend calls `optimize_module` > without `arch`, so it never creates a target machine, with or without a > plugin. If #11156 lands and passes `arch` for NVIDIA, NVIDIA plugin > users get the target machine as well. > ## Testing > - Merged current `main`; > `python/test/unit/instrumentation/test_gpuhello.py` passes on gfx950 > with `LLVM_PASS_PLUGIN_PATH` set, as in the AMD CI job. > - The table above, before and after this change. > - Omniprobe's Triton plugin with and without AMDResearch/omniprobe#44, > as described above. ``` ***Do not remove the following line from this commit*** Reactor Backport Revision: 4fd7cc5cfd0a84acfeb5dadb1578dffa409903fc --- This diff was generated by running: ``` buck run fbcode//triton/tools/reactor:reactor -- backport --pr 10849 ``` Reviewed By: agron911, plotfi, stashuk-olek Differential Revision: D122891292

…en arch is specified (#10849)' (facebookexperimental#3911) Summary: Pull Request resolved: facebookexperimental#3911 This is a backport of an upstream PR: triton-lang/triton#10849 allow-large-files bypass-github-export-checks false positive from LLVM update Upstream commit message: ``` > [LLVM] Always set the target machine when arch is specified (#10849) > `optimize_module` drops the target machine whenever an LLVM-IR pass > plugin is loaded (`LLVM_PASS_PLUGIN_PATH`). That guard was added in > facebookexperimental#4669 for instrumentation plugins that splice in externally > (Clang-)compiled code. This PR creates the target machine whenever > `arch` is known, whether or not a plugin is loaded. > ## Why > Loading a plugin currently changes how the kernel itself is optimized, > not only what the plugin adds. Without a target machine, the O3 pipeline > runs without the AMDGPU target passes (for example `AMDGPUAttributor`) > and without target cost models. Measured on gfx950 with a vector-add > kernel: > | | no plugin | in-tree `GPUHello` plugin loaded | > |---|---:|---:| > | `amdgpu-no-*` attributes in LLIR, `main` | 21 | 0 | > | `amdgpu-no-*` attributes in LLIR, this PR | 21 | 21 | > `GPUHello` only prints, yet loading it removes every attribute the > attributor infers for this kernel. How much that matters depends on the > kernel; for a plain matmul the generated code is nearly identical, while > scheduler-style plugins, such as the out-of-tree LLIR scheduler used by > the gfx950 Gluon tutorials, lose performance without the target cost > models. With this PR the optimized IR no longer depends on whether a > plugin is loaded. > ## What facebookexperimental#4669 was working around > With a target machine, the AMDGPU attributor runs at `OptimizerLastEP` > before any plugin pass, because the target machine registers its > callbacks in the `PassBuilder` constructor. It marks kernels > `amdgpu-no-*` for implicit inputs they do not use. An instrumentation > plugin that then injects calls reading those inputs (workgroup id y/z, > the implicit-argument pointer, ...) leaves the kernel with attributes > that no longer hold, and the backend does not set those inputs up. > Dropping the target machine avoided this by never running the > attributor, at the cost of disabling every target pass for all plugins. > The correct fix is in the plugin, which should remove `amdgpu-no-*` from > the functions it instruments. For Omniprobe, the instrumentation tool > this guard was written for, that fix is AMDResearch/omniprobe#44 > (reproduction and details in AMDResearch/omniprobe#43). I verified on > gfx950 that with this PR and that fix, Omniprobe's instrumented kernels > get the same implicit-input setup as today. > ## Scope > Today this only affects AMD. The NVIDIA backend calls `optimize_module` > without `arch`, so it never creates a target machine, with or without a > plugin. If #11156 lands and passes `arch` for NVIDIA, NVIDIA plugin > users get the target machine as well. > ## Testing > - Merged current `main`; > `python/test/unit/instrumentation/test_gpuhello.py` passes on gfx950 > with `LLVM_PASS_PLUGIN_PATH` set, as in the AMD CI job. > - The table above, before and after this change. > - Omniprobe's Triton plugin with and without AMDResearch/omniprobe#44, > as described above. ``` ***Do not remove the following line from this commit*** Reactor Backport Revision: 4fd7cc5cfd0a84acfeb5dadb1578dffa409903fc --- This diff was generated by running: ``` buck run fbcode//triton/tools/reactor:reactor -- backport --pr 10849 ``` Reviewed By: agron911, stashuk-olek Differential Revision: D122891292

…en arch is specified (#10849)' (facebookexperimental#3911) Summary: Pull Request resolved: facebookexperimental#3911 This is a backport of an upstream PR: triton-lang/triton#10849 allow-large-files bypass-github-export-checks false positive from LLVM update Upstream commit message: ``` > [LLVM] Always set the target machine when arch is specified (#10849) > `optimize_module` drops the target machine whenever an LLVM-IR pass > plugin is loaded (`LLVM_PASS_PLUGIN_PATH`). That guard was added in > facebookexperimental#4669 for instrumentation plugins that splice in externally > (Clang-)compiled code. This PR creates the target machine whenever > `arch` is known, whether or not a plugin is loaded. > ## Why > Loading a plugin currently changes how the kernel itself is optimized, > not only what the plugin adds. Without a target machine, the O3 pipeline > runs without the AMDGPU target passes (for example `AMDGPUAttributor`) > and without target cost models. Measured on gfx950 with a vector-add > kernel: > | | no plugin | in-tree `GPUHello` plugin loaded | > |---|---:|---:| > | `amdgpu-no-*` attributes in LLIR, `main` | 21 | 0 | > | `amdgpu-no-*` attributes in LLIR, this PR | 21 | 21 | > `GPUHello` only prints, yet loading it removes every attribute the > attributor infers for this kernel. How much that matters depends on the > kernel; for a plain matmul the generated code is nearly identical, while > scheduler-style plugins, such as the out-of-tree LLIR scheduler used by > the gfx950 Gluon tutorials, lose performance without the target cost > models. With this PR the optimized IR no longer depends on whether a > plugin is loaded. > ## What facebookexperimental#4669 was working around > With a target machine, the AMDGPU attributor runs at `OptimizerLastEP` > before any plugin pass, because the target machine registers its > callbacks in the `PassBuilder` constructor. It marks kernels > `amdgpu-no-*` for implicit inputs they do not use. An instrumentation > plugin that then injects calls reading those inputs (workgroup id y/z, > the implicit-argument pointer, ...) leaves the kernel with attributes > that no longer hold, and the backend does not set those inputs up. > Dropping the target machine avoided this by never running the > attributor, at the cost of disabling every target pass for all plugins. > The correct fix is in the plugin, which should remove `amdgpu-no-*` from > the functions it instruments. For Omniprobe, the instrumentation tool > this guard was written for, that fix is AMDResearch/omniprobe#44 > (reproduction and details in AMDResearch/omniprobe#43). I verified on > gfx950 that with this PR and that fix, Omniprobe's instrumented kernels > get the same implicit-input setup as today. > ## Scope > Today this only affects AMD. The NVIDIA backend calls `optimize_module` > without `arch`, so it never creates a target machine, with or without a > plugin. If #11156 lands and passes `arch` for NVIDIA, NVIDIA plugin > users get the target machine as well. > ## Testing > - Merged current `main`; > `python/test/unit/instrumentation/test_gpuhello.py` passes on gfx950 > with `LLVM_PASS_PLUGIN_PATH` set, as in the AMD CI job. > - The table above, before and after this change. > - Omniprobe's Triton plugin with and without AMDResearch/omniprobe#44, > as described above. ``` ***Do not remove the following line from this commit*** Reactor Backport Revision: 4fd7cc5cfd0a84acfeb5dadb1578dffa409903fc --- This diff was generated by running: ``` buck run fbcode//triton/tools/reactor:reactor -- backport --pr 10849 ``` Reviewed By: agron911, plotfi, stashuk-olek Differential Revision: D122891292

…en arch is specified (#10849)' (facebookexperimental#3911) Summary: Pull Request resolved: facebookexperimental#3911 This is a backport of an upstream PR: triton-lang/triton#10849 allow-large-files bypass-github-export-checks false positive from LLVM update Upstream commit message: ``` > [LLVM] Always set the target machine when arch is specified (#10849) > `optimize_module` drops the target machine whenever an LLVM-IR pass > plugin is loaded (`LLVM_PASS_PLUGIN_PATH`). That guard was added in > facebookexperimental#4669 for instrumentation plugins that splice in externally > (Clang-)compiled code. This PR creates the target machine whenever > `arch` is known, whether or not a plugin is loaded. > ## Why > Loading a plugin currently changes how the kernel itself is optimized, > not only what the plugin adds. Without a target machine, the O3 pipeline > runs without the AMDGPU target passes (for example `AMDGPUAttributor`) > and without target cost models. Measured on gfx950 with a vector-add > kernel: > | | no plugin | in-tree `GPUHello` plugin loaded | > |---|---:|---:| > | `amdgpu-no-*` attributes in LLIR, `main` | 21 | 0 | > | `amdgpu-no-*` attributes in LLIR, this PR | 21 | 21 | > `GPUHello` only prints, yet loading it removes every attribute the > attributor infers for this kernel. How much that matters depends on the > kernel; for a plain matmul the generated code is nearly identical, while > scheduler-style plugins, such as the out-of-tree LLIR scheduler used by > the gfx950 Gluon tutorials, lose performance without the target cost > models. With this PR the optimized IR no longer depends on whether a > plugin is loaded. > ## What facebookexperimental#4669 was working around > With a target machine, the AMDGPU attributor runs at `OptimizerLastEP` > before any plugin pass, because the target machine registers its > callbacks in the `PassBuilder` constructor. It marks kernels > `amdgpu-no-*` for implicit inputs they do not use. An instrumentation > plugin that then injects calls reading those inputs (workgroup id y/z, > the implicit-argument pointer, ...) leaves the kernel with attributes > that no longer hold, and the backend does not set those inputs up. > Dropping the target machine avoided this by never running the > attributor, at the cost of disabling every target pass for all plugins. > The correct fix is in the plugin, which should remove `amdgpu-no-*` from > the functions it instruments. For Omniprobe, the instrumentation tool > this guard was written for, that fix is AMDResearch/omniprobe#44 > (reproduction and details in AMDResearch/omniprobe#43). I verified on > gfx950 that with this PR and that fix, Omniprobe's instrumented kernels > get the same implicit-input setup as today. > ## Scope > Today this only affects AMD. The NVIDIA backend calls `optimize_module` > without `arch`, so it never creates a target machine, with or without a > plugin. If #11156 lands and passes `arch` for NVIDIA, NVIDIA plugin > users get the target machine as well. > ## Testing > - Merged current `main`; > `python/test/unit/instrumentation/test_gpuhello.py` passes on gfx950 > with `LLVM_PASS_PLUGIN_PATH` set, as in the AMD CI job. > - The table above, before and after this change. > - Omniprobe's Triton plugin with and without AMDResearch/omniprobe#44, > as described above. ``` ***Do not remove the following line from this commit*** Reactor Backport Revision: 4fd7cc5cfd0a84acfeb5dadb1578dffa409903fc --- This diff was generated by running: ``` buck run fbcode//triton/tools/reactor:reactor -- backport --pr 10849 ``` Reviewed By: agron911, stashuk-olek Differential Revision: D122891292

…en arch is specified (#10849)' (facebookexperimental#3911) Summary: Pull Request resolved: facebookexperimental#3911 This is a backport of an upstream PR: triton-lang/triton#10849 allow-large-files bypass-github-export-checks false positive from LLVM update Upstream commit message: ``` > [LLVM] Always set the target machine when arch is specified (#10849) > `optimize_module` drops the target machine whenever an LLVM-IR pass > plugin is loaded (`LLVM_PASS_PLUGIN_PATH`). That guard was added in > facebookexperimental#4669 for instrumentation plugins that splice in externally > (Clang-)compiled code. This PR creates the target machine whenever > `arch` is known, whether or not a plugin is loaded. > ## Why > Loading a plugin currently changes how the kernel itself is optimized, > not only what the plugin adds. Without a target machine, the O3 pipeline > runs without the AMDGPU target passes (for example `AMDGPUAttributor`) > and without target cost models. Measured on gfx950 with a vector-add > kernel: > | | no plugin | in-tree `GPUHello` plugin loaded | > |---|---:|---:| > | `amdgpu-no-*` attributes in LLIR, `main` | 21 | 0 | > | `amdgpu-no-*` attributes in LLIR, this PR | 21 | 21 | > `GPUHello` only prints, yet loading it removes every attribute the > attributor infers for this kernel. How much that matters depends on the > kernel; for a plain matmul the generated code is nearly identical, while > scheduler-style plugins, such as the out-of-tree LLIR scheduler used by > the gfx950 Gluon tutorials, lose performance without the target cost > models. With this PR the optimized IR no longer depends on whether a > plugin is loaded. > ## What facebookexperimental#4669 was working around > With a target machine, the AMDGPU attributor runs at `OptimizerLastEP` > before any plugin pass, because the target machine registers its > callbacks in the `PassBuilder` constructor. It marks kernels > `amdgpu-no-*` for implicit inputs they do not use. An instrumentation > plugin that then injects calls reading those inputs (workgroup id y/z, > the implicit-argument pointer, ...) leaves the kernel with attributes > that no longer hold, and the backend does not set those inputs up. > Dropping the target machine avoided this by never running the > attributor, at the cost of disabling every target pass for all plugins. > The correct fix is in the plugin, which should remove `amdgpu-no-*` from > the functions it instruments. For Omniprobe, the instrumentation tool > this guard was written for, that fix is AMDResearch/omniprobe#44 > (reproduction and details in AMDResearch/omniprobe#43). I verified on > gfx950 that with this PR and that fix, Omniprobe's instrumented kernels > get the same implicit-input setup as today. > ## Scope > Today this only affects AMD. The NVIDIA backend calls `optimize_module` > without `arch`, so it never creates a target machine, with or without a > plugin. If #11156 lands and passes `arch` for NVIDIA, NVIDIA plugin > users get the target machine as well. > ## Testing > - Merged current `main`; > `python/test/unit/instrumentation/test_gpuhello.py` passes on gfx950 > with `LLVM_PASS_PLUGIN_PATH` set, as in the AMD CI job. > - The table above, before and after this change. > - Omniprobe's Triton plugin with and without AMDResearch/omniprobe#44, > as described above. ``` ***Do not remove the following line from this commit*** Reactor Backport Revision: 4fd7cc5cfd0a84acfeb5dadb1578dffa409903fc --- This diff was generated by running: ``` buck run fbcode//triton/tools/reactor:reactor -- backport --pr 10849 ``` Reviewed By: agron911, plotfi, stashuk-olek Differential Revision: D122891292

…en arch is specified (#10849)' (facebookexperimental#3911) Summary: This is a backport of an upstream PR: triton-lang/triton#10849 allow-large-files bypass-github-export-checks false positive from LLVM update Upstream commit message: ``` > [LLVM] Always set the target machine when arch is specified (#10849) > `optimize_module` drops the target machine whenever an LLVM-IR pass > plugin is loaded (`LLVM_PASS_PLUGIN_PATH`). That guard was added in > facebookexperimental#4669 for instrumentation plugins that splice in externally > (Clang-)compiled code. This PR creates the target machine whenever > `arch` is known, whether or not a plugin is loaded. > ## Why > Loading a plugin currently changes how the kernel itself is optimized, > not only what the plugin adds. Without a target machine, the O3 pipeline > runs without the AMDGPU target passes (for example `AMDGPUAttributor`) > and without target cost models. Measured on gfx950 with a vector-add > kernel: > | | no plugin | in-tree `GPUHello` plugin loaded | > |---|---:|---:| > | `amdgpu-no-*` attributes in LLIR, `main` | 21 | 0 | > | `amdgpu-no-*` attributes in LLIR, this PR | 21 | 21 | > `GPUHello` only prints, yet loading it removes every attribute the > attributor infers for this kernel. How much that matters depends on the > kernel; for a plain matmul the generated code is nearly identical, while > scheduler-style plugins, such as the out-of-tree LLIR scheduler used by > the gfx950 Gluon tutorials, lose performance without the target cost > models. With this PR the optimized IR no longer depends on whether a > plugin is loaded. > ## What facebookexperimental#4669 was working around > With a target machine, the AMDGPU attributor runs at `OptimizerLastEP` > before any plugin pass, because the target machine registers its > callbacks in the `PassBuilder` constructor. It marks kernels > `amdgpu-no-*` for implicit inputs they do not use. An instrumentation > plugin that then injects calls reading those inputs (workgroup id y/z, > the implicit-argument pointer, ...) leaves the kernel with attributes > that no longer hold, and the backend does not set those inputs up. > Dropping the target machine avoided this by never running the > attributor, at the cost of disabling every target pass for all plugins. > The correct fix is in the plugin, which should remove `amdgpu-no-*` from > the functions it instruments. For Omniprobe, the instrumentation tool > this guard was written for, that fix is AMDResearch/omniprobe#44 > (reproduction and details in AMDResearch/omniprobe#43). I verified on > gfx950 that with this PR and that fix, Omniprobe's instrumented kernels > get the same implicit-input setup as today. > ## Scope > Today this only affects AMD. The NVIDIA backend calls `optimize_module` > without `arch`, so it never creates a target machine, with or without a > plugin. If #11156 lands and passes `arch` for NVIDIA, NVIDIA plugin > users get the target machine as well. > ## Testing > - Merged current `main`; > `python/test/unit/instrumentation/test_gpuhello.py` passes on gfx950 > with `LLVM_PASS_PLUGIN_PATH` set, as in the AMD CI job. > - The table above, before and after this change. > - Omniprobe's Triton plugin with and without AMDResearch/omniprobe#44, > as described above. ``` ***Do not remove the following line from this commit*** Reactor Backport Revision: 4fd7cc5cfd0a84acfeb5dadb1578dffa409903fc --- This diff was generated by running: ``` buck run fbcode//triton/tools/reactor:reactor -- backport --pr 10849 ``` Reviewed By: agron911, plotfi, stashuk-olek Differential Revision: D122891292