message: [XLA] Print out directories we looked in for ptxas/libdevice when we can't find them.
  This way users know what where we're looking.
  Also inform users that they can pass a flag to cause us to look in a specific directory.
diff: @@ -118,6 +118,9 @@ std::vector<string> GetCudaRootCandidates(
  const HloModuleConfig& hlo_module_config) {
  std::vector<string> potential_cuda_roots = tensorflow::CandidateCudaRoots();
  + // "." is our last resort, even though it probably won't work.
  + potential_cuda_roots.push_back(".");
  +
  // CUDA location explicitly specified by user via --xla_gpu_cuda_data_dir has
  // highest priority.
  string xla_gpu_cuda_data_dir =
  @@ -129,9 +132,23 @@ std::vector<string> GetCudaRootCandidates(
  return potential_cuda_roots;
  }
  +void PrintCantFindCudaMessage(absl::string_view msg,
  + const HloModuleConfig& hlo_module_config) {
  + LOG(WARNING) << msg;
  + LOG(WARNING) << "Searched in the following directories:";
  + for (const auto& dir : GetCudaRootCandidates(hlo_module_config)) {
  + LOG(WARNING) << " " << dir;
  + }
  + LOG(WARNING)
  + << "You can choose the search directory by setting xla_gpu_cuda_data_dir "
  + "in HloModule's DebugOptions. For most apps, setting the environment "
  + "variable XLA_FLAGS=--xla_gpu_cuda_data_dir=/path/to/cuda will work.";
  +}
  +
  // Returns the directory containing nvvm libdevice files.
  string GetLibdeviceDir(const HloModuleConfig& hlo_module_config) {
  - for (const string& cuda_root : GetCudaRootCandidates(hlo_module_config)) {
  + const auto& candidate_dirs = GetCudaRootCandidates(hlo_module_config);
  + for (const string& cuda_root : candidate_dirs) {
  string libdevice_dir =
  tensorflow::io::JoinPath(cuda_root, "nvvm", "libdevice");
  VLOG(2) << "Looking for libdevice at " << libdevice_dir;
  @@ -140,8 +157,14 @@ string GetLibdeviceDir(const HloModuleConfig& hlo_module_config) {
  return libdevice_dir;
  }
  }
  - LOG(WARNING) << "Unable to find libdevice dir. Using '.'";
  - // Last resort: maybe in the current folder.
  + PrintCantFindCudaMessage(
  + "Can't find directory containing CUDA libevice. This may result in "
  + "compilation or runtime failures, if the program we try to run uses "
  + "routines from libdevice.",
  + hlo_module_config);
  +
  + // GetCudaRotCandidates always inclues ".", but but if everything fails, we
  + // return it anyway. Better than returning the empty string.
  return ".";
  }
  @@ -843,10 +866,11 @@ std::vector<uint8> NVPTXCompiler::CompilePtxOrGetCachedResult(
  log_warning = !warning_done.exchange(true);
  }
  if (log_warning) {
  - LOG(WARNING)
  - << "Failed to compile ptx to cubin. Will attempt to let "
  - "GPU driver compile the ptx. "
  - << maybe_cubin.status();
  + PrintCantFindCudaMessage(
  + "Can't find ptxas binary. Will back to the GPU driver "
  + "for PTX -> sass compilation. This is OK so long as you don't "
  + "see a warning below about an out-of-date driver version.",
  + hlo_module_config);
  }
  // We're going to use the driver to JIT our PTX->SASS, so warn if
  