| 732 | } |
| 733 | |
| 734 | void BM_DynamicSlice(int num_iters) { |
| 735 | tensorflow::testing::StopTiming(); |
| 736 | |
| 737 | se::Platform* platform = PlatformUtil::GetDefaultPlatform().ValueOrDie(); |
| 738 | auto executors = PlatformUtil::GetStreamExecutors(platform).ValueOrDie(); |
| 739 | se::StreamExecutorMemoryAllocator allocator(platform, executors); |
| 740 | LocalClient* client = |
| 741 | ClientLibrary::GetOrCreateLocalClient(platform).ValueOrDie(); |
| 742 | auto* transfer_manager = |
| 743 | TransferManager::GetForPlatform(platform).ValueOrDie(); |
| 744 | int device_ordinal = client->default_device_ordinal(); |
| 745 | |
| 746 | XlaBuilder builder("DynamicSlice"); |
| 747 | |
| 748 | // Create input as a constant: shape [1, 2, 3, 4] |
| 749 | auto input_literal = LiteralUtil::CreateR4( |
| 750 | {{{{1, 2, 3, 4}, {5, 6, 7, 8}, {9, 10, 11, 12}}, |
| 751 | {{13, 14, 15, 16}, {17, 18, 19, 20}, {21, 22, 23, 24}}}}); |
| 752 | auto input = ConstantLiteral(&builder, input_literal); |
| 753 | |
| 754 | auto stream = |
| 755 | client->mutable_backend()->BorrowStream(device_ordinal).ValueOrDie(); |
| 756 | |
| 757 | // Create dynamic slice start indices as a parameter: shape [4] |
| 758 | auto start_indices_shape = ShapeUtil::MakeShape(S32, {}); |
| 759 | std::vector<XlaOp> start_indices(4); |
| 760 | std::vector<ScopedShapedBuffer> shaped_buffers; |
| 761 | std::vector<const Shape*> host_shapes(4); |
| 762 | for (int i = 0; i < 4; ++i) { |
| 763 | start_indices[i] = |
| 764 | Parameter(&builder, i, start_indices_shape, "start_indices"); |
| 765 | auto start_index_literal = LiteralUtil::CreateR0<int32>(i + 1); |
| 766 | // Initialize and transfer parameter buffer. |
| 767 | shaped_buffers.emplace_back( |
| 768 | client->backend() |
| 769 | .transfer_manager() |
| 770 | ->AllocateScopedShapedBuffer(start_indices_shape, &allocator, |
| 771 | /*device_ordinal=*/0) |
| 772 | .ConsumeValueOrDie()); |
| 773 | host_shapes[i] = &shaped_buffers[i].on_host_shape(); |
| 774 | ASSERT_IS_OK(transfer_manager->TransferLiteralToDevice( |
| 775 | stream.get(), start_index_literal, shaped_buffers[i])); |
| 776 | } |
| 777 | |
| 778 | // Add DynamicSlice op to the computation. |
| 779 | DynamicSlice(input, start_indices, {1, 1, 1, 1}); |
| 780 | auto computation = builder.Build().ConsumeValueOrDie(); |
| 781 | |
| 782 | TF_ASSERT_OK_AND_ASSIGN( |
| 783 | auto executables, |
| 784 | client->Compile(computation, host_shapes, ExecutableBuildOptions())); |
| 785 | auto executable = std::move(executables[0]); |
| 786 | |
| 787 | // Run some warm-up executions. |
| 788 | ExecutableRunOptions options; |
| 789 | options.set_allocator(&allocator); |
| 790 | const int kWarmups = 2; |
| 791 | std::vector<const ShapedBuffer*> shaped_buffer_ptrs; |
nothing calls this directly
no test coverage detected