tensorflow/compiler/xla/service/gpu/nvptx_compiler.cc
// Try to load textual LLVM IR from files defined in the FLAGS. If
// successful, return the llvm::Module, otherwise return nullptr.
Try to load textual LLVM IR from files defined in the FLAGS. If
successful, return the llvm::Module, otherwise return nullptr.
Try to load textual LLVM IR from files defined in the FLAGS . If successful return the llvm :: Module otherwise return nullptr .
std::unique_ptr<llvm::Module> MaybeLoadLLVMFromFile(const HloModule* module,
                                                    llvm::Module* llvm_module) {
  if (module == nullptr) {
    return nullptr;
  }
  std::string prefix = xla::FilenameFor(*module, "", "");
  auto xla_gpu_llvm_ir_file =
      module->config().debug_options().xla_gpu_llvm_ir_file();
  auto matched_filename = absl::c_find_if(
      xla_gpu_llvm_ir_file, [prefix](const std::string& full_filename) {
        return absl::StartsWith(tensorflow::io::Basename(full_filename),
                                prefix);
      });
  if (!xla_gpu_llvm_ir_file.empty() &&
      matched_filename == std::end(xla_gpu_llvm_ir_file)) {
    VLOG(1) << "RunBackend() - For module with prefix '" << prefix
            << "', we did not found a LLVM file to load.";
  }
  if (matched_filename != std::end(xla_gpu_llvm_ir_file)) {
    VLOG(1) << "RunBackend() - Will load LLVM from file: " << *matched_filename;
    llvm::LLVMContext& context = llvm_module->getContext();
    llvm::SMDiagnostic err;
    std::unique_ptr<llvm::Module> loaded_module =
        llvm::parseIRFile(*matched_filename, err, context);
    if (!loaded_module) {
      err.print("ERR", llvm::errs());
      LOG(FATAL) << "Failed to load an LLVM file. It is probably invalid LLVM.";
    }
    llvm_ir::DumpIrIfEnabled(*module, *loaded_module, false);
    return loaded_module;
  }
  return nullptr;
}
std :: unique_ptr < llvm :: Module > MaybeLoadLLVMFromFile ( const HloModule * module , llvm :: Module * llvm_module ) { if ( module == nullptr ) { return nullptr ; } std :: string prefix = xla :: FilenameFor ( * module , " " , " " ) ; auto xla_gpu_llvm_ir_file = module -> config ( ) . debug_options ( ) . xla_gpu_llvm_ir_file ( ) ; auto matched_filename = absl :: c_find_if ( xla_gpu_llvm_ir_file , [ prefix ] ( const std :: string & full_filename ) { return absl :: StartsWith ( tensorflow :: io :: Basename ( full_filename ) , prefix ) ; } ) ; if ( ! xla_gpu_llvm_ir_file . empty ( ) && matched_filename == std :: end ( xla_gpu_llvm_ir_file ) ) { VLOG ( 1 ) << " " << prefix << " " ; } if ( matched_filename != std :: end ( xla_gpu_llvm_ir_file ) ) { VLOG ( 1 ) << " " << * matched_filename ; llvm :: LLVMContext & context = llvm_module -> getContext ( ) ; llvm :: SMDiagnostic err ; std :: unique_ptr < llvm :: Module > loaded_module = llvm :: parseIRFile ( * matched_filename , err , context ) ; if ( ! loaded_module ) { err . print ( " " , llvm :: errs ( ) ) ; LOG ( FATAL ) << " " ; } llvm_ir :: DumpIrIfEnabled ( * module , * loaded_module , false ) ; return loaded_module ; } return nullptr ; }
Try to load textual LLVM IR from files defined in the FLAGS.
Try to load textual LLVM IR from files defined in the FLAGS .
// If the xla_gpu_llvm_ir_file option is set, be explicit if a file is used // and warn when a file is not used to ease catching typo in filename. // To ease comparing many LLVM versions, accept different suffixes then // the original filename. // Overwrite the dumped not optimized LLVM to show which one will be used. /*optimized=*/