| 292 | } |
| 293 | |
| 294 | void XlaWhileOp::Compile(XlaOpKernelContext* ctx) { |
| 295 | VLOG(1) << "WhileOp::Compile"; |
| 296 | |
| 297 | // Input resources need to be grouped in the end of the body function |
| 298 | // according to the convention of the XLA bridge. |
| 299 | OP_REQUIRES_OK(ctx, VerifyResourceArgsGroupedAtEnd(ctx, body_name_attr_)); |
| 300 | |
| 301 | std::vector<XlaCompiler::Argument> arguments; |
| 302 | bool has_uninitialized_vars; |
| 303 | bool has_tensor_arrays; |
| 304 | bool has_uninitialized_tensor_lists; |
| 305 | OP_REQUIRES_OK(ctx, MakeXlaCompilerArgumentsFromInputs( |
| 306 | ctx, &arguments, &has_uninitialized_vars, |
| 307 | &has_tensor_arrays, &has_uninitialized_tensor_lists)); |
| 308 | |
| 309 | xla::XlaBuilder* builder = ctx->builder(); |
| 310 | XlaCompiler* compiler = ctx->compiler(); |
| 311 | |
| 312 | // Indices of loop vars which satisfy the following conditions: |
| 313 | // 1. They are loop invariants. |
| 314 | // 2. The op inputs at these indices are compile time constants. |
| 315 | // |
| 316 | // These compile time consts do not appear as _Args in the cond/body functions |
| 317 | // and are replaced by kConstant nodes instead. As as result, the compiled |
| 318 | // body function does not have matching input and output shape. We fix this |
| 319 | // by rewriting the body computation (see body_wrapper below) to output |
| 320 | // just the non compile-time-const values and later pad up the while output |
| 321 | // with the const args. |
| 322 | std::vector<bool> compile_time_const_arg_indices(ctx->num_inputs()); |
| 323 | int num_compile_time_const_args = 0; |
| 324 | if (propagate_compile_time_consts_) { |
| 325 | OP_REQUIRES_OK(ctx, ConvertLoopInvariantsToConst( |
| 326 | ctx, body_name_attr_, &arguments, |
| 327 | &compile_time_const_arg_indices, |
| 328 | &num_compile_time_const_args, compiler->client())); |
| 329 | } |
| 330 | |
| 331 | VLOG(1) << "Compiling body"; |
| 332 | |
| 333 | // All resource that are inputs to the loop's body must also be |
| 334 | // present as loop body outputs; the signature of the loop's input and |
| 335 | // output must match. We ensure this by asking the compiler to include the |
| 336 | // current values of all resources, even if they haven't been updated by the |
| 337 | // computation. We must also ask the compiler to keep compile-time constant |
| 338 | // outputs as part of the generated computation, for the same reason. |
| 339 | // TODO(phawkins): consider adding loop-invariant inputs to XLA's While() |
| 340 | // operator. |
| 341 | XlaCompiler::CompileOptions body_options; |
| 342 | body_options.use_tuple_arg = true; |
| 343 | body_options.return_updated_values_for_all_resources = true; |
| 344 | body_options.resolve_compile_time_constants = false; |
| 345 | body_options.is_entry_computation = false; |
| 346 | body_options.add_token_input_output = has_token_input_output_; |
| 347 | XlaCompiler::CompilationResult body; |
| 348 | OP_REQUIRES_OK(ctx, compiler->CompileFunction(body_options, body_name_attr_, |
| 349 | arguments, &body)); |
| 350 | |
| 351 | // We must use a static shape for parameters to an XLA compilation. However, |
nothing calls this directly
no test coverage detected