| 22 | |
| 23 | namespace xla { |
| 24 | StatusOr<std::vector<std::unique_ptr<Executable>>> LLVMCompiler::Compile( |
| 25 | std::unique_ptr<HloModuleGroup> module_group, |
| 26 | std::vector<std::vector<se::StreamExecutor*>> stream_execs, |
| 27 | se::DeviceMemoryAllocator* device_allocator) { |
| 28 | // Tensorflow tries to enable the following behaviors in all its threads: |
| 29 | // |
| 30 | // - Denormals are zero (DAZ): roughly, operations treat denormal floats as |
| 31 | // zero. |
| 32 | // - Flush denormals to zero (FTZ): roughly, operations produce zero instead |
| 33 | // of denormal floats. |
| 34 | // |
| 35 | // In theory enabling these shouldn't matter since the compiler should ideally |
| 36 | // not leak its environment into generated code, but we turn off DAZ and FTZ |
| 37 | // to get some defense-in-depth. |
| 38 | tensorflow::port::ScopedDontFlushDenormal dont_flush_denormals; |
| 39 | |
| 40 | std::vector<std::unique_ptr<Executable>> result; |
| 41 | std::vector<std::unique_ptr<HloModule>> modules = |
| 42 | module_group->ConsumeModules(); |
| 43 | for (size_t i = 0; i < modules.size(); i++) { |
| 44 | if (stream_execs[i].size() != 1) { |
| 45 | return Unimplemented( |
| 46 | "Model partitioning not implemented for the CPU/GPU compilers!"); |
| 47 | } |
| 48 | |
| 49 | TF_ASSIGN_OR_RETURN(modules[i], |
| 50 | RunHloPasses(std::move(modules[i]), stream_execs[i][0], |
| 51 | device_allocator)); |
| 52 | TF_ASSIGN_OR_RETURN(std::unique_ptr<Executable> executable, |
| 53 | RunBackend(std::move(modules[i]), stream_execs[i][0], |
| 54 | device_allocator)); |
| 55 | result.push_back(std::move(executable)); |
| 56 | } |
| 57 | |
| 58 | return {std::move(result)}; |
| 59 | } |
| 60 | } // namespace xla |
nothing calls this directly
no test coverage detected