[smoke-dev] Add HipDevCov and OmpDevCov device-coverage drain tests

Add two smoke-dev tests that validate device-side PGO/coverage drainage via
the in-tree HSA-introspection drain (upstream LLVM #203056), independent of the
HIP host-shadow drain.

HipDevCov: builds a device-only kernel as an executable so the device link
pulls in the amdgcn profile runtime, extracts its code object (mod.co) via
llvm-objdump --offloading, then hipModuleLoad()s it from a program where that
kernel has no host shadow. The kernel's counters can only reach the merged
profile through the HSA drain; the test asserts both the host-shadow kernel and
the module-only kernel are present.

OmpDevCov: OpenMP target offload has no host-shadow drain, so a
-fprofile-generate -> llvm-profdata merge -> -fprofile-use round trip that
consumes a device function's counters exercises the HSA drain exclusively.

Both honor $AOMP/$AOMP_GPU, use the smoke-dev Makefile harness, and skip
cleanly on toolchains that do not yet ship the device profile runtime / drain.

Co-authored-by: Cursor <cursoragent@cursor.com>

Author: lfmeadow

test/smoke-dev/HipDevCov/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/HipDevCov/check.txt

@@ -0,0 +1,6 @@
+Both kernels must appear in the merged host+device profile. host_kernel is the
+host-shadow path; mod_kernel (from the hipModuleLoad'd code object, no host
+shadow) can only be present if the HSA-introspection drain collected it.
+
+CHECK-DAG: host_kernel
+CHECK-DAG: mod_kernel

test/smoke-dev/HipDevCov/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/HipDevCov/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/HipDevCov/run_and_check.sh

@@ -0,0 +1,80 @@
+#!/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"; FILECHECK="$3"; 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
+
+# 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
+co="$(ls builder*.hip-amdgcn-amd-amdhsa--*gfx* 2>/dev/null | head -1)"
+if [ -z "$co" ]; then
+  echo "FAIL HipDevCov: could not extract device code object from builder"
+  exit 1
+fi
+cp "$co" 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 | "$FILECHECK" check.txt
+
+# 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 >/dev/null 2>&1 || true
+
+echo "HipDevCov PASSED"

test/smoke-dev/OmpDevCov/Makefile

@@ -0,0 +1,26 @@
+include ../../Makefile.defs
+
+# OmpDevCov: device-side PGO round trip for OpenMP target offload. The harness
+# builds omp_pgo (TESTNAME) with -fprofile-generate; run_and_check.sh runs it,
+# verifies the device counters were drained via the HSA pass, merges, and feeds
+# them back with -fprofile-use.
+
+TESTNAME     = omp_pgo
+TESTSRC_MAIN = omp_pgo.c
+TESTSRC_AUX  =
+TESTSRC_ALL  = $(TESTSRC_MAIN) $(TESTSRC_AUX)
+
+# Extra budget: run_and_check.sh also performs the -fprofile-use rebuild + run.
+TIMEOUT      = 300s
+
+CFLAGS       = -fopenmp $(TARGET) -fprofile-generate
+LINK_FLAGS   = -Wl,-rpath,$(AOMP)/lib
+
+CC           = $(AOMP)/bin/clang $(VERBOSE)
+
+RUNCMD       = ./run_and_check.sh "$(AOMP)" "$(GPU_W_FEATURES)" "$(FILECHECK)"
+
+include ../Makefile.rules
+
+clean::
+	rm -f omp_pgo omp_pgo_use *.profraw merged.profdata

test/smoke-dev/OmpDevCov/check.txt

@@ -0,0 +1,6 @@
+The merged profile must contain device-side counters that could only have been
+collected via the HSA-introspection drain (OpenMP offload has no host-shadow
+drain). The device offload entry and the device function classify both appear.
+
+CHECK-DAG: __omp_offloading_
+CHECK-DAG: classify

test/smoke-dev/OmpDevCov/omp_pgo.c

@@ -0,0 +1,45 @@
+/* OmpDevCov: device-side PGO round trip for OpenMP target offload.
+ *
+ * OpenMP offload has no HIP-style host-shadow drain for its device code, so the
+ * only way device counters reach the .profraw is the in-tree HSA-introspection
+ * drain. A successful -fprofile-generate -> llvm-profdata merge -> -fprofile-use
+ * cycle that consumes a device function's counters therefore exercises the HSA
+ * drain exclusively.
+ *
+ * classify() runs on the device and has a data-dependent branch so there is a
+ * non-trivial counter to instrument, merge, and consume.
+ */
+
+#include <stdio.h>
+#include <stdlib.h>
+
+#define N 4096
+
+#pragma omp declare target
+static int classify(int x) {
+  if ((x & 1) == 0)
+    return x * 2; /* even path */
+  else
+    return x + 1; /* odd path  */
+}
+#pragma omp end declare target
+
+int main(void) {
+  int *a = (int *)malloc(N * sizeof(int));
+  if (!a)
+    return 1;
+  for (int i = 0; i < N; i++)
+    a[i] = i;
+
+#pragma omp target teams distribute parallel for map(tofrom : a[0:N])
+  for (int i = 0; i < N; i++)
+    a[i] = classify(a[i]);
+
+  long sum = 0;
+  for (int i = 0; i < N; i++)
+    sum += a[i];
+
+  printf("sum=%ld\n", sum);
+  free(a);
+  return 0;
+}

test/smoke-dev/OmpDevCov/run_and_check.sh

@@ -0,0 +1,76 @@
+#!/bin/bash
+# OmpDevCov check driver: device-side PGO round trip for OpenMP offload.
+#
+# Args (passed from the Makefile):
+#   $1 = AOMP       (LLVM dir containing bin/clang, llvm-profdata)
+#   $2 = GPU arch   (e.g. gfx90a, may include :features)
+#   $3 = FILECHECK  (path to FileCheck)
+#
+# omp_pgo (the -fprofile-generate binary) is built by the harness. Here we:
+#   1. run it -> host default.profraw + device amdgcn-amd-amdhsa.*.profraw
+#      (the device file is produced ONLY by the HSA drain),
+#   2. assert a device profraw was produced and merge,
+#   3. FileCheck the merged profile for the device offload function,
+#   4. rebuild with -fprofile-use and assert it consumes cleanly (no profile
+#      mismatch / out-of-date diagnostics) and runs.
+
+set -u
+AOMP="$1"; ARCH="$2"; FILECHECK="$3"
+CLANG="$AOMP/bin/clang"
+PROFDATA="$AOMP/bin/llvm-profdata"
+
+here="$(cd "$(dirname "$0")" && pwd)"
+cd "$here" || exit 1
+
+# Capability gate: skip cleanly if the toolchain has no amdgcn device profile
+# runtime (i.e. no device PGO / HSA drain yet), so CI is not red before the
+# feature lands.
+resdir="$("$CLANG" -print-resource-dir 2>/dev/null)"
+if [ ! -f "$resdir/lib/amdgcn-amd-amdhsa/libclang_rt.profile.a" ]; then
+  echo "SKIP OmpDevCov: no amdgcn device profile runtime in toolchain"
+  exit 0
+fi
+
+set -e
+rm -f ./*.profraw merged.profdata omp_pgo_use 2>/dev/null || true
+
+run_env() { LD_LIBRARY_PATH="$AOMP/lib:${LD_LIBRARY_PATH:-}" "$@"; }
+
+# 1. Run the generate binary (built by the harness as ./omp_pgo). Build it here
+#    too if missing, so the script also works standalone.
+if [ ! -x ./omp_pgo ]; then
+  "$CLANG" -fopenmp --offload-arch="$ARCH" -fprofile-generate \
+    omp_pgo.c -o omp_pgo -Wl,-rpath,"$AOMP/lib"
+fi
+run_env ./omp_pgo
+
+# 2. The HSA drain must have produced at least one device profraw.
+shopt -s nullglob
+dev=(amdgcn-amd-amdhsa*.profraw)
+if [ ${#dev[@]} -eq 0 ]; then
+  echo "FAIL OmpDevCov: no device (amdgcn-amd-amdhsa) profraw -- HSA drain did not fire"
+  ls -1 ./*.profraw 2>/dev/null || true
+  exit 1
+fi
+echo "OmpDevCov: device profraw(s): ${dev[*]}"
+
+"$PROFDATA" merge -o merged.profdata ./*.profraw
+
+# 3. The merged profile must carry the device offload function counters.
+"$PROFDATA" show --all-functions merged.profdata | "$FILECHECK" check.txt
+
+# 4. -fprofile-use round trip: must consume the profile without a mismatch /
+#    out-of-date diagnostic, then run.
+use_log="$(mktemp)"
+"$CLANG" -fopenmp --offload-arch="$ARCH" -fprofile-use=merged.profdata \
+  omp_pgo.c -o omp_pgo_use -Wl,-rpath,"$AOMP/lib" 2> "$use_log" || {
+    echo "FAIL OmpDevCov: -fprofile-use build failed"; cat "$use_log"; rm -f "$use_log"; exit 1; }
+if grep -Eiq "out of date|profile data may be out of date|no profile data available|mismatch" "$use_log"; then
+  echo "FAIL OmpDevCov: -fprofile-use reported stale/mismatched profile"
+  cat "$use_log"; rm -f "$use_log"; exit 1
+fi
+rm -f "$use_log"
+
+run_env ./omp_pgo_use
+
+echo "OmpDevCov PASSED"