1 /* Copyright 2017 The TensorFlow Authors. All Rights Reserved. 2 3 Licensed under the Apache License, Version 2.0 (the "License"); 4 you may not use this file except in compliance with the License. 5 You may obtain a copy of the License at 6 7 http://www.apache.org/licenses/LICENSE-2.0 8 9 Unless required by applicable law or agreed to in writing, software 10 distributed under the License is distributed on an "AS IS" BASIS, 11 WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. 12 See the License for the specific language governing permissions and 13 limitations 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 "absl/container/flat_hash_map.h" 20 #include "absl/strings/str_cat.h" 21 #include "llvm/IR/Module.h" 22 #include "tensorflow/compiler/xla/service/buffer_assignment.h" 23 #include "tensorflow/compiler/xla/service/hlo_instruction.h" 24 #include "tensorflow/compiler/xla/service/llvm_ir/ir_array.h" 25 #include "tensorflow/compiler/xla/types.h" 26 27 namespace xla { 28 namespace llvm_ir { 29 30 // Helper functionality used to augment the LLVM IR emitted with alias-scope 31 // metadata. 32 class AliasAnalysis { 33 public: AliasAnalysis(const HloModule & module,const BufferAssignment & assignment,llvm::LLVMContext * context)34 AliasAnalysis(const HloModule& module, const BufferAssignment& assignment, 35 llvm::LLVMContext* context) 36 : module_(module), assignment_(assignment), context_(context) {} 37 38 // Augments IrArray with aliasing information. 39 void AddAliasingInformationToIrArray(const HloInstruction& hlo, 40 llvm_ir::IrArray* array, 41 const ShapeIndex& index = {}); 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 absl::flat_hash_map<BufferAllocation::Slice, llvm::MDNode*> 80 alias_scope_metadata_; 81 82 // A map from a buffer slice to metadata corresponding to its noalias 83 // metadata. 84 absl::flat_hash_map<BufferAllocation::Slice, llvm::MDNode*> noalias_metadata_; 85 }; 86 87 } // namespace llvm_ir 88 } // namespace xla 89 90 #endif // TENSORFLOW_COMPILER_XLA_SERVICE_LLVM_IR_ALIAS_ANALYSIS_H_ 91