[ROCM] Validate that target triple is set in module and test for it (#24663)

Since LLVM commit
00a6186128d3 AMDGPU: Prefer getting the triple from the module over the
TargetMachine (#206055) the target riple of the module has to be set.

There was a suggestion for a dedicated test in
https://github.com/iree-org/iree/pull/24662#discussion_r3512817954.

Add the check for the target triple to the ROCM module validation and
add a LIT test for it.

This change has been assisted by Claude Code. The changes have been
reviewed manually before opening the PR.

Signed-off-by: Stefan Schuermans <schuermans@roofline.ai>

compiler/plugins/target/ROCM/ROCMTarget.cpp

diff --git a/compiler/plugins/target/ROCM/ROCMTarget.cpp b/compiler/plugins/target/ROCM/ROCMTarget.cpp
index 32e74b8..0b1d80f 100644
--- a/compiler/plugins/target/ROCM/ROCMTarget.cpp
+++ b/compiler/plugins/target/ROCM/ROCMTarget.cpp
@@ -655,6 +655,20 @@
   LogicalResult
   validateFinalizedModule(IREE::HAL::ExecutableVariantOp variantOp,
                           llvm::Module &module, bool allowExternalDecls) {
+    // The LLVM backend reads the target triple from the module itself (not just
+    // from the TargetMachine) to decide how to emit the code object and HSA
+    // metadata. If it is left unset, code generation crashes deep inside the
+    // AMDGPU AsmPrinter with an opaque signal (a null HSAMetadataStream
+    // dereference; see llvm-project@00a6186128d3). Fail early here with an
+    // actionable message instead. The triple is set in serializeExecutable
+    // right after the data layout.
+    if (module.getTargetTriple().str().empty()) {
+      return variantOp.emitError()
+             << "the finalized LLVM module has no target triple set; it must "
+                "be set before serialization (see ROCMTargetBackend::"
+                "serializeExecutable). Without it the LLVM AMDGPU backend "
+                "crashes while emitting the object file.";
+    }
     for (llvm::Function &func : module.functions()) {
       if (func.isDeclaration() && !func.isIntrinsic() && !func.use_empty()) {
         // In SPIR-V mode, external declarations (e.g. __ocml_*, __ockl_*,

compiler/plugins/target/ROCM/test/CMakeLists.txt

diff --git a/compiler/plugins/target/ROCM/test/CMakeLists.txt b/compiler/plugins/target/ROCM/test/CMakeLists.txt
index 9df494d..7487c87 100644
--- a/compiler/plugins/target/ROCM/test/CMakeLists.txt
+++ b/compiler/plugins/target/ROCM/test/CMakeLists.txt
@@ -47,6 +47,7 @@
   SRCS
     "emit_debug_info.mlir"
     "external_function_validation.mlir"
+    "module_target_triple.mlir"
     "smoketest.mlir"
     "smoketest_hsaco.mlir"
     "target_device_features.mlir"

compiler/plugins/target/ROCM/test/module_target_triple.mlir

diff --git a/compiler/plugins/target/ROCM/test/module_target_triple.mlir b/compiler/plugins/target/ROCM/test/module_target_triple.mlir
new file mode 100644
index 0000000..5ef20c2
--- /dev/null
+++ b/compiler/plugins/target/ROCM/test/module_target_triple.mlir
@@ -0,0 +1,28 @@
+// RUN: rm -rf %t && mkdir -p %t
+// RUN: iree-opt --iree-hal-transformation-pipeline --iree-hal-target-device=hip \
+// RUN:   --iree-rocm-target=gfx942 --iree-hal-dump-executable-intermediates-to=%t %s -o /dev/null
+// RUN: cat %t/*.optimized.ll | FileCheck %s
+
+// Regression test for the target triple being set on the finalized LLVM module.
+//
+// The ROCm target backend must set the target triple of the module, because
+// the LLVM backend expects it to be set.
+// See llvm-project@00a6186128d3: ("AMDGPU: Prefer getting the triple from the
+// module over the TargetMachine").
+// Therefore, the ROCMTargetBackend validation checks for the target triple to
+// be set and emits an explicit error if that is not the case.
+
+// CHECK: target triple = "amdgcn-amd-amdhsa"
+
+#pipeline_layout = #hal.pipeline.layout<bindings = []>
+hal.executable.source public @exe {
+  hal.executable.export public @empty ordinal(0) layout(#pipeline_layout) count(%arg0: !hal.device) -> (index, index, index) {
+    %c1 = arith.constant 1 : index
+    hal.return %c1, %c1, %c1 : index, index, index
+  } attributes {subgroup_size = 64 : index, workgroup_size = [1 : index, 1 : index, 1 : index]}
+  builtin.module {
+    llvm.func @empty() attributes {rocdl.kernel} {
+      llvm.return
+    }
+  }
+}