merge aomp-dev into amd-staging

bin/run_composable-kernels.sh

@@ -318,13 +318,7 @@ else
 fi
 
 # For some reason, CK on gfx12 wants this set.
-CKCmakeCmd+="-DBUILD_DEV=On "
-
-# TODO: Remove this workaround when possible.
-# As of 2026-JUN-12 this is needed to bypass DPP-related build errors:
-#   error: cannot compile inline asm
-#   <inline asm>:2:33: error: not a valid operand.
-CKCmakeCmd+="-DDISABLE_DPP_KERNELS=ON"
+CKCmakeCmd+="-DBUILD_DEV=On"
 
 # Ensure CK build directory is cleaned.
 if [ "${ShouldRebuildCK}" == 'yes' ]; then

bin/semaphore_isolation.src

@@ -0,0 +1,33 @@
+# shellcheck shell=bash
+#
+# Copyright © Advanced Micro Devices, Inc., or its affiliates.
+#
+# SPDX-License-Identifier:  MIT
+#
+# semaphore_isolation.src - shared helper for GNU parallel's "sem" semaphore.
+#
+# Source this file and call isolate_semaphore before any "sem" use. It gives
+# "sem" a private, self-cleaning semaphore namespace (its own PARALLEL_HOME) so
+# that a leaked/stale token can neither deadlock a run nor leak across runs, and
+# optionally bounds a stuck semaphore with a timeout.
+#
+# Background: on hosts whose name contains uppercase letters, GNU parallel's
+# dead-lock reaper (remove_dead_locks) never matches its own token files, so a
+# token left behind by an unclean death is never reclaimed and every later
+# "sem --wait" deadlocks. A per-run private namespace sidesteps this entirely.
+#
+# Design:
+#   - No-op when PARALLEL_HOME is already set, so an outer caller never collides.
+#   - AOMP_SEMAPHORE_TIMEOUT (seconds), if set, is passed to every "sem" call as
+#     --semaphoretimeout: a wedged sem/sem --wait then steals and proceeds
+#     instead of hanging. Left unset => pure isolation, no timeout.
+
+isolate_semaphore() {
+  [ -n "${PARALLEL_HOME:-}" ] && return 0
+  _semaphore_home=$(mktemp -d "${TMPDIR:-/tmp}/aomp-sem.XXXXXX") || return 0
+  export PARALLEL_HOME="$_semaphore_home"
+  if [ -n "${AOMP_SEMAPHORE_TIMEOUT:-}" ]; then
+    export PARALLEL="--semaphoretimeout ${AOMP_SEMAPHORE_TIMEOUT}${PARALLEL:+ $PARALLEL}"
+  fi
+  trap 'rm -rf "$_semaphore_home"' EXIT
+}

test/smoke-dev/OutputHipDevCov/Makefile

@@ -0,0 +1,30 @@
+include ../../Makefile.defs
+
+# HipDevCov: device-side coverage drained via the in-tree HSA introspection
+# pass. main (TESTNAME) is built by the harness with source-based coverage; it
+# hipModuleLoad()s mod.co (built+extracted by run_and_check.sh) whose kernel has
+# no host shadow and is therefore only reachable by the HSA drain.
+
+TESTNAME     = main
+TESTSRC_MAIN = main.hip
+TESTSRC_AUX  =
+TESTSRC_ALL  = $(TESTSRC_MAIN) $(TESTSRC_AUX)
+
+AOMPHIP     ?= $(AOMP)
+ROCM_PATH   ?= $(AOMPHIP)
+
+# Extra budget: run_and_check.sh also compiles mod.hip and extracts mod.co.
+TIMEOUT      = 300s
+
+CFLAGS       = -x hip $(TARGET) -fno-gpu-rdc -fprofile-instr-generate -fcoverage-mapping --rocm-path=$(ROCM_PATH)
+LINK_FLAGS   = -L$(ROCM_PATH)/lib -lamdhip64 -Wl,-rpath,$(ROCM_PATH)/lib
+
+CC           = $(AOMP)/bin/clang $(VERBOSE)
+
+RUNCMD       = ./run_and_check.sh "$(AOMP)" "$(GPU_W_FEATURES)" "$(FILECHECK)" "$(ROCM_PATH)"
+
+include ../Makefile.rules
+
+clean::
+	rm -f mod.co builder builder.*.host-* builder.*.hip-amdgcn-amd-amdhsa--* \
+	      *.profraw merged.profdata

test/smoke-dev/OutputHipDevCov/main.hip

@@ -0,0 +1,62 @@
+// HipDevCov: prove the HSA-introspection device-coverage drain captures device
+// code that the HIP host-shadow drain cannot see.
+//
+// This program contains exactly one statically-registered kernel (host_kernel)
+// -- it has a host-side shadow, so the host-shadow drain covers it. It then
+// hipModuleLoad()s mod.co (built+extracted by run_and_check.sh) and launches
+// mod_kernel, which is NOT compiled into this program and therefore has no host
+// shadow here. mod_kernel's device counters can only reach the profile via the
+// HSA drain. run_and_check.sh asserts BOTH kernels appear in the merged
+// profile, so a passing run requires the HSA drain to be working.
+
+#include <hip/hip_runtime.h>
+#include <cstdio>
+
+__global__ void host_kernel(int *p, int n) {
+  int i = blockIdx.x * blockDim.x + threadIdx.x;
+  if (i < n)
+    p[i] = p[i] * 3;   // taken region
+  else
+    p[0] = p[0] - 1;   // untaken region
+}
+
+int main() {
+  const int n = 64;
+  int *d = nullptr;
+  if (hipMalloc(&d, n * sizeof(int)) != hipSuccess)
+    return 1;
+  if (hipMemset(d, 0, n * sizeof(int)) != hipSuccess)
+    return 2;
+
+  // (1) host-shadow path: a normally-registered kernel.
+  host_kernel<<<dim3(1), dim3(64)>>>(d, n);
+  if (hipDeviceSynchronize() != hipSuccess)
+    return 3;
+
+  // (2) HSA-only path: a kernel from a separately loaded code object, with no
+  // host shadow in this process.
+  hipModule_t mod;
+  if (hipModuleLoad(&mod, "mod.co") != hipSuccess) {
+    fprintf(stderr, "hipModuleLoad(mod.co) failed\n");
+    return 4;
+  }
+  hipFunction_t fn;
+  if (hipModuleGetFunction(&fn, mod, "mod_kernel") != hipSuccess) {
+    fprintf(stderr, "hipModuleGetFunction(mod_kernel) failed\n");
+    return 5;
+  }
+  int nn = n;
+  void *args[] = {&d, &nn};
+  if (hipModuleLaunchKernel(fn, 1, 1, 1, 64, 1, 1, 0, nullptr, args, nullptr) !=
+      hipSuccess)
+    return 6;
+  if (hipDeviceSynchronize() != hipSuccess)
+    return 7;
+
+  // Intentionally do NOT hipModuleUnload(mod): the HSA-introspection drain runs
+  // at process exit and walks the code objects still loaded on the agent.
+  // Unloading mod.co here would remove its counters before the drain sees them.
+  (void)hipFree(d);
+  printf("HipDevCov: ran host_kernel + module mod_kernel\n");
+  return 0;
+}

test/smoke-dev/OutputHipDevCov/mod.hip

@@ -0,0 +1,42 @@
+// Device-only kernel used to produce a standalone, loadable code object
+// (mod.co) that carries its own instrumentation (__llvm_profile_sections +
+// the device profile runtime). main.hip hipModuleLoad()s this code object, so
+// in that process mod_kernel has NO host-side shadow registration. The HIP
+// host-shadow drain therefore cannot see it; only the HSA-introspection drain
+// (which walks every code object loaded on the agent) can collect its
+// counters. That is exactly what this test asserts.
+//
+// When built with -DBUILD_MODULE_EXE this file is a normal HIP executable; the
+// full executable link makes clang's driver add the amdgcn device profile
+// runtime to the device link (addProfileRTLibs), so the embedded device code
+// object is fully instrumented. run_and_check.sh extracts that object as
+// mod.co. (The standalone --offload-device-only / hipcc --genco paths do NOT
+// add the device profile RT, leaving __llvm_profile_instrument_gpu /
+// __llvm_profile_raw_version unresolved -- hence the extract-from-executable
+// approach.)
+
+#include <hip/hip_runtime.h>
+
+extern "C" __global__ void mod_kernel(int *p, int n) {
+  int i = blockIdx.x * blockDim.x + threadIdx.x;
+  if (i < n)
+    p[i] = p[i] + 7;   // taken region
+  else
+    p[0] = p[0] - 1;   // untaken region (non-trivial coverage)
+}
+
+#ifdef BUILD_MODULE_EXE
+int main() {
+  const int n = 64;
+  int *d = nullptr;
+  if (hipMalloc(&d, n * sizeof(int)) != hipSuccess)
+    return 1;
+  if (hipMemset(d, 0, n * sizeof(int)) != hipSuccess)
+    return 2;
+  mod_kernel<<<dim3(1), dim3(64)>>>(d, n);
+  if (hipDeviceSynchronize() != hipSuccess)
+    return 3;
+  (void)hipFree(d);
+  return 0;
+}
+#endif

test/smoke-dev/OutputHipDevCov/run_and_check.sh

@@ -0,0 +1,84 @@
+#!/bin/bash
+# HipDevCov check driver.
+#
+# Args (passed from the Makefile):
+#   $1 = AOMP        (LLVM dir containing bin/clang, llvm-profdata, llvm-objdump)
+#   $2 = GPU arch    (e.g. gfx90a, may include :features)
+#   $3 = FILECHECK   (path to FileCheck)
+#   $4 = ROCM_PATH   (HIP/ROCm install for --rocm-path and libamdhip64)
+#
+# Returns 0 only if BOTH the host-shadow kernel and the module-only kernel are
+# present in the merged profile (the latter requires the HSA drain).
+
+set -u
+AOMP="$1"; ARCH="$2"; ROCM="$4"
+CLANG="$AOMP/bin/clang"
+PROFDATA="$AOMP/bin/llvm-profdata"
+OBJDUMP="$AOMP/bin/llvm-objdump"
+COV="$AOMP/bin/llvm-cov"
+
+here="$(cd "$(dirname "$0")" && pwd)"
+cd "$here" || exit 1
+
+set -x
+
+# Capability gate: skip cleanly on toolchains that do not ship the device
+# profile runtime / HSA drain yet (so this does not red-fail CI before the
+# feature lands). A real, drain-capable toolchain has the amdgcn device profile
+# RT and the host drain symbol.
+resdir="$("$CLANG" -print-resource-dir 2>/dev/null)"
+devrt="$resdir/lib/amdgcn-amd-amdhsa/libclang_rt.profile.a"
+if [ ! -f "$devrt" ]; then
+  echo "SKIP HipDevCov: no device profile runtime at $devrt"
+  exit 0
+fi
+if ! ls "$resdir"/lib/*/libclang_rt.profile_rocm*.a >/dev/null 2>&1; then
+  echo "SKIP HipDevCov: no host profile_rocm runtime (HSA drain) in toolchain"
+  exit 0
+fi
+
+set -e
+rm -f ./*.profraw mod.co merged.profdata builder builder.*.host-* \
+      builder.*.hip-amdgcn-amd-amdhsa--* 2>/dev/null || true
+
+# 1. Build mod.hip as a full executable so the device link gets the device
+#    profile RT, then extract its device code object as the loadable mod.co.
+"$CLANG" -x hip --offload-arch="$ARCH" -fno-gpu-rdc -DBUILD_MODULE_EXE \
+  -fprofile-instr-generate -fcoverage-mapping --rocm-path="$ROCM" \
+  mod.hip -o builder -L"$ROCM/lib" -lamdhip64 -Wl,-rpath,"$ROCM/lib"
+"$OBJDUMP" --offloading builder >/dev/null 2>&1 || true
+shopt -s nullglob
+extracted=(builder*.hip-amdgcn-amd-amdhsa--*gfx*)
+if [ ${#extracted[@]} -eq 0 ]; then
+  echo "FAIL HipDevCov: could not extract device code object from builder"
+  exit 1
+fi
+cp "${extracted[0]}" mod.co
+
+# 2. main is built by the smoke harness (TESTNAME=main). Build it here too if it
+#    is missing, so the script also works when run standalone.
+if [ ! -x ./main ]; then
+  "$CLANG" -x hip --offload-arch="$ARCH" -fno-gpu-rdc \
+    -fprofile-instr-generate -fcoverage-mapping --rocm-path="$ROCM" \
+    main.hip -o main -L"$ROCM/lib" -lamdhip64 -Wl,-rpath,"$ROCM/lib"
+fi
+
+# 3. Run. Device .profraw files are written to CWD (arch-prefixed for the
+#    host-shadow drain, arch.hsa<N>-prefixed for the HSA drain); the host one
+#    goes to LLVM_PROFILE_FILE.
+rm -f ./*.profraw
+LLVM_PROFILE_FILE="$here/host.profraw" \
+  LD_LIBRARY_PATH="$ROCM/lib:${LD_LIBRARY_PATH:-}" \
+  ./main
+
+# 4. Merge host + all device profraws and assert both kernels are present.
+"$PROFDATA" merge -sparse -o merged.profdata ./*.profraw
+"$PROFDATA" show --all-functions merged.profdata 
+
+# 5. Sanity (non-fatal): llvm-cov can consume the merged device+host profile.
+#    main's covmap only describes host_kernel/main, so llvm-cov warns about the
+#    module-only function; that is expected, hence non-fatal.
+"$COV" report ./main -instr-profile=merged.profdata 
+"$COV" show ./main -instr-profile=merged.profdata
+
+echo "HipDevCov PASSED"