logical_buffer.h revision 7b401106a488eba759b6cce370393dce05d1d173
1/* Copyright 2017 The TensorFlow Authors. All Rights Reserved. 2 3Licensed under the Apache License, Version 2.0 (the "License"); 4you may not use this file except in compliance with the License. 5You may obtain a copy of the License at 6 7 http://www.apache.org/licenses/LICENSE-2.0 8 9Unless required by applicable law or agreed to in writing, software 10distributed under the License is distributed on an "AS IS" BASIS, 11WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. 12See the License for the specific language governing permissions and 13limitations under the License. 14==============================================================================*/ 15 16#ifndef TENSORFLOW_COMPILER_XLA_SERVICE_LOGICAL_BUFFER_H_ 17#define TENSORFLOW_COMPILER_XLA_SERVICE_LOGICAL_BUFFER_H_ 18 19#include <iosfwd> 20#include <string> 21#include <vector> 22 23#include "tensorflow/compiler/xla/service/hlo.pb.h" 24#include "tensorflow/compiler/xla/service/hlo_instruction.h" 25#include "tensorflow/compiler/xla/shape_util.h" 26#include "tensorflow/compiler/xla/types.h" 27#include "tensorflow/compiler/xla/xla_data.pb.h" 28#include "tensorflow/core/lib/gtl/array_slice.h" 29#include "tensorflow/core/platform/macros.h" 30#include "tensorflow/core/platform/types.h" 31 32namespace xla { 33 34// Class describing a contiguous sequence of elements (ie, C array) which form 35// the components of Shaped values in XLA. XLA arrays are trivially a 36// single LogicalBuffer. Tuple values are made up of more than one 37// LogicalBuffer: a LogicalBuffer for the pointers to elements, and a 38// LogicalBuffer for each child element. 39// 40// Every buffer is defined by a particular instruction and most instructions 41// define only a single buffer. Instructions which define a single buffer 42// include array-shaped instructions such as Add but also includes Tuple-shaped 43// instructions such as Tuple. The Tuple instruction defines a single buffer 44// which is a vector of pointers to the buffers containing the Tuple 45// instruction's operands. Though the result of the Tuple instruction includes 46// multiple buffers only the top-level buffer (the vector of pointers) is 47// defined by the Tuple instruction. The buffers containing the tuple elements 48// are defined by earlier instructions, usually the operands of the Tuple 49// instruction. 50// 51// Instructions which construct both the tuple *and* the tuple elements define 52// more than one buffer. This includes (at least) tuple-shaped Constant, 53// Parameter, Infeed and While instructions. The tuple-shaped instructions do 54// not assemble a tuple from existing buffers like the Tuple instruction does, 55// but rather define the entire tuple. 56// 57// Some instructions, such as Bitcast, define no buffers. These instructions 58// simply forward buffers from their operands. 59// 60// The LogicalBuffer object describes which HLO instruction defines a buffer and 61// where within that instruction's output shape the buffer is defined. The 62// location within the output shape is indicated by LogicalBuffer::index() which 63// is defined identically to the index used in 64// ShapeUtil::GetSubshape(). Examples: 65// 66// %add = Add(%foo, %bar) 67// %tuple_constant = Constant({1, {42, 43}}) 68// 69// %add defines a single array-shaped buffer LogicalBuffer(%add, {}) which holds 70// the array result of the add operation. The nested-tuple-shaped 71// %tuple_constant defines 5 buffers described by the following LogicalBuffer 72// objects: 73// 74// LogicalBuffer(%tuple_constant, {}) // "Top-level" buffer: vector of 75// // pointers to LogicalBuffers at 76// // indices {0} and {1} 77// LogicalBuffer(%tuple_constant, {0}) // Holds value "1" 78// LogicalBuffer(%tuple_constant, {1}) // Holds nested tuple: vector of 79// // pointers to LogicalBuffers at 80// // indices {1, 0} and {1, 1} 81// LogicalBuffer(%tuple_constant, {1, 0}) // Holds value "42" 82// LogicalBuffer(%tuple_constant, {1, 1}) // Holds value "43" 83class LogicalBuffer { 84 public: 85 // Id is a unique identifier for the LogicalBuffer to facilitate efficient 86 // collections of LogicalBuffers with stable iteration order. 87 // LogicalBuffers are typically created and accessed through 88 // TuplePointsToAnalysis, and points-to analysis assigns each LogicalBuffer a 89 // unique value. 90 using Id = int64; 91 92 // Functions which return the size and alignment of a logical buffer in bytes. 93 using SizeFunction = std::function<int64(const LogicalBuffer&)>; 94 using AlignmentFunction = std::function<int64(const LogicalBuffer&)>; 95 96 LogicalBuffer(HloInstruction* instruction, const ShapeIndex& index, Id id) 97 : instruction_(instruction), index_(index), id_(id) {} 98 99 Id id() const { return id_; } 100 101 // Return the instruction that defines the buffer. 102 HloInstruction* instruction() const { return instruction_; } 103 104 // Return the index within the output of the instruction where the buffer is 105 // defined. Index used defined as in ShapeUtil::GetSubshape() 106 const ShapeIndex& index() const { return index_; } 107 108 // Return the shape of the buffer. This reference points into the shape field 109 // of the instruction defining the buffer. Therefore, the returned shape will 110 // contain the layout of instruction, if any. 111 const Shape& shape() const { 112 return ShapeUtil::GetSubshape(instruction_->shape(), index_); 113 } 114 115 // Returns true if this buffer is the top-level output buffer of the defining 116 // HLO instruction. This is equivalent to index == {}. 117 bool IsTopLevel() const { return index_.empty(); } 118 119 // Whether this buffer contains a tuple. 120 bool IsTuple() const { return ShapeUtil::IsTuple(shape()); } 121 122 // operator< is required for std::set. 123 bool operator<(const LogicalBuffer& other) const { return id_ < other.id_; } 124 125 // Whether this buffer contains an array. 126 bool IsArray() const { return ShapeUtil::IsArray(shape()); } 127 128 string ToString() const; 129 LogicalBufferProto ToProto(const SizeFunction& size_fn) const; 130 131 // Returns the LogicalBufferProto::Location that serializes the given 132 // instruction and index. 133 static LogicalBufferProto::Location ToLocationProto( 134 const HloInstruction& instruction, const ShapeIndex& index); 135 136 private: 137 HloInstruction* instruction_; 138 ShapeIndex index_; 139 Id id_; 140 141 // Similar to HLO constructs (HloInstruction, etc), pointers are used for 142 // comparison to equality, so disable all copying. 143 TF_DISALLOW_COPY_AND_ASSIGN(LogicalBuffer); 144}; 145 146std::ostream& operator<<(std::ostream& out, const LogicalBuffer& buffer); 147 148} // namespace xla 149 150#endif // TENSORFLOW_COMPILER_XLA_SERVICE_LOGICAL_BUFFER_H_ 151