rocmPackages.miopen: fix bnorm and dropout kernels (#507733)

pkgs/development/rocm-modules/miopen/default.nix

@@ -23,6 +23,7 @@
   half,
   boost,
   sqlite,
+  symlinkJoin,
   bzip2,
   lbzip2,
   nlohmann_json,
@@ -82,6 +83,15 @@ let
     )
   );
 
+  # for hiprtcCompileProgram (dropout kernels require rocrand in -I at runtime)
+  hiprtcCompileRocmPath = symlinkJoin {
+    name = "miopen-hiprtc-compile-rocm-path";
+    paths = [
+      clr
+      rocrand
+    ];
+  };
+
   # Kernel databases moved from Git LFS to DVC (anonymous s3 bucket s3://therock-dvc/rocm-libraries)
   fetchKdb =
     name:
@@ -203,7 +213,6 @@ stdenv.mkDerivation (finalAttrs: {
     "-DAMDGPU_TARGETS=${lib.concatStringsSep ";" supportedTargets}"
     "-DGPU_TARGETS=${lib.concatStringsSep ";" supportedTargets}"
     "-DGPU_ARCHS=${lib.concatStringsSep ";" supportedTargets}"
-    "-DMIOPEN_USE_SQLITE_PERFDB=ON"
     "-DCMAKE_VERBOSE_MAKEFILE=ON"
     "-DCMAKE_MODULE_PATH=${clr}/hip/cmake"
     "-DCMAKE_BUILD_TYPE=Release"
@@ -244,7 +253,7 @@ stdenv.mkDerivation (finalAttrs: {
     patchShebangs test src/composable_kernel fin utils install_deps.cmake
 
     substituteInPlace src/comgr.cpp \
-      --replace-fail '"/opt/rocm"' '"${clr}"'
+      --replace-fail '"/opt/rocm"' '"${hiprtcCompileRocmPath}"'
   ''
   + linkKDBsTo "src/kernels"
   + ''
@@ -285,12 +294,22 @@ stdenv.mkDerivation (finalAttrs: {
   requiredSystemFeatures = [ "big-parallel" ];
 
   passthru.impureTests = {
-    # bash $(nix-build -A rocmPackages.miopen.passthru.impureTests.conv)
+    # bash $(nix-build -A rocmPackages.miopen.passthru.impureTests.conv) etc
+    bnorm = callPackage ./test-runtime-compilation.nix {
+      miopen = finalAttrs.finalPackage;
+      name = "bnorm";
+      testScript = "MIOpenDriver bnorm -n 16 -c 16 -H 512 -W 512 -m 1 -F 1 -s 1 -r 1";
+    };
     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";
     };
+    dropout = callPackage ./test-runtime-compilation.nix {
+      miopen = finalAttrs.finalPackage;
+      name = "dropout";
+      testScript = "MIOpenDriver dropout -d 64,32,14,14";
+    };
     pool = callPackage ./test-runtime-compilation.nix {
       miopen = finalAttrs.finalPackage;
       name = "pool";