| 211 | } |
| 212 | |
| 213 | xla::StatusOr<xla::XlaComputation> BuildWrappedBody( |
| 214 | XlaOpKernelContext* ctx, const XlaCompiler::CompilationResult& body, |
| 215 | const std::vector<bool>& compile_time_const_arg_indices, |
| 216 | int num_compile_time_const_args, bool has_token_input_output) { |
| 217 | if (num_compile_time_const_args <= 0) { |
| 218 | return xla::XlaComputation(body.computation->proto()); |
| 219 | } |
| 220 | xla::XlaComputation body_wrapper; |
| 221 | std::unique_ptr<xla::XlaBuilder> cb = |
| 222 | ctx->builder()->CreateSubBuilder("body_wrapper"); |
| 223 | xla::Shape body_input_shape = body.xla_input_shapes[0]; |
| 224 | auto inputs = xla::Parameter(cb.get(), 0, body_input_shape, "inputs"); |
| 225 | // Call the original body function which has mismatched inputs and outputs |
| 226 | // and strip the compile time consts from the list of outputs. While requires |
| 227 | // the inputs and outputs of its body function to match. |
| 228 | auto outputs = xla::Call(cb.get(), *body.computation, {inputs}); |
| 229 | std::vector<xla::XlaOp> non_compile_time_const_outputs; |
| 230 | for (int i = 0; i < compile_time_const_arg_indices.size(); i++) { |
| 231 | if (!compile_time_const_arg_indices[i]) { |
| 232 | non_compile_time_const_outputs.push_back( |
| 233 | xla::GetTupleElement(outputs, i)); |
| 234 | } |
| 235 | } |
| 236 | // If `body` has a token output, append it to |
| 237 | // `non_compile_time_const_outputs`. |
| 238 | if (has_token_input_output) { |
| 239 | non_compile_time_const_outputs.push_back( |
| 240 | xla::GetTupleElement(outputs, ctx->num_outputs())); |
| 241 | } |
| 242 | xla::Tuple(cb.get(), non_compile_time_const_outputs); |
| 243 | return cb->Build(); |
| 244 | } |
| 245 | |
| 246 | xla::XlaOp BuildWhile(XlaOpKernelContext* ctx, |
| 247 | const xla::XlaComputation& wrapped_cond, |
no test coverage detected