fix(rocisa): stop reading a closed pipe fd in POSIX run()

projects/hipblaslt/tensilelite/rocisa/rocisa/src/helper.cpp

@@ -234,6 +234,7 @@ std::string nullDevicePath()
 
 #else
 // POSIX implementation.
+#include <fcntl.h>
 #include <sys/wait.h>
 #include <unistd.h>
 #ifdef __GNUG__
@@ -242,24 +243,43 @@ std::string nullDevicePath()
 
 std::pair<int, std::string> run(const std::vector<char*>& cmd, const std::string& input, bool debug)
 {
+    // The pipe only feeds `input` to the child's stdin. O_CLOEXEC keeps it out of
+    // children forked concurrently by other threads: an inherited write end would
+    // keep this child's stdin from ever seeing EOF.
     int   p[2];
     pid_t pid;
-    if(pipe(p) == -1)
+    if(pipe2(p, O_CLOEXEC) == -1)
     {
         throw std::runtime_error("cmd failed!");
     }
 
     pid = fork();
-    if(pid == 0)
+    if(pid == -1)
     {
+        close(p[0]);
         close(p[1]);
-        dup2(p[0], STDIN_FILENO);
+        throw std::runtime_error("fork failed!");
+    }
+    if(pid == 0)
+    {
+        // dup2 clears FD_CLOEXEC on the new descriptor; the pipe fds themselves are
+        // closed by execvp.
+        if(p[0] == STDIN_FILENO)
+            fcntl(p[0], F_SETFD, 0);
+        else
+            dup2(p[0], STDIN_FILENO);
         if(!debug)
         {
-            dup2(p[0], STDERR_FILENO);
-            dup2(p[0], STDOUT_FILENO);
+            // Discard the child's output; only its exit status is reported.
+            int devNull = open("/dev/null", O_WRONLY);
+            if(devNull != -1)
+            {
+                dup2(devNull, STDOUT_FILENO);
+                dup2(devNull, STDERR_FILENO);
+                if(devNull > STDERR_FILENO)
+                    close(devNull);
+            }
         }
-        close(p[0]);
         execvp(cmd[0], cmd.data());
         perror("execvp");
         // 127 is the conventional "command not found" status; _exit (not exit) is
@@ -271,11 +291,9 @@ std::pair<int, std::string> run(const std::vector<char*>& cmd, const std::string
         close(p[0]);
         write(p[1], input.c_str(), input.size());
         close(p[1]);
-        char        buf[128] = {0};
+        // The child's output is not captured (see above), so there is no result text.
         std::string result;
-        read(p[0], buf, 128);
-        result += buf;
-        int rcode;
+        int         rcode;
         waitpid(pid, &rcode, 0);
         if(WIFEXITED(rcode))
         {

projects/hipblaslt/tensilelite/rocisa/test/test_asm_caps_closed_stdin.py

@@ -0,0 +1,61 @@
+# Copyright Advanced Micro Devices, Inc., or its affiliates.
+# SPDX-License-Identifier: MIT
+
+import json
+import os
+import shutil
+import subprocess
+import sys
+
+import pytest
+
+_ISA = (9, 0, 10)
+
+# rocIsa caches capabilities per ISA for the life of the process, so each probe
+# runs in a fresh interpreter that imports the same rocisa as this one.
+_PROBE = """
+import json, sys
+sys.path[:] = json.loads(sys.argv[1])
+import rocisa
+isa = tuple(json.loads(sys.argv[2]))
+ti = rocisa.rocIsa.getInstance()
+ti.init(isa, sys.argv[3], False)
+ti.setKernel(isa, 64)
+print(json.dumps(dict(ti.getAsmCaps())))
+"""
+
+
+def _assembler():
+    rocm_path = os.environ.get("ROCM_PATH", "/opt/rocm")
+    search_path = os.pathsep.join(
+        [
+            os.path.join(rocm_path, "bin"),
+            os.path.join(rocm_path, "lib", "llvm", "bin"),
+        ]
+    )
+    return shutil.which("amdclang++", path=search_path) or "amdclang++"
+
+
+def _probe_asm_caps(close_stdin):
+    parent_path = json.dumps([os.path.abspath(p) for p in sys.path])
+    result = subprocess.run(
+        [sys.executable, "-c", _PROBE, parent_path, json.dumps(_ISA), _assembler()],
+        preexec_fn=(lambda: os.close(0)) if close_stdin else None,
+        capture_output=True,
+        text=True,
+    )
+    assert result.returncode == 0, result.stderr
+    return json.loads(result.stdout)
+
+
+@pytest.mark.skipif(os.name != "posix", reason="closes fd 0 in the probe process")
+def test_asm_caps_probe_matches_with_closed_stdin():
+    # With fd 0 closed, the probe pipe's read end is allocated as fd 0. The
+    # assembler must still get the probe source on stdin; otherwise it assembles
+    # nothing, succeeds, and every capability is reported as supported.
+    normal_caps = _probe_asm_caps(close_stdin=False)
+    closed_stdin_caps = _probe_asm_caps(close_stdin=True)
+
+    assert normal_caps["SupportedISA"] == 1
+    assert 0 in normal_caps.values()
+    assert closed_stdin_caps == normal_caps