| 2929 | } |
| 2930 | |
| 2931 | Status IrEmitter::FinishVisit(HloInstruction* root) { |
| 2932 | // When this method is called, we should have already emitted an IR value for |
| 2933 | // the root (return) op. The IR value holds the address of the buffer holding |
| 2934 | // the value. If the root is a constant or parameter, we perform a memcpy from |
| 2935 | // this buffer to the retval buffer of the computation. Otherwise, there's |
| 2936 | // nothing to do since the result was already written directly into the output |
| 2937 | // buffer. |
| 2938 | VLOG(2) << "FinishVisit root: " << root->ToString(); |
| 2939 | if (root->opcode() == HloOpcode::kOutfeed) { |
| 2940 | VLOG(2) << " outfeed with value: " |
| 2941 | << llvm_ir::DumpToString(*GetEmittedValueFor(root->operand(0))); |
| 2942 | } else { |
| 2943 | VLOG(2) << " value: " << llvm_ir::DumpToString(*GetEmittedValueFor(root)); |
| 2944 | } |
| 2945 | |
| 2946 | auto record_complete_computation = [&](llvm::Value* prof_counter) { |
| 2947 | if (prof_counter) { |
| 2948 | profiling_state_.RecordCompleteComputation(&b_, prof_counter); |
| 2949 | } |
| 2950 | }; |
| 2951 | |
| 2952 | // For the entry computation this increment is cumulative of embedded |
| 2953 | // computations since it includes cycles spent in computations invoked by |
| 2954 | // While, Call etc. |
| 2955 | record_complete_computation(GetProfileCounterFor(*root->parent())); |
| 2956 | return Status::OK(); |
| 2957 | } |
| 2958 | |
| 2959 | template <typename T> |
| 2960 | llvm::Value* IrEmitter::GetProfileCounterCommon( |
no test coverage detected