Home
last modified time | relevance | path

Searched refs:llvm_ir (Results 1 – 25 of 74) sorted by relevance

123

/external/tensorflow/tensorflow/compiler/xla/service/gpu/
Dir_emitter_unnested.h135 const llvm_ir::IrArray::Index& index, llvm::Value* y_loc,
142 const ThreadIdInfo& thread_id_info, const llvm_ir::IrArray::Index& index,
228 const llvm_ir::ElementGenerator& body_emitter) override;
234 const HloInstruction& hlo, const llvm_ir::ElementGenerator& body_emitter,
355 return llvm_ir::ByteSizeOf( in ByteSizeOf()
366 absl::Span<const llvm_ir::IrArray> result_ir_arrays,
367 const llvm_ir::IrArray::Index& index, bool use_linear_index,
368 absl::Span<const std::pair<llvm_ir::ElementGenerator, int>>
458 absl::Span<const llvm_ir::IrArray> ir_arrays,
459 const llvm_ir::IrArray::Index& index);
[all …]
Dparallel_loop_emitter.cc43 const llvm_ir::ElementGenerator& target_element_generator, in ParallelLoopEmitter()
44 absl::Span<const llvm_ir::IrArray> target_arrays, in ParallelLoopEmitter()
52 const llvm_ir::ElementGenerator& target_element_generator, in ParallelLoopEmitter()
53 const llvm_ir::IrArray& target_array, in ParallelLoopEmitter()
60 std::vector<llvm_ir::IrArray::Index>
77 std::vector<llvm_ir::IrArray::Index> array_indices; in EmitIndexAndSetExitBasicBlock()
80 llvm_ir::AddRangeMetadata(0, launch_dimensions_.block_counts().x, in EmitIndexAndSetExitBasicBlock()
90 llvm_ir::AddRangeMetadata(0, launch_dimensions_.thread_counts_per_block().x, in EmitIndexAndSetExitBasicBlock()
111 llvm_ir::EmitCallToIntrinsic( in EmitIndexAndSetExitBasicBlock()
142 auto if_in_bounds = llvm_ir::EmitIfThenElse( in EmitIndexAndSetExitBasicBlock()
[all …]
Dir_emitter_nested.cc65 llvm_ir::ShapeToIrType(param_shape, module_)->getPointerTo()); in CodegenNestedComputation()
67 llvm_ir::ByteSizeOf(param_shape, module_->getDataLayout()); in CodegenNestedComputation()
75 llvm_ir::ShapeToIrType(root_shape, module_)->getPointerTo()); in CodegenNestedComputation()
76 int64 root_size = llvm_ir::ByteSizeOf( in CodegenNestedComputation()
87 llvm_ir::SanitizeFunctionName( in CodegenNestedComputation()
139 llvm::Type* tuple_type = llvm_ir::ShapeToIrType(return_shape, module_); in CodegenNestedComputation()
146 llvm_ir::EmitGetTupleElement(element_shape, in CodegenNestedComputation()
150 llvm_ir::EmitGetTupleElement(element_shape, in CodegenNestedComputation()
168 const llvm_ir::ElementGenerator& element_generator) { in EmitTargetElementLoop()
172 std::vector<llvm_ir::IrArray> target_arrays = in EmitTargetElementLoop()
[all …]
Dhlo_to_ir_bindings.cc85 llvm_ir::ConstantHloToGlobalName(*non_io_hlo)); in EmitBasePointersForHlos()
87 << llvm_ir::ConstantHloToGlobalName(*non_io_hlo); in EmitBasePointersForHlos()
91 llvm_ir::ShapeToIrType(non_io_hlo->shape(), module_); in EmitBasePointersForHlos()
93 llvm_ir::EmitAllocaAtFunctionEntry( in EmitBasePointersForHlos()
105 return llvm_ir::EmitGetTupleElement( in EmitGetTupleElement()
109 return llvm_ir::EmitGetTupleElement( in EmitGetTupleElement()
125 llvm_ir::ShapeToIrType(shape, b->GetInsertBlock()->getModule()); in CastToTypedValue()
146 ir_value->setName(llvm_ir::IrName(&hlo, "raw")); in GetTypedIrValue()
149 typed_ir_value->setName(llvm_ir::IrName(&hlo, "typed")); in GetTypedIrValue()
191 llvm_ir::IrArray HloToIrBindings::GetIrArray(const HloInstruction& hlo, in GetIrArray()
[all …]
Dparallel_loop_emitter.h31 class ParallelLoopEmitter : public llvm_ir::LoopEmitter {
40 ParallelLoopEmitter(const llvm_ir::ElementGenerator& target_element_generator,
41 const llvm_ir::IrArray& target_array,
50 ParallelLoopEmitter(const llvm_ir::ElementGenerator& target_element_generator,
51 absl::Span<const llvm_ir::IrArray> target_arrays,
59 std::vector<llvm_ir::IrArray::Index> EmitIndexAndSetExitBasicBlock(
Dir_emitter.cc74 using llvm_ir::IrName;
75 using llvm_ir::SetToFirstInsertPoint;
93 operand_to_generator[operand] = [=](const llvm_ir::IrArray::Index& index) { in DefaultAction()
116 ? llvm_ir::ConvertLiteralToIrConstant(literal, module_) in EmitConstants()
131 unsigned global_address_space = llvm_ir::GetGlobalMemoryAddressSpace( in EmitConstants()
134 std::string global_name = llvm_ir::ConstantHloToGlobalName(*instr); in EmitConstants()
187 llvm_ir::EmitGetTupleElement( in HandleGetTupleElement()
220 llvm_ir::EmitTuple(GetIrArray(*tuple, *tuple), base_ptrs, &b_); in HandleTuple()
400 int element_size = llvm_ir::GetSizeInBits(element_type); in EmitAtomicOperationUsingCAS()
410 llvm::Value* cas_old_output_address = llvm_ir::EmitAllocaAtFunctionEntry( in EmitAtomicOperationUsingCAS()
[all …]
Dir_emitter.h73 std::function<std::vector<llvm_ir::IrArray>()>;
125 llvm_ir::IrArray GetIrArray(const HloInstruction& inst,
138 std::vector<llvm_ir::IrArray> ConstructIrArrayForOutputs(
147 const llvm_ir::ElementGenerator& body_emitter) = 0;
207 const llvm_ir::IrArray::Index& keys_index,
208 const llvm_ir::IrArray::Index& compare_keys_index,
209 const llvm_ir::IrArray& keys_array);
Dir_emitter_unnested.cc138 using llvm_ir::IrArray;
139 using llvm_ir::IrName;
563 llvm_ir::SanitizeFunctionName(std::string(name))); in BuildKernelPrototype()
779 llvm_ir::GetGlobalMemoryAddressSpace(*ir_emitter_context_->llvm_module()); in EmitConstant()
820 std::vector<llvm_ir::IrArray> ir_arrays; in EmitPadToStaticFromMlir()
826 const llvm_ir::IrArray source_array = ir_arrays[0]; in EmitPadToStaticFromMlir()
827 const llvm_ir::IrArray output_array = ir_arrays[1]; in EmitPadToStaticFromMlir()
829 absl::Span<const llvm_ir::IrArray>(ir_arrays).subspan(2); in EmitPadToStaticFromMlir()
906 llvm_ir::LoopEmitter::BodyEmitter body_generator = in EmitPadToStaticFromMlir()
907 [&](const llvm_ir::IrArray::Index& array_index) -> Status { in EmitPadToStaticFromMlir()
[all …]
Delemental_ir_emitter.cc56 using llvm_ir::IrArray;
57 using llvm_ir::IrName;
58 using llvm_ir::SetToFirstInsertPoint;
171 return llvm_ir::EmitCallToIntrinsic( in EmitFloatBinaryOp()
289 llvm_ir::EmitCallToIntrinsic(llvm::Intrinsic::fabs, {input}, {type}, b()); in EmitTanh()
291 llvm::Value* fast_tanh = llvm_ir::EmitFastTanh(b(), input); in EmitTanh()
293 auto one_with_sign = llvm_ir::EmitCallToIntrinsic(llvm::Intrinsic::copysign, in EmitTanh()
/external/tensorflow/tensorflow/compiler/xla/service/cpu/
Ddot_op_emitter.cc56 using llvm_ir::SetToFirstInsertPoint;
125 const llvm_ir::IrArray& target_array,
126 const llvm_ir::IrArray& lhs_array,
127 const llvm_ir::IrArray& rhs_array,
128 const llvm_ir::IrArray* addend_array,
223 const llvm_ir::IrArray& target_array_;
224 const llvm_ir::IrArray& lhs_array_;
225 const llvm_ir::IrArray& rhs_array_;
226 const llvm_ir::IrArray* addend_array_;
236 DotInfo dot_info, string dot_hlo_name, const llvm_ir::IrArray& target_array, in DotOpEmitter()
[all …]
Dparallel_loop_emitter.cc26 const llvm_ir::ElementGenerator& target_element_generator, in ParallelLoopEmitter()
27 const llvm_ir::IrArray& target_array, in ParallelLoopEmitter()
32 std::vector<llvm_ir::IrArray::Index>
46 llvm_ir::ForLoopNest loop_nest(loop_name, b_); in EmitIndexAndSetExitBasicBlock()
60 std::unique_ptr<llvm_ir::ForLoop> loop = loop_nest.AddLoop( in EmitIndexAndSetExitBasicBlock()
66 std::unique_ptr<llvm_ir::ForLoop> loop = loop_nest.AddLoop( in EmitIndexAndSetExitBasicBlock()
74 llvm_ir::SetToFirstInsertPoint(loop_nest.GetInnerLoopBodyBasicBlock(), b_); in EmitIndexAndSetExitBasicBlock()
80 llvm_ir::IrArray::Index array_index(array_multi_index, shape_, index_type); in EmitIndexAndSetExitBasicBlock()
Dir_emitter.cc85 using llvm_ir::IrName;
86 using llvm_ir::SetToFirstInsertPoint;
110 b_.setFastMathFlags(llvm_ir::GetCpuFastMathFlags(hlo_module_config_)); in IrEmitter()
120 llvm_ir::IrArray root_value = GetIrArrayFor(computation->root_instruction()); in EmitThreadLocalFunctionEpilogue()
131 llvm::Type* tuple_type = llvm_ir::ShapeToIrType(return_shape, module_); in EmitThreadLocalFunctionEpilogue()
137 llvm::Value* destination = llvm_ir::EmitGetTupleElement( in EmitThreadLocalFunctionEpilogue()
143 llvm::Value* source = llvm_ir::EmitGetTupleElement( in EmitThreadLocalFunctionEpilogue()
236 llvm_ir::ConvertLiteralToIrConstant(literal, module_); in EmitGlobalForLiteral()
257 const Literal& literal = llvm_ir::LiteralForConstantAllocation(allocation); in EmitConstantGlobals()
311 return llvm_ir::ByteSizeOf(shape, module_->getDataLayout()); in ByteSizeOf()
[all …]
Ddot_op_emitter.h62 const llvm_ir::IrArray& target_array,
63 const llvm_ir::IrArray& lhs_array,
64 const llvm_ir::IrArray& rhs_array,
65 const llvm_ir::IrArray* addend_array,
Dparallel_loop_emitter.h48 class ParallelLoopEmitter : public llvm_ir::LoopEmitter {
54 ParallelLoopEmitter(const llvm_ir::ElementGenerator& target_element_generator,
55 const llvm_ir::IrArray& target_array,
63 std::vector<llvm_ir::IrArray::Index> EmitIndexAndSetExitBasicBlock(
DBUILD223 "//tensorflow/compiler/xla/service/llvm_ir:llvm_util",
396 "//tensorflow/compiler/xla/service/llvm_ir:alias_analysis",
397 "//tensorflow/compiler/xla/service/llvm_ir:buffer_assignment_util",
398 "//tensorflow/compiler/xla/service/llvm_ir:dynamic_update_slice_util",
399 "//tensorflow/compiler/xla/service/llvm_ir:fused_ir_emitter",
400 "//tensorflow/compiler/xla/service/llvm_ir:ir_array",
401 "//tensorflow/compiler/xla/service/llvm_ir:ir_builder_mixin",
402 "//tensorflow/compiler/xla/service/llvm_ir:llvm_loop",
403 "//tensorflow/compiler/xla/service/llvm_ir:llvm_util",
404 "//tensorflow/compiler/xla/service/llvm_ir:loop_emitter",
[all …]
Dir_emitter.h69 std::function<std::vector<llvm_ir::IrArray>()>;
225 llvm_ir::IrArray GetIrArrayFor(const HloInstruction* hlo);
228 std::vector<llvm_ir::IrArray> GetIrArraysForOperandsOf(
237 llvm_ir::IrArray* array) { in AddAliasingInformationToIrArray()
313 const llvm_ir::ElementGenerator& element_generator);
316 const llvm_ir::ElementGenerator& element_generator);
383 const llvm_ir::IrArray& containing_array);
399 const llvm_ir::IrArray::Index& output_index,
415 const llvm_ir::IrArray& target_array,
416 const llvm_ir::IrArray& source_array);
[all …]
/external/tensorflow/tensorflow/compiler/xla/service/
Delemental_ir_emitter.cc54 using llvm_ir::IrArray;
55 using llvm_ir::IrName;
56 using llvm_ir::SetToFirstInsertPoint;
225 llvm_ir::PrimitiveTypeToIrType(to_type, module)); in EmitIntegralToFloating()
230 llvm_ir::PrimitiveTypeToIrType(to_type, module)); in EmitIntegralToFloating()
263 llvm_ir::PrimitiveTypeToIrType(PRED, module_)); in EmitIntegerUnaryOp()
267 llvm_ir::PrimitiveTypeToIrType(to_type, module_), in EmitIntegerUnaryOp()
280 auto to_ir_component_type = llvm_ir::PrimitiveTypeToIrType( in EmitIntegerUnaryOp()
306 llvm_ir::PrimitiveTypeToIrType(to_type, module_)); in EmitIntegerUnaryOp()
320 llvm_ir::PrimitiveTypeToIrType(op->shape().element_type(), module_); in EmitIntegerUnaryOp()
[all …]
Delemental_ir_emitter.h39 std::unordered_map<const HloInstruction*, llvm_ir::ElementGenerator>;
48 llvm_ir::ElementGenerator MakeElementGenerator(
206 const llvm_ir::IrArray::Index& index);
211 const llvm_ir::IrArray::Index& index);
216 const llvm_ir::IrArray::Index& target_index);
221 const llvm_ir::IrArray::Index& index);
226 const llvm_ir::IrArray::Index& index);
231 const llvm_ir::IrArray::Index& index);
236 const llvm_ir::IrArray::Index& padded_index);
241 const llvm_ir::IrArray::Index& dot_result_index);
[all …]
/external/tensorflow/tensorflow/compiler/xla/service/cpu/tests/
Dcpu_noalias_test.cc73 llvm_ir::AliasAnalysis aa(*hlo_module, *status_or_buffer_assn.ValueOrDie(), in TEST_F()
93 llvm_ir::IrArray param_x_array(param_x_val, param_shape); in TEST_F()
95 llvm_ir::IrArray::Index zero_2d({zero, zero}, param_shape, zero->getType()); in TEST_F()
104 llvm_ir::IrArray concat1_array(concat1_val, shape); in TEST_F()
106 llvm_ir::IrArray::Index zero_2d({zero, zero}, shape, zero->getType()); in TEST_F()
115 llvm_ir::IrArray concat2_array(concat2_val, shape); in TEST_F()
117 llvm_ir::IrArray::Index zero_2d({zero, zero}, shape, zero->getType()); in TEST_F()
125 llvm_ir::IrArray add_array(concat2_val, shape); in TEST_F()
127 llvm_ir::IrArray::Index zero_2d({zero, zero}, shape, zero->getType()); in TEST_F()
146 RunFileCheck(llvm_ir::DumpModuleToString(ir_module), filecheck_pattern)); in TEST_F()
/external/tensorflow/tensorflow/compiler/xla/service/llvm_ir/
Dkernel_support_library.cc45 std::unique_ptr<llvm_ir::ForLoop> loop = llvm_ir::ForLoop::EmitForLoop( in ForWithStatus()
54 llvm_ir::SetToLastInsertPoint(loop->GetExitBasicBlock(), b_); in ForWithStatus()
63 llvm_ir::LlvmIfData if_data = in IfWithStatus()
64 llvm_ir::EmitIfThenElse(condition, name, b_, in IfWithStatus()
72 llvm_ir::SetToLastInsertPoint(if_data.after_block, b_); in IfWithStatus()
84 module->getFunction(llvm_ir::AsStringRef(kernel_name)); in EmitAndCallOutlinedKernel()
108 function = llvm_ir::CreateCpuFunction(function_type, in EmitAndCallOutlinedKernel()
139 b->CreateCall(function, llvm_ir::AsArrayRef(sanitized_args)); in EmitAndCallOutlinedKernel()
Dllvm_loop.h35 namespace llvm_ir {
86 UnrollMode unroll_mode = llvm_ir::UnrollMode::kDefaultUnroll,
199 UnrollMode unroll_mode = xla::llvm_ir::UnrollMode::kDefaultUnroll,
206 UnrollMode unroll_mode = xla::llvm_ir::UnrollMode::kDefaultUnroll,
214 UnrollMode unroll_mode = xla::llvm_ir::UnrollMode::kDefaultUnroll,
220 UnrollMode unroll_mode = xla::llvm_ir::UnrollMode::kDefaultUnroll,
254 const llvm_ir::IrArray& operand_array, int64 dimension_to_skip,
Dmath_ops.cc20 namespace llvm_ir { namespace
29 llvm_ir::EmitCallToIntrinsic(llvm::Intrinsic::fabs, {input}, {type}, b); in EmitFastTanh()
37 llvm::Value* input_clamped = llvm_ir::EmitFloatMin( in EmitFastTanh()
38 llvm_ir::EmitFloatMax(input, llvm::ConstantFP::get(type, -9.0), b, in EmitFastTanh()
Dllvm_loop.cc32 namespace llvm_ir { namespace
150 if (unroll_mode_ == xla::llvm_ir::UnrollMode::kNoUnroll) { in GetLoopMetadata()
161 if (unroll_mode_ == xla::llvm_ir::UnrollMode::kFullyUnroll) { in GetLoopMetadata()
169 return llvm_ir::IrName(prefix_, llvm_ir::IrName(name, suffix_)); in GetQualifiedName()
247 std::unique_ptr<llvm_ir::ForLoop> loop = AddLoop( in AddLoopsForShapeOnDimensions()
251 llvm_ir::IrName(suffix, absl::StrCat(dimension))); in AddLoopsForShapeOnDimensions()
258 const llvm_ir::IrArray& operand_array, int64 dimension_to_skip, in EmitOperandArrayLoopNest()
Dfused_ir_emitter.cc41 using llvm_ir::IrArray;
84 llvm_ir::GetGlobalMemoryAddressSpace(*module_); in HandleConstant()
88 llvm_ir::ConvertLiteralToIrConstant(literal, module_); in HandleConstant()
103 llvm_ir::ShapeToIrType(literal.shape(), module_)->getPointerTo()); in HandleConstant()
127 operand_elemental_ir_types.push_back(llvm_ir::PrimitiveTypeToIrType( in HandleTuple()
/external/mesa3d/src/amd/compiler/
Daco_interface.cpp109 std::string llvm_ir; in aco_compile_shader() local
121 llvm_ir = std::string(data, data + size); in aco_compile_shader()
174 size_t size = llvm_ir.size(); in aco_compile_shader()
230 legacy_binary->ir_size = llvm_ir.size(); in aco_compile_shader()
232 …llvm_ir.copy((char*) legacy_binary->data + legacy_binary->stats_size + legacy_binary->code_size, l… in aco_compile_shader()
235 …cy_binary->data + legacy_binary->stats_size + legacy_binary->code_size + llvm_ir.size(), disasm.si… in aco_compile_shader()

123