Invokes the given computation passing arbitrary data for every (unbound) parameter if use_fake_data, Otherwise use recorded data if available. Similarly, infeeds fake data of shape fake_infeed_shape if it is provided. If generate_fake_infeed is true, the required infeed shape is derived from the computation and then used to provide a fake infeed shape. If neither generate_fake_infeed is true nor
| 219 | // If neither generate_fake_infeed is true nor a fake_infeed_shape is provided, |
| 220 | // no infeed is performed. |
| 221 | StatusOr<Literal> ReplayComputation(const HloSnapshot& module, |
| 222 | LocalExecutable* executable, |
| 223 | LocalClient* client, const Options& opts) { |
| 224 | XlaComputation computation(module.hlo().hlo_module()); |
| 225 | |
| 226 | // Build the `argument_ptrs` vector, which contains ShapedBuffer*s to our |
| 227 | // arguments. This is a bit involved, because we may have to convert from |
| 228 | // GlobalData to ShapedBuffer*, and we have to manage the lifetime of all our |
| 229 | // objects. |
| 230 | std::vector<ScopedShapedBuffer> scoped_shaped_buffer_arguments; |
| 231 | std::vector<std::unique_ptr<GlobalData>> global_data_arguments; |
| 232 | std::vector<const ShapedBuffer*> argument_ptrs; |
| 233 | if (opts.use_fake_data) { |
| 234 | // Run fake computations with debug options ignoring XLA_FLAGS. Users very |
| 235 | // likely want XLA_FLAGS only to apply to the "real" computation being run, |
| 236 | // not to the fake computations we use for generating arguments. |
| 237 | auto debug_opts_flags = GetDebugOptionsFromFlags(); |
| 238 | auto debug_opts = DefaultDebugOptionsIgnoringFlags(); |
| 239 | |
| 240 | // ptxas can be called during the generation of fake data. |
| 241 | // As it is cached, we want it to not ignore this flag. |
| 242 | debug_opts.set_xla_gpu_asm_extra_flags(debug_opts_flags.xla_gpu_asm_extra_flags()); |
| 243 | |
| 244 | global_data_arguments = |
| 245 | MakeFakeArgumentsOrDie(computation, client, &debug_opts); |
| 246 | for (const auto& data : global_data_arguments) { |
| 247 | argument_ptrs.push_back( |
| 248 | client->GlobalDataToShapedBuffer(data->handle(), /*device_ordinal=*/0) |
| 249 | .ValueOrDie()); |
| 250 | } |
| 251 | } else { // use recorded data if available |
| 252 | for (const auto& proto : module.arguments()) { |
| 253 | TF_ASSIGN_OR_RETURN(Literal literal, Literal::CreateFromProto(proto)); |
| 254 | TF_ASSIGN_OR_RETURN( |
| 255 | ScopedShapedBuffer data, |
| 256 | client->LiteralToShapedBuffer(literal, /*device_ordinal=*/0)); |
| 257 | scoped_shaped_buffer_arguments.push_back(std::move(data)); |
| 258 | } |
| 259 | for (const auto& argument : scoped_shaped_buffer_arguments) { |
| 260 | argument_ptrs.push_back(&argument); |
| 261 | } |
| 262 | } |
| 263 | |
| 264 | if (absl::optional<Shape> infeed_shape = GetXfeedShape( |
| 265 | /*is_infeed=*/true, computation.proto(), opts)) { |
| 266 | auto infeed_data = std::make_shared<Literal>( |
| 267 | std::move(MakeFakeLiteral(*infeed_shape)).ValueOrDie()); |
| 268 | xla::gpu::GetOrCreateInfeedManager() |
| 269 | ->RegisterBeforeGetNextDestinationCallback([infeed_data, client] { |
| 270 | TF_CHECK_OK(client->TransferToInfeed(*infeed_data)); |
| 271 | }); |
| 272 | } |
| 273 | |
| 274 | absl::optional<tensorflow::thread::ThreadPool> outfeed_thread_pool; |
| 275 | if (absl::optional<Shape> outfeed_shape = GetXfeedShape( |
| 276 | /*is_infeed=*/false, computation.proto(), opts)) { |
| 277 | // For each an outfeed that runs, enqueue a task that will consume it. We |
| 278 | // need a thread pool because the act of running an outfeed blocks on there |
no test coverage detected