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_LLVM_IR_ALIAS_ANALYSIS_H_
17#define TENSORFLOW_COMPILER_XLA_SERVICE_LLVM_IR_ALIAS_ANALYSIS_H_
18
19#include "llvm/IR/Module.h"
20#include "tensorflow/compiler/xla/service/buffer_assignment.h"
21#include "tensorflow/compiler/xla/service/hlo_instruction.h"
22#include "tensorflow/compiler/xla/service/llvm_ir/ir_array.h"
23#include "tensorflow/compiler/xla/types.h"
24#include "tensorflow/core/lib/gtl/flatmap.h"
25#include "tensorflow/core/lib/gtl/flatset.h"
26#include "tensorflow/core/lib/strings/strcat.h"
27
28namespace xla {
29namespace llvm_ir {
30
31// Helper functionality used to augment the LLVM IR emitted with alias-scope
32// metadata.
33class AliasAnalysis {
34 public:
35  AliasAnalysis(const HloModule& module, const BufferAssignment& assignment,
36                llvm::LLVMContext* context)
37      : module_(module), assignment_(assignment), context_(context) {}
38
39  // Augments IrArray with aliasing information.
40  void AddAliasingInformationToIrArray(const HloInstruction& hlo,
41                                       llvm_ir::IrArray* array);
42
43 private:
44  // Returns a unique alias domain for this emitter.
45  llvm::MDNode* GetAliasDomain();
46
47  // Returns an alias.scope metadata node corresponding to a given buffer slice.
48  llvm::MDNode* GetAliasScopeMetadataForBuffer(
49      const BufferAllocation::Slice& buffer_slice, llvm::MDNode* domain);
50
51  // Returns a noalias metadata node corresponding to a given buffer slice.
52  //
53  // |buffer_slice| is the buffer slice.
54  //
55  // |domain| corresponds to the alias scope domain as documented at
56  // http://llvm.org/docs/LangRef.html#noalias-and-alias-scope-metadata
57  //
58  // |hlo| is the instruction we are computing a noalias set for.
59  llvm::MDNode* GetNoaliasMetadataForBuffer(
60      const BufferAllocation::Slice& buffer_slice, llvm::MDNode* domain,
61      const BufferAssignment& assignment, const HloInstruction& hlo);
62
63  // The HLO module we are compiling for.
64  const HloModule& module_;
65
66  // Assignment of the temporary buffers needed by the computation and their
67  // shape information.
68  const BufferAssignment& assignment_;
69
70  // The LLVM context which we are using for IR emission.
71  llvm::LLVMContext* context_;
72
73  // Holds the alias domain for this computation.
74  llvm::MDNode* alias_domain_ = nullptr;
75
76  // A map from a buffer slice to metadata corresponding to its alias.scope
77  // metadata.  The index kParameterAliasSet is used to hold aliasing
78  // information for parameters.
79  tensorflow::gtl::FlatMap<BufferAllocation::Slice, llvm::MDNode*,
80                           BufferAllocation::Slice::Hasher>
81      alias_scope_metadata_;
82
83  // A map from a buffer slice to metadata corresponding to its noalias
84  // metadata.
85  tensorflow::gtl::FlatMap<BufferAllocation::Slice, llvm::MDNode*,
86                           BufferAllocation::Slice::Hasher>
87      noalias_metadata_;
88};
89
90}  // namespace llvm_ir
91}  // namespace xla
92
93#endif  // TENSORFLOW_COMPILER_XLA_SERVICE_LLVM_IR_ALIAS_ANALYSIS_H_
94