diff --git a/pkgs/development/rocm-modules/clr/default.nix b/pkgs/development/rocm-modules/clr/default.nix index b70d9ec68c12..8b76a12d9351 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 = [ @@ -294,6 +295,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 ]; + }; +} 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 ]; + }; +} diff --git a/pkgs/development/rocm-modules/miopen/default.nix b/pkgs/development/rocm-modules/miopen/default.nix index 90845f67582e..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" + '' @@ -274,6 +277,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 ]; + }; +} 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