| 28 | : Thunk(Kind::kInfeed, hlo_instruction), infeed_slices_(infeed_slices) {} |
| 29 | |
| 30 | Status InfeedThunk::ExecuteOnStream(const ExecuteParams& params) { |
| 31 | auto& stream = *params.stream; |
| 32 | auto& buffer_allocations = *params.buffer_allocations; |
| 33 | |
| 34 | VLOG(2) << "Infeeding to GPU: " << hlo_instruction()->ToString(); |
| 35 | |
| 36 | auto op_profiler = |
| 37 | params.profiler->MakeScopedInstructionProfiler(hlo_instruction()); |
| 38 | ShapeTree<InfeedBuffer> infeed_buffers = |
| 39 | GetOrCreateInfeedManager()->BlockingGetNextDestination(); |
| 40 | |
| 41 | // infeed_slices_'s shape should be a tuple of shape (buffers, token). |
| 42 | const auto& infeed_shape = infeed_slices_.shape(); |
| 43 | TF_RET_CHECK(infeed_shape.IsTuple()) |
| 44 | << ShapeUtil::HumanStringWithLayout(infeed_shape); |
| 45 | TF_RET_CHECK(infeed_shape.tuple_shapes().size() == 2) |
| 46 | << ShapeUtil::HumanStringWithLayout(infeed_shape); |
| 47 | TF_RET_CHECK(infeed_shape.tuple_shapes(1).IsToken()) |
| 48 | << ShapeUtil::HumanStringWithLayout(infeed_shape); |
| 49 | TF_RET_CHECK( |
| 50 | ShapeUtil::Equal(infeed_buffers.shape(), infeed_shape.tuple_shapes(0))) |
| 51 | << "Expected infeed of shape " |
| 52 | << ShapeUtil::HumanStringWithLayout(infeed_shape.tuple_shapes(0)) |
| 53 | << " but was " |
| 54 | << ShapeUtil::HumanStringWithLayout(infeed_buffers.shape()); |
| 55 | |
| 56 | { |
| 57 | // The infeed buffer has an extra outer tuple with a token. Adjust the index |
| 58 | // accordingly. |
| 59 | ShapeIndex index = {0}; |
| 60 | std::function<void(std::vector<void*>*)> copy_tuple_contents = |
| 61 | [&](std::vector<void*>* tuple_element_addresses) { |
| 62 | const Shape& shape = ShapeUtil::GetSubshape(infeed_buffers.shape(), |
| 63 | ShapeIndexView(index, 1)); |
| 64 | // For the leaf buffers of the tuple copy the elements directly. |
| 65 | if (shape.IsArray()) { |
| 66 | const BufferAllocation::Slice& tuple_element_buffer = |
| 67 | infeed_slices_.element(index); |
| 68 | se::DeviceMemoryBase tuple_element_address = |
| 69 | buffer_allocations.GetDeviceAddress(tuple_element_buffer); |
| 70 | |
| 71 | InfeedBuffer* buffer = |
| 72 | infeed_buffers.mutable_element(ShapeIndexView(index, 1)); |
| 73 | stream.ThenMemcpy(&tuple_element_address, |
| 74 | *(buffer->device_memory()), buffer->length()); |
| 75 | tuple_element_addresses->push_back(tuple_element_address.opaque()); |
| 76 | return; |
| 77 | } |
| 78 | |
| 79 | const int64 tuple_element_count = ShapeUtil::TupleElementCount(shape); |
| 80 | index.push_back(0); |
| 81 | std::vector<void*> inner_tuple_element_addresses; |
| 82 | for (int64 i = 0; i < tuple_element_count; ++i) { |
| 83 | index.back() = i; |
| 84 | copy_tuple_contents(&inner_tuple_element_addresses); |
| 85 | } |
| 86 | index.pop_back(); |
| 87 |
nothing calls this directly
no test coverage detected