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