| 85 | } // namespace |
| 86 | |
| 87 | string GetLibdeviceDir(const HloModuleConfig& hlo_module_config) { |
| 88 | for (const string& cuda_root : CandidateCudaRoots(hlo_module_config)) { |
| 89 | string libdevice_dir = |
| 90 | tensorflow::io::JoinPath(cuda_root, "nvvm", "libdevice"); |
| 91 | VLOG(2) << "Looking for libdevice at " << libdevice_dir; |
| 92 | if (tensorflow::Env::Default()->IsDirectory(libdevice_dir).ok()) { |
| 93 | VLOG(2) << "Found libdevice dir " << libdevice_dir; |
| 94 | return libdevice_dir; |
| 95 | } |
| 96 | } |
| 97 | PrintCantFindCudaMessage( |
| 98 | "Can't find libdevice directory ${CUDA_DIR}/nvvm/libdevice. This may " |
| 99 | "result in compilation or runtime failures, if the program we try to run " |
| 100 | "uses routines from libdevice.", |
| 101 | hlo_module_config); |
| 102 | |
| 103 | // GetCudaRootCandidates always includes ".", but but if everything fails, we |
| 104 | // return it anyway. Better than returning the empty string. |
| 105 | return "."; |
| 106 | } |
| 107 | |
| 108 | Status NVPTXCompiler::OptimizeHloConvolutionCanonicalization( |
| 109 | HloModule* hlo_module, se::StreamExecutor* stream_exec, |