[AMD][Gluon] support inline when amdgpu ops in the callee (triton-lang#8041)

Before this change, when there are amdgpu ops in the kernel
as callee, inlining will fail.
A`gluon-inline` pass runs after ttgir is generated by the
frontend, so `run_parser` in test_frontend will not trigger it.
Therefore, a runtime testing using `warmup` is added.

Author: zwu-2025

python/test/gluon/test_core.py

@@ -811,3 +811,26 @@ def kernel(N, out):
     out = torch.empty(1, dtype=torch.int32, device="cuda")
     compiled_kernel = kernel.warmup(N=100, out=out, grid=(1, ))
     assert compiled_kernel.asm["llir"].count("define") == 1
+
+
+@pytest.mark.skipif(not is_hip_cdna3() and not is_hip_cdna4(), reason="Requires CDNA3 or CDNA4")
+def test_inline_with_amdgpu_dialect():
+
+    @gluon.jit
+    def buffer_load(x, offsets):
+        return ttgl.amd.cdna3.buffer_load(ptr=x, offsets=offsets)
+
+    @gluon.jit
+    def kernel(x, y):
+        layout: ttgl.constexpr = ttgl.BlockedLayout(size_per_thread=[1], threads_per_warp=[64], warps_per_cta=[4],
+                                                    order=[0])
+        offsets = ttgl.arange(0, 64, layout=layout)
+
+        a = buffer_load(x, offsets)
+        ttgl.amd.cdna3.buffer_store(stored_value=a, ptr=y, offsets=offsets)
+
+    input = torch.arange(64, device="cuda").to(torch.int32)
+    output = torch.empty_like(input)
+
+    compiled_kernel = kernel.warmup(input, output, grid=(1, ))
+    assert compiled_kernel.asm["ttgir"].count("tt.func private") == 0

third_party/amd/lib/Dialect/TritonAMDGPU/IR/Dialect.cpp

@@ -27,6 +27,7 @@
 #include "mlir/IR/OpImplementation.h"
 #include "third_party/amd/include/Utils/Utility.h"
 #include "triton/Conversion/TritonGPUToLLVM/Utility.h"
+#include "triton/Dialect/Triton/IR/Interfaces.h"
 #include "triton/Tools/LayoutUtils.h"
 #include "llvm/ADT/TypeSwitch.h"
 
@@ -50,6 +51,8 @@ void mlir::triton::amdgpu::TritonAMDGPUDialect::initialize() {
 #define GET_OP_LIST
 #include "Dialect/TritonAMDGPU/IR/Ops.cpp.inc"
       >();
+
+  addInterfaces<TritonInlinerInterface>();
 }
 
 #include "Dialect/TritonAMDGPU/IR/TritonAMDGPUEnums.cpp.inc"