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