From b9cfeaffa4671dd8481340ab62b3ebdc42dc6e8a Mon Sep 17 00:00:00 2001 From: Luna Nova Date: Mon, 9 Mar 2026 17:13:53 -0700 Subject: [PATCH 1/6] rocmPackages.miopen: add impureTests that verify conv/pool functionality --- .../rocm-modules/miopen/default.nix | 13 +++++++ .../miopen/test-runtime-compilation.nix | 37 +++++++++++++++++++ 2 files changed, 50 insertions(+) create mode 100644 pkgs/development/rocm-modules/miopen/test-runtime-compilation.nix diff --git a/pkgs/development/rocm-modules/miopen/default.nix b/pkgs/development/rocm-modules/miopen/default.nix index 90845f67582e..56e6cbb66c48 100644 --- a/pkgs/development/rocm-modules/miopen/default.nix +++ b/pkgs/development/rocm-modules/miopen/default.nix @@ -274,6 +274,19 @@ stdenv.mkDerivation (finalAttrs: { requiredSystemFeatures = [ "big-parallel" ]; + passthru.impureTests = { + # bash $(nix-build -A rocmPackages.miopen.passthru.impureTests.conv) + conv = callPackage ./test-runtime-compilation.nix { + miopen = finalAttrs.finalPackage; + name = "conv"; + testScript = "MIOpenDriver conv -n 1 -c 1 -H 4 -W 4 -k 1 -y 3 -x 3 -p 0 -q 0 -V 0"; + }; + pool = callPackage ./test-runtime-compilation.nix { + miopen = finalAttrs.finalPackage; + name = "pool"; + testScript = "MIOpenDriver pool -W 1x1x4x4 -y 2 -x 2 -p 0 -q 0 -F 1 -V 0"; + }; + }; passthru.tests = { # Ensure all .tn.model files can be loaded by whatever version of frugally-deep we have # This is otherwise hard to verify as MIOpen will only use these models on specific, diff --git a/pkgs/development/rocm-modules/miopen/test-runtime-compilation.nix b/pkgs/development/rocm-modules/miopen/test-runtime-compilation.nix new file mode 100644 index 000000000000..c781d3646251 --- /dev/null +++ b/pkgs/development/rocm-modules/miopen/test-runtime-compilation.nix @@ -0,0 +1,37 @@ +{ + lib, + makeImpureTest, + writableTmpDirAsHomeHook, + miopen, + clr, + rocm-smi, + name, + testScript, +}: + +makeImpureTest { + inherit name; + testedPackage = "rocmPackages.miopen"; + + sandboxPaths = [ + "/sys" + "/dev/dri" + "/dev/kfd" + ]; + + nativeBuildInputs = [ + writableTmpDirAsHomeHook + miopen + clr + rocm-smi + ]; + + testScript = '' + rocm-smi + ${testScript} + ''; + + meta = { + teams = [ lib.teams.rocm ]; + }; +} From e7cf3ae7169859effd03fb7e575d324d2b2932c1 Mon Sep 17 00:00:00 2001 From: Luna Nova Date: Mon, 9 Mar 2026 17:01:15 -0700 Subject: [PATCH 2/6] rocmPackages.miopen: fix hardcoded /opt/rocm path in comgr.cpp --- pkgs/development/rocm-modules/miopen/default.nix | 3 +++ 1 file changed, 3 insertions(+) diff --git a/pkgs/development/rocm-modules/miopen/default.nix b/pkgs/development/rocm-modules/miopen/default.nix index 56e6cbb66c48..8e173f4ff40e 100644 --- a/pkgs/development/rocm-modules/miopen/default.nix +++ b/pkgs/development/rocm-modules/miopen/default.nix @@ -235,6 +235,9 @@ stdenv.mkDerivation (finalAttrs: { return()' patchShebangs test src/composable_kernel fin utils install_deps.cmake + + substituteInPlace src/comgr.cpp \ + --replace-fail '"/opt/rocm"' '"${clr}"' '' + linkKDBsTo "src/kernels" + '' From 235ea64d880ba27fd940d66ef6c7ef6641a35bd7 Mon Sep 17 00:00:00 2001 From: Luna Nova Date: Mon, 9 Mar 2026 17:13:35 -0700 Subject: [PATCH 3/6] rocmPackages.clr: apply own rpath workaround to all libs Same issue applies to libhiprtc.so, seems safer to apply to all since it's a recurring pattern in clr. --- pkgs/development/rocm-modules/clr/default.nix | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) diff --git a/pkgs/development/rocm-modules/clr/default.nix b/pkgs/development/rocm-modules/clr/default.nix index 71451c81e443..9e5b7abbf099 100644 --- a/pkgs/development/rocm-modules/clr/default.nix +++ b/pkgs/development/rocm-modules/clr/default.nix @@ -207,12 +207,13 @@ stdenv.mkDerivation (finalAttrs: { ln -s ${hipClang} $out/llvm ''; - # libamdhip64.so dlopens its own bare name for hipGetProcAddress symbol resolution. - # Add its own directory to its RPATH so it can find itself + # libamdhip64.so dlopens its own bare name for hipGetProcAddress symbol resolution, + # same pattern with libhiprtc.so, so add own lib directory to all .so's + # RPATHs so they can find themselves and neighbouring libs # Must be in postFixup so it runs after patchelf --shrink-rpath which removes # the apparently useless rpath postFixup = '' - patchelf --add-rpath "$out/lib" "$out/lib/libamdhip64.so" + patchelf --add-rpath "$out/lib" "$out"/lib/*.so ''; disallowedRequisites = [ From 6df5d367483ec56869ca847568517cb0ffe82d2a Mon Sep 17 00:00:00 2001 From: Luna Nova Date: Tue, 10 Mar 2026 08:38:50 -0700 Subject: [PATCH 4/6] rocmPackages.clr: add impureTest for C++ stdlib availability --- pkgs/development/rocm-modules/clr/default.nix | 4 ++ .../clr/test-hiprtc-type-traits.cpp | 68 +++++++++++++++++++ .../clr/test-hiprtc-type-traits.nix | 57 ++++++++++++++++ 3 files changed, 129 insertions(+) create mode 100644 pkgs/development/rocm-modules/clr/test-hiprtc-type-traits.cpp create mode 100644 pkgs/development/rocm-modules/clr/test-hiprtc-type-traits.nix diff --git a/pkgs/development/rocm-modules/clr/default.nix b/pkgs/development/rocm-modules/clr/default.nix index 9e5b7abbf099..d7202b6190ef 100644 --- a/pkgs/development/rocm-modules/clr/default.nix +++ b/pkgs/development/rocm-modules/clr/default.nix @@ -297,6 +297,10 @@ stdenv.mkDerivation (finalAttrs: { "amdgcnspirv" ]; }; + hiprtc-type-traits = callPackage ./test-hiprtc-type-traits.nix { + clr = finalAttrs.finalPackage; + inherit rocm-smi; + }; }; selectGpuTargets = diff --git a/pkgs/development/rocm-modules/clr/test-hiprtc-type-traits.cpp b/pkgs/development/rocm-modules/clr/test-hiprtc-type-traits.cpp new file mode 100644 index 000000000000..39a99f39eae1 --- /dev/null +++ b/pkgs/development/rocm-modules/clr/test-hiprtc-type-traits.cpp @@ -0,0 +1,68 @@ +#include +#include +#include +#include + +#define CHECK_HIP(expr) do { \ + if ((expr) != hipSuccess) { \ + std::cerr << #expr << " failed" << std::endl; \ + return 1; \ + } \ +} while(0) + +#define CHECK_HIPRTC(expr) do { \ + hiprtcResult _res = (expr); \ + if (_res != HIPRTC_SUCCESS) { \ + std::cerr << #expr << " failed: " << hiprtcGetErrorString(_res) << std::endl; \ + hiprtcGetProgramLogSize(prog, &log_size); \ + if (log_size > 0) { \ + std::string log(log_size, '\0'); \ + hiprtcGetProgramLog(prog, log.data()); \ + std::cerr << "Compile log:\n" << log << std::endl; \ + } \ + return 1; \ + } \ +} while(0) + +static const char* kernelSource = R"( + #include + + extern "C" __global__ void test_kernel(int* out) { + static_assert(std::is_same::type>::value, + "type_traits not working"); + out[0] = 5; + } +)"; + +int main() { + hiprtcProgram prog; + size_t log_size = 0; + CHECK_HIPRTC(hiprtcCreateProgram(&prog, kernelSource, "test.hip", 0, nullptr, nullptr)); + CHECK_HIPRTC(hiprtcCompileProgram(prog, 0, nullptr)); + + size_t code_size; + CHECK_HIPRTC(hiprtcGetCodeSize(prog, &code_size)); + std::string code(code_size, '\0'); + CHECK_HIPRTC(hiprtcGetCode(prog, code.data())); + hiprtcDestroyProgram(&prog); + + hipModule_t module; + hipFunction_t kernel; + CHECK_HIP(hipModuleLoadData(&module, code.data())); + CHECK_HIP(hipModuleGetFunction(&kernel, module, "test_kernel")); + + int* d_out; + int h_out = 0; + CHECK_HIP(hipMalloc(&d_out, sizeof(int))); + void* args[] = { &d_out }; + CHECK_HIP(hipModuleLaunchKernel(kernel, 1, 1, 1, 1, 1, 1, 0, nullptr, args, nullptr)); + CHECK_HIP(hipMemcpy(&h_out, d_out, sizeof(int), hipMemcpyDeviceToHost)); + + if (h_out != 5) { + std::cerr << "Kernel output mismatch: expected 5, got " << h_out << std::endl; + return 1; + } + + std::cout << "HIPRTC type_traits test passed (output=" << h_out << ")" << std::endl; + return 0; +} diff --git a/pkgs/development/rocm-modules/clr/test-hiprtc-type-traits.nix b/pkgs/development/rocm-modules/clr/test-hiprtc-type-traits.nix new file mode 100644 index 000000000000..f0f3628fcdaa --- /dev/null +++ b/pkgs/development/rocm-modules/clr/test-hiprtc-type-traits.nix @@ -0,0 +1,57 @@ +{ + lib, + stdenv, + makeImpureTest, + clr, + rocm-smi, +}: +# minimal hiprtc test that compiles a kernel using at runtime +# mirrors an migraphx workload, better test/iteration UX to be able to confirm +# with just a build up to clr +let + hiprtc-test = stdenv.mkDerivation { + pname = "hiprtc-type-traits-test"; + version = "0"; + + dontUnpack = true; + + nativeBuildInputs = [ clr ]; + + buildPhase = '' + runHook preBuild + hipcc -o hiprtc-test ${./test-hiprtc-type-traits.cpp} -lhiprtc + runHook postBuild + ''; + + installPhase = '' + runHook preInstall + mkdir -p $out/bin + cp hiprtc-test $out/bin/ + runHook postInstall + ''; + }; +in +makeImpureTest { + name = "hiprtc-type-traits"; + testedPackage = "rocmPackages.clr"; + + sandboxPaths = [ + "/sys" + "/dev/dri" + "/dev/kfd" + ]; + + nativeBuildInputs = [ + hiprtc-test + rocm-smi + ]; + + testScript = '' + rocm-smi + hiprtc-test + ''; + + meta = { + teams = [ lib.teams.rocm ]; + }; +} From 4781916d7059582a5d60f102ba3f87fbbe1b5c53 Mon Sep 17 00:00:00 2001 From: Luna Nova Date: Tue, 10 Mar 2026 08:39:46 -0700 Subject: [PATCH 5/6] rocmPackages.rocm-comgr: bake fallback LLVM PATH to fix C++ stdlib search --- .../rocm-modules/rocm-comgr/default.nix | 17 +++++++++++++---- 1 file changed, 13 insertions(+), 4 deletions(-) diff --git a/pkgs/development/rocm-modules/rocm-comgr/default.nix b/pkgs/development/rocm-modules/rocm-comgr/default.nix index 0a49cb0a24fa..e7663a1f7130 100644 --- a/pkgs/development/rocm-modules/rocm-comgr/default.nix +++ b/pkgs/development/rocm-modules/rocm-comgr/default.nix @@ -45,10 +45,19 @@ stdenv.mkDerivation (finalAttrs: { ./fix-ccob-compat-output-filename.patch ]; - postPatch = '' - substituteInPlace cmake/opencl_header.cmake \ - --replace-fail "\''${CLANG_CMAKE_DIR}/../../../" "${llvm.clang-unwrapped.lib}" - ''; + postPatch = + # Fix relative path assumption for libllvm + '' + substituteInPlace cmake/opencl_header.cmake \ + --replace-fail "\''${CLANG_CMAKE_DIR}/../../../" "${llvm.clang-unwrapped.lib}" + '' + # Bake LLVM root for cfg/includes or HIPRTC can't find C++ stdlib headers (e.g. ). + + '' + substituteInPlace src/comgr-env.cpp \ + --replace-fail \ + 'return EnvLLVMPath;' \ + 'return EnvLLVMPath ? EnvLLVMPath : "${llvm.rocm-toolchain}";' + ''; nativeBuildInputs = [ cmake From c1c21ed52282e642a16ad0154039bdc77a31dbd9 Mon Sep 17 00:00:00 2001 From: Luna Nova Date: Tue, 10 Mar 2026 08:59:36 -0700 Subject: [PATCH 6/6] rocmPackages.migraphx: add migraphx-driver impureTest using small resnet --- .../rocm-modules/migraphx/default.nix | 7 +++ .../migraphx/test-migraphx-driver.nix | 47 +++++++++++++++++++ 2 files changed, 54 insertions(+) create mode 100644 pkgs/development/rocm-modules/migraphx/test-migraphx-driver.nix diff --git a/pkgs/development/rocm-modules/migraphx/default.nix b/pkgs/development/rocm-modules/migraphx/default.nix index 0e6f84ca1487..8b602bcdf362 100644 --- a/pkgs/development/rocm-modules/migraphx/default.nix +++ b/pkgs/development/rocm-modules/migraphx/default.nix @@ -1,6 +1,7 @@ { lib, stdenv, + callPackage, fetchFromGitHub, rocmUpdateScript, pkg-config, @@ -183,6 +184,12 @@ stdenv.mkDerivation (finalAttrs: { patchelf $test/bin/test_* --shrink-rpath --allowed-rpath-prefixes "$NIX_STORE" ''; + passthru.impureTests = { + # NIXPKGS_ALLOW_UNFREE=1 bash $(nix-build -A rocmPackages.migraphx.impureTests.migraphx-driver) + migraphx-driver = callPackage ./test-migraphx-driver.nix { + migraphx = finalAttrs.finalPackage; + }; + }; passthru.updateScript = rocmUpdateScript { name = finalAttrs.pname; inherit (finalAttrs.src) owner; diff --git a/pkgs/development/rocm-modules/migraphx/test-migraphx-driver.nix b/pkgs/development/rocm-modules/migraphx/test-migraphx-driver.nix new file mode 100644 index 000000000000..c34c5a8a71d4 --- /dev/null +++ b/pkgs/development/rocm-modules/migraphx/test-migraphx-driver.nix @@ -0,0 +1,47 @@ +{ + lib, + fetchurl, + makeImpureTest, + writableTmpDirAsHomeHook, + migraphx, + clr, + rocm-smi, +}: + +# Verify that a ≈50MiB resnet onnx can run with migraphx +let + resnet18 = fetchurl { + url = "https://huggingface.co/onnxmodelzoo/resnet18_Opset18_timm/resolve/main/resnet18_Opset18_timm.onnx"; + hash = "sha256-u2Io20n72qoA9atRsFIWb0zHF1WdJYgHQdMWfJhJGHA="; + meta.license = lib.licenses.unfree; + }; +in +makeImpureTest { + name = "migraphx-driver"; + testedPackage = "rocmPackages.migraphx"; + + sandboxPaths = [ + "/sys" + "/dev/dri" + "/dev/kfd" + ]; + + nativeBuildInputs = [ + writableTmpDirAsHomeHook + migraphx + clr + rocm-smi + ]; + + # FIXME(@LunNova): tol values are set too high - was seeing high divergence on iGPU + # want this test to be useful for verifying workloads run at all + # and will investigate what's broken for accuracy + testScript = '' + rocm-smi + migraphx-driver verify -O --rms-tol 0.03 --atol 1.0 --rtol 0.01 ${resnet18} + ''; + + meta = { + teams = [ lib.teams.rocm ]; + }; +}