[HIP] Fix rocm detection

[yaxunl]

Summary

Do not detect device library by default in rocm detector. Only detect device library in HIP toolchain.

Detect rocm path by version file in host toolchains.

Also added detecting rocm version and printing rocm installation path and version with -v.

Fixed include path and device library detection for ROCm 3.5.

[arsenm]

clang/lib/Driver/ToolChains/AMDGPU.cpp

196

I don't think the detection should fail if this is missing

[yaxunl]

clang/lib/Driver/ToolChains/AMDGPU.cpp

196

why? this file is unique to HIP installation. If it is missing, the installation is broken.

[arsenm]

clang/lib/Driver/ToolChains/AMDGPU.cpp

196

Because I should be able to use this without a complete hip installation. Without a specified version, it should assume the most modern layout. This will for example break pointing --rocm-path at the device library build directory for library tests

[arsenm]

clang/lib/Driver/ToolChains/AMDGPU.cpp

196

I also don't see what value checking the version really provides; it may be informative to print it, but I don't think it's useful to derive information from it

[yaxunl]

clang/lib/Driver/ToolChains/AMDGPU.cpp

196

what is the directory structure of your most modern layout?

[arsenm]

clang/lib/Driver/ToolChains/AMDGPU.cpp

196

/opt/rocm/amdgcn/bitcode/foo.bc

The plan is to remove this and rely on symlinks in the resource directory though

[tra]

clang/lib/Driver/ToolChains/AMDGPU.cpp

196

I also don't see what value checking the version really provides; it may be informative to print it, but I don't think it's useful to derive information from it

In CUDA it's used to detect which back-end features to enable (they depend on what's supported by ptxas supplied by that CUDA version). I don't think that would be relevant to AMDGPU as it does not need external dependencies to generate GPU code.

It may be useful for filesystem layout detection. E.g. where to find bitcode files and what names to expect. This part seems to be relevant for ROCm, but I'm not sure if the layout is closely tied to the version.

[yaxunl]

clang/lib/Driver/ToolChains/AMDGPU.cpp

196

We are required to support previous ROCm releases for certain time range. To do that we need to detect HIP version and enable/disable certain HIP features based on HIP version when necessary.

Therefore if we have a new directory structure for ROCm, that directory structure should contain a HIP version file so that we can detect HIP version.

[yaxunl]

clang/lib/Driver/ToolChains/AMDGPU.cpp

196

Currently clang includes some wrapper headers by default, which does not work with ROCm 3.5 since those device functions are defined in HIP headers of ROCm 3.5. To support ROCm 3.5, we have to disable including those wrapper headers. This is one example why we need detect HIP version.

196

I think we need to separate HIP runtime detection and device library detection and also separate hasValidHIPRuntime and hasValidDeviceLibrary. ROCm toolchain only need to make sure hasValidDeviceLibrary but HIP toolchain need both hasValidDeviceLibrary and hasValidHIPRuntime.

[arsenm]

clang/lib/Driver/ToolChains/AMDGPU.cpp

196

Regardless of whether there's a version file or if does anything, I think the absence of one implies do the most modern thing

[tra]

clang/lib/Driver/ToolChains/AMDGPU.cpp

196

"The most modern thing" is an ever moving target. If the failure modes keep changing it will create a source of new problems every time something changes. Probably not a big issue in this case, but with (eventually) wide enough user base there will be some.

I assume that AMD does not have much of legacy user base yet for compilations with clang, so defaulting to a recent version may be sensible (unlike CUDA where clang had to deal with a lot of existing CUDA users using old CUDA versions). That said, I'd recommend pinning it to a specific version, so the expectations remain stable. Then we can document the failure/fallback in a way users can deal with -- if no version file is found, clang assumes version x.y and expects to find files X, Y and Z in directory A, B, C. You can specify the locations manually with --something-something-X=A or specify the version with --hip-version=3.7.

Considering that clang is being used from various derived tools, you may not always have the luxury of having the whole ROCm installation around and need to be able to override the expectations via command line flags.

If we have a way to override the version, then it does not matter all that much which version we pick as a fallback. If the guess is wrong, we can always correct it with a flag.

[arsenm]

clang/lib/Driver/ToolChains/AMDGPU.cpp

196

Unlike CUDA we're not really chasing some 3rd party thing that periodically drops unpredictable changes. We know what the most modern thing is on tip of tree

[tra]

clang/lib/Driver/ToolChains/AMDGPU.cpp

196

Unlike CUDA we're not really chasing some 3rd party thing that periodically drops unpredictable changes. We know what the most modern thing is on tip of tree

We're in agreement on this part. If lack of the version file is non-fatal, I'm fine with whatever version AMD picks as the default.

I'm just pointing out that you'll likely need a way to keep compiler working even if the filesystem does not have everything you expect for the full end-to-end compilation, be it the version file or bitcode. Reasonable default is fine. Overrides can be implemented later, if necessary.

[yaxunl]

Separate detecting ROCm path and device library path.

Assume 3.5 as default HIP version. If -rocm-path is specified, do not require version file to exist.

Added --hip-version to override detected HIP version.

Fixed default option for -std and -fhip-use-new-launch-api.

[tra]

LGTM in principle. Few style comments.

clang/lib/Driver/ToolChains/AMDGPU.cpp

161

A comment about expected version format would be helpful here.

Perhaps it could be simplified by splitting the string on "." into an array in one call instead of doing it one element at a time.

244–257

It would be nice to collect the list of directories to check into an array and then loop over it -- the same way you already do for the InstallPath candidates. Makes it easier to add/change search logic.

E.g. in the end we would get something like this:

for (path: {InstallPath, make_path(InstallPath, "lib"), make_path(InstallPath, "lib", "bitcode"), ... }) {
   if (CheckDeviceLib(path)) {
      // record path as the correct location
     return;
   }
}

248–249

Creation of the path from base+suffixes seems to be a common pattern to be extracted into a lambda. I'm surprised it's not provided by LLVM itself.

248–250

CheckDeviceLib should probably take the path to check as an argument and so should scanLibDevicePath.

299

Nit: IsRocm35 is a bit misleading as it would be true for versions other than 3.5.
Rename it to something closer matching the intent? IsRocm35OrOlder ot, perhaps, hipUsesRuntimeWrapper.

333

You could use CC1Args.append({"foo", "bar"}) here.

clang/lib/Driver/ToolChains/HIP.cpp

283

Nit: Path was fine.

[yaxunl]

thanks. will fix when commit

[njames93]

This is causing a test case to fail on mac http://45.33.8.238/mac/17048/step_7.txt

[thakis]

This breaks check-clang on macOS: http://45.33.8.238/mac/17053/step_7.txt

The output contains

5: clang: warning: argument unused during compilation: '--rocm-path=/Users/thakis/src/llvm-project/clang/test/Driver/Inputs/rocm' [-Wunused-command-line-argument]

which if true would explain the failure.

That bot uses the GN build, but I verified that the test fails the same way in the CMake build.

Please take a look, and if it takes a while to fix, please revert while you investigate.

clang/test/Driver/Inputs/rocm/bin/.hipVersion

2

Is this true? I don't see any cmake code generating this. Did I miss it?

[thakis]

…oh, njames said that already 5h ago :) I guess I'll wait another hour or two, and then I'll revert.

[yaxunl]

In D82930#2145914, @thakis wrote:

This breaks check-clang on macOS: http://45.33.8.238/mac/17053/step_7.txt

The output contains

5: clang: warning: argument unused during compilation: '--rocm-path=/Users/thakis/src/llvm-project/clang/test/Driver/Inputs/rocm' [-Wunused-command-line-argument]

which if true would explain the failure.

That bot uses the GN build, but I verified that the test fails the same way in the CMake build.

Please take a look, and if it takes a while to fix, please revert while you investigate.

should be fixed by 5d2c3e031a6861b3e95673d0e238c09938dd9c0d

[yaxunl]

In D82930#2145827, @njames93 wrote:

This is causing a test case to fail on mac http://45.33.8.238/mac/17048/step_7.txt

5d2c3e031a6861b3e95673d0e238c09938dd9c0d should fix it

[thakis]

Great, thanks :)

clang/include/clang/Basic/DiagnosticDriverKinds.td

 def err_drv_cuda_bad_gpu_arch : Error<"Unsupported CUDA gpu architecture: %0">;
 def err_drv_no_cuda_installation : Error<
   "cannot find CUDA installation. Provide its path via --cuda-path, or pass "
   "-nocudainc to build without CUDA includes.">;
 def err_drv_no_cuda_libdevice : Error<
   "cannot find libdevice for %0. Provide path to different CUDA installation "
   "via --cuda-path, or pass -nocudalib to build without linking with libdevice.">;
 
-def err_drv_no_rocm_installation : Error<
-  "cannot find ROCm installation.  Provide its path via --rocm-path, or pass "
-  "-nogpulib and -nogpuinc to build without ROCm device library and HIP includes.">;
 def err_drv_no_rocm_device_lib : Error<
-  "cannot find device library for %0. Provide path to different ROCm installation "
-  "via --rocm-path, or pass -nogpulib to build without linking default libraries.">;
+  "cannot find ROCm device library%select{| for %1}0. Provide its path via --rocm-path or "
+  "--rocm-device-lib-path, or pass -nogpulib to build without ROCm device library.">;
+def err_drv_no_hip_runtime : Error<
+  "cannot find HIP runtime. Provide its path via --rocm-path, or pass "
+  "-nogpuinc to build without HIP runtime.">;
 
 def err_drv_cuda_version_unsupported : Error<
   "GPU arch %0 is supported by CUDA versions between %1 and %2 (inclusive), "
   "but installation at %3 is %4. Use --cuda-path to specify a different CUDA "
   "install, pass a different GPU arch with --cuda-gpu-arch, or pass "
   "--no-cuda-version-check.">;
 def warn_drv_unknown_cuda_version: Warning<
   "Unknown CUDA version %0. Assuming the latest supported version %1">,

clang/include/clang/Driver/Options.td

 defm gpu_rdc : OptInFFlag<"gpu-rdc",
   "Generate relocatable device code, also known as separate compilation mode", "", "">;
 def : Flag<["-"], "fcuda-rdc">, Alias<fgpu_rdc>;
 def : Flag<["-"], "fno-cuda-rdc">, Alias<fno_gpu_rdc>;
 defm cuda_short_ptr : OptInFFlag<"cuda-short-ptr",
   "Use 32-bit pointers for accessing const/local/shared address spaces">;
 def rocm_path_EQ : Joined<["--"], "rocm-path=">, Group<i_Group>,
   HelpText<"ROCm installation path, used for finding and automatically linking required bitcode libraries.">;
-def hip_device_lib_path_EQ : Joined<["--"], "hip-device-lib-path=">, Group<Link_Group>,
-  HelpText<"HIP device library path. Alternative to rocm-path.">;
+def rocm_device_lib_path_EQ : Joined<["--"], "rocm-device-lib-path=">, Group<Link_Group>,
+  HelpText<"ROCm device library path. Alternative to rocm-path.">;
+def : Joined<["--"], "hip-device-lib-path=">, Alias<rocm_device_lib_path_EQ>;
 def hip_device_lib_EQ : Joined<["--"], "hip-device-lib=">, Group<Link_Group>,
   HelpText<"HIP device library">;
+def hip_version_EQ : Joined<["--"], "hip-version=">,
+  HelpText<"HIP version in the format of major.minor.patch">;
 def fhip_dump_offload_linker_script : Flag<["-"], "fhip-dump-offload-linker-script">,
   Group<f_Group>, Flags<[NoArgumentUnused, HelpHidden]>;
 defm hip_new_launch_api : OptInFFlag<"hip-new-launch-api",
-  "Use new kernel launching API for HIP">;
-defm gpu_allow_device_init : OptInFFlag<"gpu-allow-device-init", "Allow device side init function in HIP">;
+  "Use", "Don't use", " new kernel launching API for HIP">;
+defm gpu_allow_device_init : OptInFFlag<"gpu-allow-device-init",
+  "Allow", "Don't allow", " device side init function in HIP">;
 def gpu_max_threads_per_block_EQ : Joined<["--"], "gpu-max-threads-per-block=">,
   Flags<[CC1Option]>,
   HelpText<"Default max threads per block for kernel launch bounds for HIP">;
 def libomptarget_nvptx_path_EQ : Joined<["--"], "libomptarget-nvptx-path=">, Group<i_Group>,
   HelpText<"Path to libomptarget-nvptx libraries">;
 def dD : Flag<["-"], "dD">, Group<d_Group>, Flags<[CC1Option]>,
   HelpText<"Print macro definitions in -E mode in addition to normal output">;
 def dI : Flag<["-"], "dI">, Group<d_Group>, Flags<[CC1Option]>,

clang/lib/CodeGen/CGCUDANV.cpp

std::string CGNVCUDARuntime::getDeviceSideName(const NamedDecl *ND) {
   return DeviceSideName;
 }
 
 void CGNVCUDARuntime::emitDeviceStub(CodeGenFunction &CGF,
                                      FunctionArgList &Args) {
   EmittedKernels.push_back({CGF.CurFn, CGF.CurFuncDecl});
   if (CudaFeatureEnabled(CGM.getTarget().getSDKVersion(),
                          CudaFeature::CUDA_USES_NEW_LAUNCH) ||
-      CGF.getLangOpts().HIPUseNewLaunchAPI)
+      (CGF.getLangOpts().HIP && CGF.getLangOpts().HIPUseNewLaunchAPI))
     emitDeviceStubBodyNew(CGF, Args);
   else
     emitDeviceStubBodyLegacy(CGF, Args);
 }
 
 // CUDA 9.0+ uses new way to launch kernels. Parameters are packed in a local
 // array and kernels are launched using cudaLaunchKernel().
 void CGNVCUDARuntime::emitDeviceStubBodyNew(CodeGenFunction &CGF,

clang/lib/Driver/ToolChains/AMDGPU.cpp

if (BaseName == "ocml") {
       llvm::Twine GfxName = Twine("gfx") + IsaVersionNumber;
       SmallString<8> Tmp;
       LibDeviceMap.insert(
         std::make_pair(GfxName.toStringRef(Tmp), FilePath.str()));
     }
   }
 }
 
-RocmInstallationDetector::RocmInstallationDetector(
-    const Driver &D, const llvm::Triple &HostTriple,
-    const llvm::opt::ArgList &Args)
-    : D(D) {
-  struct Candidate {
-    std::string Path;
-    bool StrictChecking;
-
-    Candidate(std::string Path, bool StrictChecking = false)
-        : Path(Path), StrictChecking(StrictChecking) {}
-  };
-
+void RocmInstallationDetector::ParseHIPVersionFile(llvm::StringRef V) {
+  SmallVector<StringRef, 4> VersionParts;
+  V.split(VersionParts, '\n');
+  unsigned Major;
+  unsigned Minor;
+  for (auto Part : VersionParts) {
+    auto Splits = Part.split('=');
+    if (Splits.first == "HIP_VERSION_MAJOR")
+      Splits.second.getAsInteger(0, Major);
+    else if (Splits.first == "HIP_VERSION_MINOR")
+      Splits.second.getAsInteger(0, Minor);
+    else if (Splits.first == "HIP_VERSION_PATCH")
+      VersionPatch = Splits.second.str();
+  }
+  VersionMajorMinor = llvm::VersionTuple(Major, Minor);
+  DetectedVersion =
+      (Twine(Major) + "." + Twine(Minor) + "." + VersionPatch).str();
+}
+
+// For candidate specified by --rocm-path we do not do strict check.
+SmallVector<RocmInstallationDetector::Candidate, 4>
+RocmInstallationDetector::getInstallationPathCandidates() {
   SmallVector<Candidate, 4> Candidates;
+  if (!RocmPathArg.empty()) {
+    Candidates.emplace_back(RocmPathArg.str());
+    return Candidates;
+  }
 
-  if (Args.hasArg(clang::driver::options::OPT_rocm_path_EQ)) {
-    Candidates.emplace_back(
-        Args.getLastArgValue(clang::driver::options::OPT_rocm_path_EQ).str());
-  } else {
-    // Try to find relative to the compiler binary.
-    const char *InstallDir = D.getInstalledDir();
 
-    // Check both a normal Unix prefix position of the clang binary, as well as
-    // the Windows-esque layout the ROCm packages use with the host architecture
-    // subdirectory of bin.
 
-    // Strip off directory (usually bin)
-    StringRef ParentDir = llvm::sys::path::parent_path(InstallDir);
-    StringRef ParentName = llvm::sys::path::filename(ParentDir);
 
-    // Some builds use bin/{host arch}, so go up again.
-    if (ParentName == "bin") {
-      ParentDir = llvm::sys::path::parent_path(ParentDir);
-      ParentName = llvm::sys::path::filename(ParentDir);
-    }
 
-    if (ParentName == "llvm") {
-      // Some versions of the rocm llvm package install to /opt/rocm/llvm/bin
-      Candidates.emplace_back(llvm::sys::path::parent_path(ParentDir).str(),
-                              /*StrictChecking=*/true);
-    }
+  if (ParentName == "llvm")
+    ParentDir = llvm::sys::path::parent_path(ParentDir);
+
+  Candidates.emplace_back(ParentDir.str(), /*StrictChecking=*/true);
+
+  // Device library may be installed in clang resource directory.
+  Candidates.emplace_back(D.ResourceDir, /*StrictChecking=*/true);
+
+  Candidates.emplace_back(D.SysRoot + "/opt/rocm", /*StrictChecking=*/true);
+  return Candidates;
 
-    Candidates.emplace_back(D.SysRoot + "/opt/rocm");
+RocmInstallationDetector::RocmInstallationDetector(
+    const Driver &D, const llvm::Triple &HostTriple,
+    const llvm::opt::ArgList &Args, bool DetectHIPRuntime, bool DetectDeviceLib)
+    : D(D) {
+  RocmPathArg = Args.getLastArgValue(clang::driver::options::OPT_rocm_path_EQ);
+  RocmDeviceLibPathArg =
+      Args.getAllArgValues(clang::driver::options::OPT_rocm_device_lib_path_EQ);
+  if (auto *A = Args.getLastArg(clang::driver::options::OPT_hip_version_EQ)) {
+    HIPVersionArg = A->getValue();
+    unsigned Major = 0;
+    unsigned Minor = 0;
+    auto Splits = HIPVersionArg.split('.');
+    Splits.first.getAsInteger(0, Major);
+    Splits = Splits.second.split('.');
+    Splits.first.getAsInteger(0, Minor);
+    VersionPatch = Splits.second.str();
+    if (VersionPatch.empty())
+      VersionPatch = "0";
+    if (Major == 0 || Minor == 0)
+      D.Diag(diag::err_drv_invalid_value)
+          << A->getAsString(Args) << HIPVersionArg;
+
+    VersionMajorMinor = llvm::VersionTuple(Major, Minor);
+    DetectedVersion =
+        (Twine(Major) + "." + Twine(Minor) + "." + VersionPatch).str();
+  } else {
+    VersionPatch = DefaultVersionPatch;
+    VersionMajorMinor =
+        llvm::VersionTuple(DefaultVersionMajor, DefaultVersionMinor);
+    DetectedVersion = (Twine(DefaultVersionMajor) + "." +
+                       Twine(DefaultVersionMinor) + "." + VersionPatch)
+                          .str();
   }
 
-  bool NoBuiltinLibs = Args.hasArg(options::OPT_nogpulib);
+  if (DetectHIPRuntime)
+    detectHIPRuntime();
+  if (DetectDeviceLib)
+    detectDeviceLibrary();
+}
 
+void RocmInstallationDetector::detectDeviceLibrary() {
   assert(LibDevicePath.empty());
 
-  if (Args.hasArg(clang::driver::options::OPT_hip_device_lib_path_EQ)) {
-    LibDevicePath
-      = Args.getLastArgValue(clang::driver::options::OPT_hip_device_lib_path_EQ);
-  } else if (const char *LibPathEnv = ::getenv("HIP_DEVICE_LIB_PATH")) {
+  if (!RocmDeviceLibPathArg.empty())
+    LibDevicePath = RocmDeviceLibPathArg[RocmDeviceLibPathArg.size() - 1];
+  else if (const char *LibPathEnv = ::getenv("HIP_DEVICE_LIB_PATH"))
     LibDevicePath = LibPathEnv;
-  }
 
   auto &FS = D.getVFS();
   if (!LibDevicePath.empty()) {
     // Maintain compatability with HIP flag/envvar pointing directly at the
     // bitcode library directory. This points directly at the library path instead
     // of the rocm root installation.
     if (!FS.exists(LibDevicePath))
       return;
 
     scanLibDevicePath();
-    IsValid = allGenericLibsValid() && !LibDeviceMap.empty();
+    HasDeviceLibrary = allGenericLibsValid() && !LibDeviceMap.empty();
     return;
   }
 
-  for (const auto &Candidate : Candidates) {
-    InstallPath = Candidate.Path;
-    if (InstallPath.empty() || !FS.exists(InstallPath))
-      continue;
-
-    // The install path situation in old versions of ROCm is a real mess, and
-    // use a different install layout. Multiple copies of the device libraries
-    // exist for each frontend project, and differ depending on which build
-    // system produced the packages. Standalone OpenCL builds also have a
-    // different directory structure from the ROCm OpenCL package.
-    //
-    // The desired structure is (${ROCM_ROOT} or
-    // ${OPENCL_ROOT})/amdgcn/bitcode/*, so try to detect this layout.
-
-    // BinPath = InstallPath + "/bin";
-    llvm::sys::path::append(IncludePath, InstallPath, "include");
-    llvm::sys::path::append(LibDevicePath, InstallPath, "amdgcn", "bitcode");
-
-    // We don't need the include path for OpenCL, since clang already ships with
-    // the default header.
-
+  // The desired structure is ${ROCM_ROOT})/amdgcn/bitcode/* or
+  // ${ROCM_ROOT}/lib/* or ${ROCM_ROOT}/lib/bitcode/* so try to detect this
+  // layout.
+  auto Candidates = getInstallationPathCandidates();
+  for (const auto &Candidate : Candidates) {
+    InstallPath = Candidate.Path;
+    auto CheckDeviceLib = [&]() {
       bool CheckLibDevice = (!NoBuiltinLibs || Candidate.StrictChecking);
       if (CheckLibDevice && !FS.exists(LibDevicePath))
-      continue;
+        return false;
 
       scanLibDevicePath();
 
       if (!NoBuiltinLibs) {
         // Check that the required non-target libraries are all available.
         if (!allGenericLibsValid())
-        continue;
+          return false;
 
-      // Check that we have found at least one libdevice that we can link in if
-      // -nobuiltinlib hasn't been specified.
+        // Check that we have found at least one libdevice that we can link in
+        // if -nobuiltinlib hasn't been specified.
         if (LibDeviceMap.empty())
-        continue;
+          return false;
       }
+      return true;
+    };
 
-    IsValid = true;
-    break;
+    LibDevicePath = InstallPath;
+    llvm::sys::path::append(LibDevicePath, "amdgcn", "bitcode");
+    HasDeviceLibrary = CheckDeviceLib();
+    if (HasDeviceLibrary)
+      return;
+    LibDevicePath = InstallPath;
+    llvm::sys::path::append(LibDevicePath, "lib");
+    HasDeviceLibrary = CheckDeviceLib();
+    if (HasDeviceLibrary)
+      return;
+    LibDevicePath = InstallPath;
+    llvm::sys::path::append(LibDevicePath, "lib", "bitcode");
+    HasDeviceLibrary = CheckDeviceLib();
+    if (HasDeviceLibrary)
+      return;
   }
 }
 
+void RocmInstallationDetector::detectHIPRuntime() {
+  auto Candidates = getInstallationPathCandidates();
+  auto &FS = D.getVFS();
+
+  for (const auto &Candidate : Candidates) {
+    InstallPath = Candidate.Path;
+    if (InstallPath.empty() || !FS.exists(InstallPath))
+      continue;
+
+    BinPath = InstallPath;
+    llvm::sys::path::append(BinPath, "bin");
+    IncludePath = InstallPath;
+    llvm::sys::path::append(IncludePath, "include");
+    LibPath = InstallPath;
+    llvm::sys::path::append(LibPath, "lib");
+
+    llvm::ErrorOr<std::unique_ptr<llvm::MemoryBuffer>> VersionFile =
+        FS.getBufferForFile(BinPath + "/.hipVersion");
+    if (!VersionFile && Candidate.StrictChecking)
+      continue;
+
+    if (HIPVersionArg.empty() && VersionFile)
+      ParseHIPVersionFile((*VersionFile)->getBuffer());
+
+    HasHIPRuntime = true;
+    return;
+  }
+  HasHIPRuntime = false;
+}
+
 void RocmInstallationDetector::print(raw_ostream &OS) const {
-  if (isValid())
-    OS << "Found ROCm installation: " << InstallPath << '\n';
+  if (hasHIPRuntime())
+    OS << "Found HIP installation: " << InstallPath << ", version "
+       << DetectedVersion << '\n';
 }
 
 void RocmInstallationDetector::AddHIPIncludeArgs(const ArgList &DriverArgs,
                                                  ArgStringList &CC1Args) const {
+  bool IsRocm35 = VersionMajorMinor <= llvm::VersionTuple(3, 5);
+
   if (!DriverArgs.hasArg(options::OPT_nobuiltininc)) {
     // HIP header includes standard library wrapper headers under clang
     // cuda_wrappers directory. Since these wrapper headers include_next
     // standard C++ headers, whereas libc++ headers include_next other clang
     // headers. The include paths have to follow this order:
     // - wrapper include path
     // - standard C++ include path
     // - other clang include path
     // Since standard C++ and other clang include paths are added in other
     // places after this function, here we only need to make sure wrapper
     // include path is added.
+    //
+    // ROCm 3.5 does not fully support the wrapper headers. Therefore it needs
+    // a workaround.
     SmallString<128> P(D.ResourceDir);
-    llvm::sys::path::append(P, "include");
-    llvm::sys::path::append(P, "cuda_wrappers");
+    if (!IsRocm35)
+      llvm::sys::path::append(P, "include", "cuda_wrappers");
     CC1Args.push_back("-internal-isystem");
     CC1Args.push_back(DriverArgs.MakeArgString(P));
   }
 
   if (DriverArgs.hasArg(options::OPT_nogpuinc))
     return;
 
-  if (!isValid()) {
-    D.Diag(diag::err_drv_no_rocm_installation);
+  if (!hasHIPRuntime()) {
+    D.Diag(diag::err_drv_no_hip_runtime);
     return;
   }
 
   CC1Args.push_back("-internal-isystem");
   CC1Args.push_back(DriverArgs.MakeArgString(getIncludePath()));
+  if (!IsRocm35) {
     CC1Args.push_back("-include");
     CC1Args.push_back("__clang_hip_runtime_wrapper.h");
   }
+}
 
 void amdgpu::Linker::ConstructJob(Compilation &C, const JobAction &JA,
                                   const InputInfo &Output,
                                   const InputInfoList &Inputs,
                                   const ArgList &Args,
                                   const char *LinkingOutput) const {
 
   std::string Linker = getToolChain().GetProgramPath(getShortName());
bool AMDGPUToolChain::isWave64(const llvm::opt::ArgList &DriverArgs,
   return !HasWave32 || DriverArgs.hasFlag(
     options::OPT_mwavefrontsize64, options::OPT_mno_wavefrontsize64, false);
 }
 
 
 /// ROCM Toolchain
 ROCMToolChain::ROCMToolChain(const Driver &D, const llvm::Triple &Triple,
                              const ArgList &Args)
     : AMDGPUToolChain(D, Triple, Args),
-    RocmInstallation(D, Triple, Args) { }
+      RocmInstallation(D, Triple, Args, /*DetectHIPRuntime=*/false,
+                       /*DetectDeviceLib=*/true) {}
 
 void AMDGPUToolChain::addClangTargetOptions(
     const llvm::opt::ArgList &DriverArgs,
     llvm::opt::ArgStringList &CC1Args,
     Action::OffloadKind DeviceOffloadingKind) const {
   // Default to "hidden" visibility, as object level linking will not be
   // supported for the foreseeable future.
   if (!DriverArgs.hasArg(options::OPT_fvisibility_EQ,
void ROCMToolChain::addClangTargetOptions(
   // disable bitcode linking.
   if (DeviceOffloadingKind == Action::OFK_None &&
       DriverArgs.hasArg(options::OPT_nostdlib))
     return;
 
   if (DriverArgs.hasArg(options::OPT_nogpulib))
     return;
 
-  if (!RocmInstallation.isValid()) {
-    getDriver().Diag(diag::err_drv_no_rocm_installation);
+  if (!RocmInstallation.hasDeviceLibrary()) {
+    getDriver().Diag(diag::err_drv_no_rocm_device_lib) << 0;
     return;
   }
 
   // Get the device name and canonicalize it
   const StringRef GpuArch = DriverArgs.getLastArgValue(options::OPT_mcpu_EQ);
   auto Kind = llvm::AMDGPU::parseArchAMDGCN(GpuArch);
   const StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind);
   std::string LibDeviceFile = RocmInstallation.getLibDeviceFile(CanonArch);
   if (LibDeviceFile.empty()) {
-    getDriver().Diag(diag::err_drv_no_rocm_device_lib) << GpuArch;
+    getDriver().Diag(diag::err_drv_no_rocm_device_lib) << 1 << GpuArch;
     return;
   }
 
   bool Wave64 = isWave64(DriverArgs, Kind);
 
   // TODO: There are way too many flags that change this. Do we need to check
   // them all?
   bool DAZ = DriverArgs.hasArg(options::OPT_cl_denorms_are_zero) ||

clang/lib/Driver/ToolChains/Clang.cpp

if (&I == &Input) {
       CudaDeviceInput = &I;
     } else if (IsOpenMPDevice && !OpenMPDeviceInput) {
       OpenMPDeviceInput = &I;
     } else {
       llvm_unreachable("unexpectedly given multiple inputs");
     }
   }
 
-  const llvm::Triple *AuxTriple = IsCuda ? TC.getAuxTriple() : nullptr;
+  const llvm::Triple *AuxTriple =
+      (IsCuda || IsHIP) ? TC.getAuxTriple() : nullptr;
   bool IsWindowsMSVC = RawTriple.isWindowsMSVCEnvironment();
   bool IsIAMCU = RawTriple.isOSIAMCU();
 
   // Adjust IsWindowsXYZ for CUDA/HIP compilations.  Even when compiling in
   // device mode (i.e., getToolchain().getTriple() is NVPTX/AMDGCN, not
   // Windows), we need to pass Windows-specific flags to cc1.
   if (IsCuda || IsHIP)
     IsWindowsMSVC |= AuxTriple && AuxTriple->isWindowsMSVCEnvironment();
#endif
   if (Args.hasFlag(options::OPT_funique_basic_block_section_names,
                    options::OPT_fno_unique_basic_block_section_names, false))
     CmdArgs.push_back("-funique-basic-block-section-names");
 
   Args.AddLastArg(CmdArgs, options::OPT_finstrument_functions,
                   options::OPT_finstrument_functions_after_inlining,
                   options::OPT_finstrument_function_entry_bare);
 
-  // NVPTX doesn't support PGO or coverage. There's no runtime support for
-  // sampling, overhead of call arc collection is way too high and there's no
-  // way to collect the output.
-  if (!Triple.isNVPTX())
+  // NVPTX/AMDGCN doesn't support PGO or coverage. There's no runtime support
+  // for sampling, overhead of call arc collection is way too high and there's
+  // no way to collect the output.
+  if (!Triple.isNVPTX() && !Triple.isAMDGCN())
     addPGOAndCoverageFlags(TC, C, D, Output, Args, CmdArgs);
 
   Args.AddLastArg(CmdArgs, options::OPT_fclang_abi_compat_EQ);
 
   // Add runtime flag for PS4 when PGO, coverage, or sanitizers are enabled.
   if (RawTriple.isPS4CPU() &&
       !Args.hasArg(options::OPT_nostdlib, options::OPT_nodefaultlibs)) {
     PS4cpu::addProfileRTArgs(TC, Args, CmdArgs);
if (Std) {
     if (!types::isCXX(InputType))
       Args.AddAllArgsTranslated(CmdArgs, options::OPT_std_default_EQ, "-std=",
                                 /*Joined=*/true);
     else if (IsWindowsMSVC)
       ImplyVCPPCXXVer = true;
 
     Args.AddLastArg(CmdArgs, options::OPT_ftrigraphs,
                     options::OPT_fno_trigraphs);
+
+    // HIP headers has minimum C++ standard requirements. Therefore set the
+    // default language standard.
+    if (IsHIP)
+      CmdArgs.push_back(IsWindowsMSVC ? "-std=c++14" : "-std=c++11");
   }
 
   // GCC's behavior for -Wwrite-strings is a bit strange:
   //  * In C, this "warning flag" changes the types of string literals from
   //    'char[N]' to 'const char[N]', and thus triggers an unrelated warning
   //    for the discarded qualifier.
   //  * In C++, this is just a normal warning flag.
   //
if (Arg *A = Args.getLastArg(options::OPT_mrestrict_it,
     // Windows on ARM expects restricted IT blocks
     CmdArgs.push_back("-mllvm");
     CmdArgs.push_back("-arm-restrict-it");
   }
 
   // Forward -cl options to -cc1
   RenderOpenCLOptions(Args, CmdArgs);
 
-  if (Args.hasFlag(options::OPT_fhip_new_launch_api,
-                   options::OPT_fno_hip_new_launch_api, false))
+  if (IsHIP && Args.hasFlag(options::OPT_fhip_new_launch_api,
+                            options::OPT_fno_hip_new_launch_api, true))
     CmdArgs.push_back("-fhip-new-launch-api");
 
   if (Arg *A = Args.getLastArg(options::OPT_fcf_protection_EQ)) {
     CmdArgs.push_back(
         Args.MakeArgString(Twine("-fcf-protection=") + A->getValue()));
   }
 
   // Forward -f options with positive and negative forms; we translate

clang/lib/Driver/ToolChains/Gnu.cpp

 }
 
 Tool *Generic_GCC::buildLinker() const { return new tools::gcc::Linker(*this); }
 
 void Generic_GCC::printVerboseInfo(raw_ostream &OS) const {
   // Print the information about how we detected the GCC installation.
   GCCInstallation.print(OS);
   CudaInstallation.print(OS);
+  RocmInstallation.print(OS);
 }
 
 bool Generic_GCC::IsUnwindTablesDefault(const ArgList &Args) const {
   return getArch() == llvm::Triple::x86_64;
 }
 
 bool Generic_GCC::isPICDefault() const {
   switch (getArch()) {

clang/lib/Driver/ToolChains/HIP.cpp

 }
 
 HIPToolChain::HIPToolChain(const Driver &D, const llvm::Triple &Triple,
                              const ToolChain &HostTC, const ArgList &Args)
     : ROCMToolChain(D, Triple, Args), HostTC(HostTC) {
   // Lookup binaries into the driver directory, this is used to
   // discover the clang-offload-bundler executable.
   getProgramPaths().push_back(getDriver().Dir);
+  RocmInstallation.detectHIPRuntime();
 }
 
 void HIPToolChain::addClangTargetOptions(
     const llvm::opt::ArgList &DriverArgs,
     llvm::opt::ArgStringList &CC1Args,
     Action::OffloadKind DeviceOffloadingKind) const {
   HostTC.addClangTargetOptions(DriverArgs, CC1Args, DeviceOffloadingKind);
 
if (!DriverArgs.hasArg(options::OPT_fvisibility_EQ,
     CC1Args.push_back("-fapply-global-visibility-to-externs");
   }
 
   if (DriverArgs.hasArg(options::OPT_nogpulib))
     return;
   ArgStringList LibraryPaths;
 
   // Find in --hip-device-lib-path and HIP_LIBRARY_PATH.
-  for (auto Path :
-       DriverArgs.getAllArgValues(options::OPT_hip_device_lib_path_EQ))
-    LibraryPaths.push_back(DriverArgs.MakeArgString(Path));
+  for (auto P : RocmInstallation.getRocmDeviceLibPathArg())
+    LibraryPaths.push_back(DriverArgs.MakeArgString(P));
 
   addDirectoryList(DriverArgs, LibraryPaths, "", "HIP_DEVICE_LIB_PATH");
 
   // Maintain compatability with --hip-device-lib.
   auto BCLibs = DriverArgs.getAllArgValues(options::OPT_hip_device_lib_EQ);
   if (!BCLibs.empty()) {
     for (auto Lib : BCLibs)
       addBCLib(getDriver(), DriverArgs, CC1Args, LibraryPaths, Lib);
   } else {
-    if (!RocmInstallation.isValid()) {
-      getDriver().Diag(diag::err_drv_no_rocm_installation);
+    if (!RocmInstallation.hasDeviceLibrary()) {
+      getDriver().Diag(diag::err_drv_no_rocm_device_lib) << 0;
       return;
     }
 
     std::string LibDeviceFile = RocmInstallation.getLibDeviceFile(CanonArch);
     if (LibDeviceFile.empty()) {
-      getDriver().Diag(diag::err_drv_no_rocm_device_lib) << GpuArch;
+      getDriver().Diag(diag::err_drv_no_rocm_device_lib) << 1 << GpuArch;
       return;
     }
 
     // If --hip-device-lib is not set, add the default bitcode libraries.
     // TODO: There are way too many flags that change this. Do we need to check
     // them all?
     bool DAZ = DriverArgs.hasFlag(options::OPT_fcuda_flush_denormals_to_zero,
                                   options::OPT_fno_cuda_flush_denormals_to_zero,

clang/lib/Driver/ToolChains/MSVC.cpp

 
 void MSVCToolChain::AddHIPIncludeArgs(const ArgList &DriverArgs,
                                       ArgStringList &CC1Args) const {
   RocmInstallation.AddHIPIncludeArgs(DriverArgs, CC1Args);
 }
 
 void MSVCToolChain::printVerboseInfo(raw_ostream &OS) const {
   CudaInstallation.print(OS);
+  RocmInstallation.print(OS);
 }
 
 // Windows SDKs and VC Toolchains group their contents into subdirectories based
 // on the target architecture. This function converts an llvm::Triple::ArchType
 // to the corresponding subdirectory name.
 static const char *llvmArchToWindowsSDKArch(llvm::Triple::ArchType Arch) {
   using ArchType = llvm::Triple::ArchType;
   switch (Arch) {

clang/lib/Driver/ToolChains/ROCm.h

 #include "clang/Basic/LLVM.h"
 #include "clang/Driver/Driver.h"
 #include "clang/Driver/Options.h"
 #include "llvm/ADT/SmallSet.h"
 #include "llvm/ADT/SmallString.h"
 #include "llvm/ADT/StringMap.h"
 #include "llvm/ADT/Triple.h"
 #include "llvm/Option/ArgList.h"
+#include "llvm/Support/VersionTuple.h"
 
 namespace clang {
 namespace driver {
 
 /// A class to find a viable ROCM installation
 /// TODO: Generalize to handle libclc.
 class RocmInstallationDetector {
 private:
   struct ConditionalLibrary {
     SmallString<0> On;
     SmallString<0> Off;
 
     bool isValid() const { return !On.empty() && !Off.empty(); }
 
     StringRef get(bool Enabled) const {
       assert(isValid());
       return Enabled ? On : Off;
     }
   };
 
+  // Installation path candidate.
+  struct Candidate {
+    std::string Path;
+    bool StrictChecking;
+
+    Candidate(std::string Path, bool StrictChecking = false)
+        : Path(Path), StrictChecking(StrictChecking) {}
+  };
+
   const Driver &D;
-  bool IsValid = false;
-  // RocmVersion Version = RocmVersion::UNKNOWN;
+  bool HasHIPRuntime = false;
+  bool HasDeviceLibrary = false;
+
+  // Default version if not detected or specified.
+  const unsigned DefaultVersionMajor = 3;
+  const unsigned DefaultVersionMinor = 5;
+  const char *DefaultVersionPatch = "0";
+
+  // The version string in Major.Minor.Patch format.
+  std::string DetectedVersion;
+  // Version containing major and minor.
+  llvm::VersionTuple VersionMajorMinor;
+  // Version containing patch.
+  std::string VersionPatch;
+
+  // ROCm path specified by --rocm-path.
+  StringRef RocmPathArg;
+  // ROCm device library paths specified by --rocm-device-lib-path.
+  std::vector<std::string> RocmDeviceLibPathArg;
+  // HIP version specified by --hip-version.
+  StringRef HIPVersionArg;
+  // Wheter -nogpulib is specified.
+  bool NoBuiltinLibs = false;
+
+  // Paths
   SmallString<0> InstallPath;
-  // SmallString<0> BinPath;
+  SmallString<0> BinPath;
   SmallString<0> LibPath;
   SmallString<0> LibDevicePath;
   SmallString<0> IncludePath;
   llvm::StringMap<std::string> LibDeviceMap;
 
   // Libraries that are always linked.
   SmallString<0> OCML;
   SmallString<0> OCKL;
return !OCML.empty() && !OCKL.empty() && !OpenCL.empty() && !HIP.empty() &&
            CorrectlyRoundedSqrt.isValid();
   }
 
   // GPU architectures for which we have raised an error in
   // CheckRocmVersionSupportsArch.
   mutable llvm::SmallSet<CudaArch, 4> ArchsWithBadVersion;
 
   void scanLibDevicePath();
+  void ParseHIPVersionFile(llvm::StringRef V);
+  SmallVector<Candidate, 4> getInstallationPathCandidates();
 
 public:
   RocmInstallationDetector(const Driver &D, const llvm::Triple &HostTriple,
-                           const llvm::opt::ArgList &Args);
+                           const llvm::opt::ArgList &Args,
+                           bool DetectHIPRuntime = true,
+                           bool DetectDeviceLib = false);
 
   /// Add arguments needed to link default bitcode libraries.
   void addCommonBitcodeLibCC1Args(const llvm::opt::ArgList &DriverArgs,
                                   llvm::opt::ArgStringList &CC1Args,
                                   StringRef LibDeviceFile, bool Wave64,
                                   bool DAZ, bool FiniteOnly, bool UnsafeMathOpt,
                                   bool FastRelaxedMath, bool CorrectSqrt) const;
 
   /// Emit an error if Version does not support the given Arch.
   ///
   /// If either Version or Arch is unknown, does not emit an error.  Emits at
   /// most one error per Arch.
   void CheckRocmVersionSupportsArch(CudaArch Arch) const;
 
-  /// Check whether we detected a valid Rocm install.
-  bool isValid() const { return IsValid; }
+  /// Check whether we detected a valid HIP runtime.
+  bool hasHIPRuntime() const { return HasHIPRuntime; }
+
+  /// Check whether we detected a valid ROCm device library.
+  bool hasDeviceLibrary() const { return HasDeviceLibrary; }
+
   /// Print information about the detected ROCm installation.
   void print(raw_ostream &OS) const;
 
   /// Get the detected Rocm install's version.
   // RocmVersion version() const { return Version; }
 
   /// Get the detected Rocm installation path.
   StringRef getInstallPath() const { return InstallPath; }
public:
 
   /// Get libdevice file for given architecture
   std::string getLibDeviceFile(StringRef Gpu) const {
     return LibDeviceMap.lookup(Gpu);
   }
 
   void AddHIPIncludeArgs(const llvm::opt::ArgList &DriverArgs,
                          llvm::opt::ArgStringList &CC1Args) const;
+
+  void detectDeviceLibrary();
+  void detectHIPRuntime();
+
+  /// Get the values for --rocm-device-lib-path arguments
+  std::vector<std::string> getRocmDeviceLibPathArg() const {
+    return RocmDeviceLibPathArg;
+  }
+
+  /// Get the value for --rocm-path argument
+  StringRef getRocmPathArg() const { return RocmPathArg; }
+
+  /// Get the value for --hip-version argument
+  StringRef getHIPVersionArg() const { return HIPVersionArg; }
+
+  std::string getHIPVersion() const { return DetectedVersion; }
 };
 
 } // end namespace driver
 } // end namespace clang
 
 #endif // LLVM_CLANG_LIB_DRIVER_TOOLCHAINS_ROCM_H

clang/test/Driver/Inputs/rocm/bin/.hipVersion

+# Auto-generated by cmake
+HIP_VERSION_MAJOR=3
+HIP_VERSION_MINOR=6
+HIP_VERSION_PATCH=20214-a2917cd

clang/test/Driver/hip-launch-api.hip

+// REQUIRES: clang-driver
+// REQUIRES: x86-registered-target
+// REQUIRES: amdgpu-registered-target
+
+// By default FE assumes -fhip-new-launch-api.
+
+// RUN: %clang -### -target x86_64-unknown-linux-gnu -offload-arch=gfx906 %s \
+// RUN:   2>&1 | FileCheck -check-prefixes=NEW %s
+// NEW: "-fhip-new-launch-api"
+
+// RUN: %clang -### -target x86_64-unknown-linux-gnu -offload-arch=gfx906 %s \
+// RUN:   -fhip-new-launch-api 2>&1 | FileCheck -check-prefixes=NEW %s
+// NEW: "-fhip-new-launch-api"
+
+// RUN: %clang -### -target x86_64-unknown-linux-gnu -offload-arch=gfx906 %s \
+// RUN:   -fno-hip-new-launch-api 2>&1 | FileCheck -check-prefixes=OLD %s
+// OLD-NOT: "-fhip-new-launch-api"

clang/test/Driver/hip-std.hip

+// REQUIRES: clang-driver
+// REQUIRES: x86-registered-target
+// REQUIRES: amdgpu-registered-target
+
+// RUN: %clang -### -target x86_64-unknown-linux-gnu -offload-arch=gfx906 %s \
+// RUN:   2>&1 | FileCheck -check-prefixes=DEFAULT %s
+// DEFAULT: "{{.*}}clang{{.*}}" "-cc1"{{.*}}"-fcuda-is-device"{{.*}}"-std=c++11"
+// DEFAULT: "{{.*}}clang{{.*}}" "-cc1"{{.*}}"-std=c++11"
+
+// RUN: %clang -### -target x86_64-unknown-linux-gnu -offload-arch=gfx906 %s \
+// RUN:   -std=c++17 %s 2>&1 | FileCheck -check-prefixes=SPECIFIED %s
+// SPECIFIED: "{{.*}}clang{{.*}}" "-cc1"{{.*}}"-fcuda-is-device"{{.*}}"-std=c++17"
+// SPECIFIED: "{{.*}}clang{{.*}}" "-cc1"{{.*}}"-std=c++17"
+
+// RUN: %clang -### -target x86_64-pc-windows-msvc -offload-arch=gfx906 %s \
+// RUN:   2>&1 | FileCheck -check-prefixes=MSVC-DEF %s
+// MSVC-DEF: "{{.*}}clang{{.*}}" "-cc1"{{.*}}"-fcuda-is-device"{{.*}}"-std=c++14"
+// MSVC-DEF: "{{.*}}clang{{.*}}" "-cc1"{{.*}}"-std=c++14"
+
+// RUN: %clang -### -target x86_64-pc-windows-msvc -offload-arch=gfx906 %s \
+// RUN:   -std=c++17 %s 2>&1 | FileCheck -check-prefixes=MSVC-SPEC %s
+// MSVC-SPEC: "{{.*}}clang{{.*}}" "-cc1"{{.*}}"-fcuda-is-device"{{.*}}"-std=c++17"
+// MSVC-SPEC: "{{.*}}clang{{.*}}" "-cc1"{{.*}}"-std=c++17"

clang/test/Driver/hip-version.hip

+// REQUIRES: clang-driver
+// REQUIRES: x86-registered-target
+// REQUIRES: amdgpu-registered-target
+
+// RUN: %clang -v --rocm-path=%S/Inputs/rocm 2>&1 \
+// RUN:   | FileCheck -check-prefixes=FOUND %s
+
+// FOUND: Found HIP installation: {{.*Inputs.*rocm}}, version 3.6.20214-a2917cd
+
+// When --rocm-path is set and .hipVersion is not found, use default version
+
+// RUN: %clang -v --rocm-path=%S 2>&1 \
+// RUN:   | FileCheck -check-prefixes=DEFAULT %s
+
+// DEFAULT: Found HIP installation: {{.*Driver}}, version 3.5.
+
+// RUN: %clang -v --rocm-path=%S --hip-version=3.7.0 2>&1 \
+// RUN:   | FileCheck -check-prefixes=SPECIFIED %s
+
+// SPECIFIED: Found HIP installation: {{.*Driver}}, version 3.7.0
+
+// RUN: %clang -v --rocm-path=%S --hip-version=3.7 2>&1 \
+// RUN:   | FileCheck -check-prefixes=SPECIFIED2 %s
+
+// SPECIFIED2: Found HIP installation: {{.*Driver}}, version 3.7.0
+
+// RUN: not %clang -v --rocm-path=%S --hip-version=x.y 2>&1 \
+// RUN:   | FileCheck -check-prefixes=INVALID %s
+
+// INVALID: error: invalid value 'x.y' in '--hip-version=x.y'

clang/test/Driver/rocm-detect.hip

 // RUN:   | FileCheck -check-prefixes=COMMON,GFX902-DEFAULTLIBS %s
 
 
 // RUN: %clang -### -v -target x86_64-linux-gnu --cuda-gpu-arch=gfx902 -nogpulib \
 // RUN:   --rocm-path=%S/Inputs/rocm %s 2>&1 \
 // RUN:   | FileCheck -check-prefixes=COMMON,GFX902,NODEFAULTLIBS %s
 
 
-// GFX902-DEFAULTLIBS: error: cannot find device library for gfx902. Provide path to different ROCm installation via --rocm-path, or pass -nogpulib to build without linking default libraries.
+// GFX902-DEFAULTLIBS: error: cannot find ROCm device library for gfx902. Provide its path via --rocm-path or --rocm-device-lib-path, or pass -nogpulib to build without ROCm device library
 
 // NODEFAULTLIBS-NOT: error: cannot find

clang/test/Driver/rocm-detect.cl

 // RUN:   | FileCheck -check-prefixes=COMMON,GFX902-DEFAULTLIBS %s
 
 
 // RUN: %clang -### -v -target amdgcn-amd-amdhsa -mcpu=gfx902 -nogpulib \
 // RUN:   --rocm-path=%S/Inputs/rocm %s 2>&1 \
 // RUN:   | FileCheck -check-prefixes=COMMON,GFX902,NODEFAULTLIBS %s
 
 
-// GFX902-DEFAULTLIBS: error: cannot find device library for gfx902. Provide path to different ROCm installation via --rocm-path, or pass -nogpulib to build without linking default libraries.
+// GFX902-DEFAULTLIBS: error: cannot find ROCm device library for gfx902. Provide its path via --rocm-path or --rocm-device-lib-path, or pass -nogpulib to build without ROCm device library
 
 // NODEFAULTLIBS-NOT: error: cannot find

clang/test/Driver/rocm-not-found.cl

 // REQUIRES: clang-driver
 
 // Check that we raise an error if we're trying to compile OpenCL for amdhsa code but can't
 // find a ROCm install, unless -nogpulib was passed.
 
 // RUN: %clang -### --sysroot=%s/no-rocm-there -target amdgcn--amdhsa %s 2>&1 | FileCheck %s --check-prefix ERR
 // RUN: %clang -### --rocm-path=%s/no-rocm-there -target amdgcn--amdhsa %s 2>&1 | FileCheck %s --check-prefix ERR
-// ERR: cannot find ROCm installation. Provide its path via --rocm-path, or pass -nogpulib and -nogpuinc to build without ROCm device library and HIP includes.
+// ERR: cannot find ROCm device library. Provide its path via --rocm-path or --rocm-device-lib-path, or pass -nogpulib to build without ROCm device library
 
 // Accept nogpulib or nostdlib for OpenCL.
 // RUN: %clang -### -nogpulib --rocm-path=%s/no-rocm-there %s 2>&1 | FileCheck %s --check-prefix OK
 // RUN: %clang -### -nostdlib --rocm-path=%s/no-rocm-there %s 2>&1 | FileCheck %s --check-prefix OK
 // OK-NOT: cannot find ROCm installation.