Home
last modified time | relevance | path

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

1234

/external/tensorflow/tensorflow/compiler/xla/service/gpu/
Dir_emitter_nested.cc74 llvm_ir::ShapeToIrType(param_shape, module_)->getPointerTo()); in CodegenNestedComputation()
76 llvm_ir::ByteSizeOf(param_shape, module_->getDataLayout()); in CodegenNestedComputation()
84 llvm_ir::ShapeToIrType(root_shape, module_)->getPointerTo()); in CodegenNestedComputation()
85 int64_t root_size = llvm_ir::ByteSizeOf( in CodegenNestedComputation()
96 llvm_ir::SanitizeFunctionName( in CodegenNestedComputation()
144 Load(llvm_ir::ShapeToIrType(return_shape, module_), root_value, in CodegenNestedComputation()
150 llvm::Type* tuple_type = llvm_ir::ShapeToIrType(return_shape, module_); in CodegenNestedComputation()
156 llvm::Value* destination = llvm_ir::EmitGetTupleElement( in CodegenNestedComputation()
160 llvm::Value* source = llvm_ir::EmitGetTupleElement( in CodegenNestedComputation()
164 llvm_ir::ShapeToIrType(root_instruction->shape(), module_), &b_); in CodegenNestedComputation()
[all …]
Dhlo_to_ir_bindings.cc84 llvm_ir::ConstantHloToGlobalName(*non_io_hlo)); in EmitBasePointersForHlos()
86 << llvm_ir::ConstantHloToGlobalName(*non_io_hlo); in EmitBasePointersForHlos()
90 llvm_ir::ShapeToIrType(non_io_hlo->shape(), module_); in EmitBasePointersForHlos()
92 llvm_ir::EmitAllocaAtFunctionEntry( in EmitBasePointersForHlos()
104 return llvm_ir::EmitGetTupleElement( in EmitGetTupleElement()
107 llvm_ir::ShapeToIrType(gte->operand(0)->shape(), module_), b_); in EmitGetTupleElement()
109 return llvm_ir::EmitGetTupleElement( in EmitGetTupleElement()
112 llvm_ir::ShapeToIrType(gte->operand(0)->shape(), module_), b_); in EmitGetTupleElement()
126 llvm_ir::ShapeToIrType(shape, b->GetInsertBlock()->getModule()); in CastToTypedValue()
147 ir_value->setName(llvm_ir::IrName(&hlo, "raw")); in GetTypedIrValue()
[all …]
Dparallel_loop_emitter.cc35 llvm_ir::BodyEmitter body_emitter, const Shape& shape, in ParallelLoopEmitter()
45 const llvm_ir::ElementGenerator& target_element_generator, in ParallelLoopEmitter()
46 absl::Span<const llvm_ir::IrArray> target_arrays, in ParallelLoopEmitter()
53 llvm_ir::MakeBodyEmitter(target_element_generator, target_arrays, b, in ParallelLoopEmitter()
63 llvm_ir::AddRangeMetadata(0, launch_dimensions_.block_counts().x, in EmitLinearBaseAndThreadIdx()
71 llvm_ir::AddRangeMetadata(0, launch_dimensions_.thread_counts_per_block().x, in EmitLinearBaseAndThreadIdx()
98 llvm_ir::AddRangeMetadata(0, launch_dimensions_.thread_counts_per_block().y, in EmitLinearBaseAndThreadIdx()
121 llvm_ir::EmitCallToIntrinsic( in EmitLinearBaseAndThreadIdx()
146 std::vector<llvm_ir::IrArray::Index>
150 std::vector<llvm_ir::IrArray::Index> array_indices; in EmitLogicalIndexAndSetExitBasicBlock()
[all …]
Dir_emitter_unnested.h137 const ThreadIdInfo& thread_id_info, const llvm_ir::IrArray::Index& index,
147 const ThreadIdInfo& thread_id_info, const llvm_ir::IrArray::Index& index,
152 ConstHloInstructionMap<absl::Span<llvm_ir::IrArray const>>;
154 using ExtraOutputGensMap = ConstHloInstructionMap<llvm_ir::ElementGenerator>;
231 const llvm_ir::ElementGenerator& body_emitter) override;
358 return llvm_ir::ByteSizeOf( in ByteSizeOf()
369 const llvm_ir::IrArray::Index& index,
460 absl::Span<const llvm_ir::IrArray> ir_arrays,
461 const llvm_ir::IrArray::Index& index);
472 const llvm_ir::IrArray& output,
[all …]
Dparallel_loop_emitter.h36 ParallelLoopEmitter(llvm_ir::BodyEmitter body_emitter, const Shape& shape,
46 ParallelLoopEmitter(const llvm_ir::ElementGenerator& target_element_generator,
47 absl::Span<const llvm_ir::IrArray> target_arrays,
55 std::vector<llvm_ir::IrArray::Index> EmitIndexAndSetExitBasicBlock(
63 std::vector<llvm_ir::IrArray::Index> EmitLogicalIndexAndSetExitBasicBlock(
86 llvm_ir::BodyEmitter body_emitter_;
Dir_emitter.cc67 using llvm_ir::IrName;
68 using llvm_ir::SetToFirstInsertPoint;
83 operand_to_generator[operand] = [=](const llvm_ir::IrArray::Index& index) { in DefaultAction()
115 llvm_ir::EmitGetTupleElement( in HandleGetTupleElement()
120 llvm_ir::ShapeToIrType(operand->shape(), module_), &b_)); in HandleGetTupleElement()
149 llvm_ir::EmitTuple(GetIrArray(*tuple, *tuple), base_ptrs, &b_); in HandleTuple()
195 Load(llvm_ir::PrimitiveTypeToIrType(element_type, module_), in MaybeEmitDirectAtomicOperation()
341 int element_size = llvm_ir::GetSizeInBits(element_type); in EmitAtomicOperationUsingCAS()
351 llvm::AllocaInst* cas_old_output_address = llvm_ir::EmitAllocaAtFunctionEntry( in EmitAtomicOperationUsingCAS()
353 llvm::AllocaInst* cas_new_output_address = llvm_ir::EmitAllocaAtFunctionEntry( in EmitAtomicOperationUsingCAS()
[all …]
Dir_emitter_unnested.cc144 using llvm_ir::IrArray;
145 using llvm_ir::IrName;
467 return llvm_ir::SanitizeConstantName(mlir::GetNameFromLoc(loc)); in GetIrNameFromLoc()
498 llvm_ir::SanitizeFunctionName(std::string(name))); in BuildKernelPrototype()
738 std::vector<llvm_ir::IrArray> ir_arrays; in EmitPadToStatic()
743 const llvm_ir::IrArray source_array = ir_arrays[0]; in EmitPadToStatic()
744 const llvm_ir::IrArray output_array = ir_arrays[1]; in EmitPadToStatic()
746 absl::Span<const llvm_ir::IrArray>(ir_arrays).subspan(2); in EmitPadToStatic()
824 llvm_ir::BodyEmitter body_generator = in EmitPadToStatic()
825 [&](const llvm_ir::IrArray::Index& array_index) -> Status { in EmitPadToStatic()
[all …]
Dir_emitter.h114 llvm_ir::IrArray GetIrArray(const HloInstruction& inst,
127 std::vector<llvm_ir::IrArray> ConstructIrArrayForOutputs(
136 const llvm_ir::ElementGenerator& body_emitter) = 0;
206 const llvm_ir::IrArray::Index& keys_index,
207 const llvm_ir::IrArray::Index& compare_keys_index,
208 const llvm_ir::IrArray& keys_array);
Delemental_ir_emitter.cc55 using llvm_ir::IrArray;
56 using llvm_ir::IrName;
57 using llvm_ir::SetToFirstInsertPoint;
161 llvm_ir::IrArray::Index GpuElementalIrEmitter::GetSourceIndexOfBitcast( in GetSourceIndexOfBitcast()
162 const llvm_ir::IrArray::Index& index, const HloInstruction* hlo) { in GetSourceIndexOfBitcast()
187 return llvm_ir::EmitCallToIntrinsic( in EmitFloatBinaryOp()
305 llvm_ir::EmitCallToIntrinsic(llvm::Intrinsic::fabs, {input}, {type}, b()); in EmitTanh()
307 llvm::Value* fast_tanh = llvm_ir::EmitFastTanh(b(), input); in EmitTanh()
309 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,
234 const llvm_ir::IrArray& target_array_;
235 const llvm_ir::IrArray& lhs_array_;
236 const llvm_ir::IrArray& rhs_array_;
237 const llvm_ir::IrArray* addend_array_;
248 const llvm_ir::IrArray& target_array, const llvm_ir::IrArray& lhs_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.cc89 using llvm_ir::IrName;
90 using llvm_ir::SetToFirstInsertPoint;
121 b_.setFastMathFlags(llvm_ir::GetCpuFastMathFlags(hlo_module_config_)); in IrEmitter()
131 llvm_ir::IrArray root_value = GetIrArrayFor(computation->root_instruction()); in EmitThreadLocalFunctionEpilogue()
143 llvm::Type* tuple_type = llvm_ir::ShapeToIrType(return_shape, module_); in EmitThreadLocalFunctionEpilogue()
149 llvm::Value* destination = llvm_ir::EmitGetTupleElement( in EmitThreadLocalFunctionEpilogue()
155 llvm::Value* source = llvm_ir::EmitGetTupleElement( in EmitThreadLocalFunctionEpilogue()
263 llvm_ir::ConvertLiteralToIrConstant(literal, module_); in EmitGlobalForLiteral()
284 const Literal& literal = llvm_ir::LiteralForConstantAllocation(allocation); in EmitConstantGlobals()
338 return llvm_ir::ByteSizeOf(shape, module_->getDataLayout()); in ByteSizeOf()
[all …]
Ddot_op_emitter.h61 const llvm_ir::IrArray& target_array,
62 const llvm_ir::IrArray& lhs_array,
63 const llvm_ir::IrArray& rhs_array,
64 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(
Dir_emitter.h68 std::function<std::vector<llvm_ir::IrArray>()>;
234 llvm_ir::IrArray GetIrArrayFor(const HloInstruction* hlo);
237 std::vector<llvm_ir::IrArray> GetIrArraysForOperandsOf(
246 llvm_ir::IrArray* array) { in AddAliasingInformationToIrArray()
333 const llvm_ir::ElementGenerator& element_generator);
336 const llvm_ir::ElementGenerator& element_generator);
403 const llvm_ir::IrArray& containing_array);
419 const llvm_ir::IrArray::Index& output_index,
435 const llvm_ir::IrArray& target_array,
436 const llvm_ir::IrArray& source_array);
[all …]
DBUILD317 "//tensorflow/compiler/xla/service/llvm_ir:llvm_command_line_options",
321 "//tensorflow/compiler/xla/service/llvm_ir:llvm_util",
526 "//tensorflow/compiler/xla/service/llvm_ir:alias_analysis",
527 "//tensorflow/compiler/xla/service/llvm_ir:buffer_assignment_util",
528 "//tensorflow/compiler/xla/service/llvm_ir:dynamic_update_slice_util",
529 "//tensorflow/compiler/xla/service/llvm_ir:fused_ir_emitter",
530 "//tensorflow/compiler/xla/service/llvm_ir:ir_array",
531 "//tensorflow/compiler/xla/service/llvm_ir:ir_builder_mixin",
532 "//tensorflow/compiler/xla/service/llvm_ir:llvm_loop",
533 "//tensorflow/compiler/xla/service/llvm_ir:llvm_type_conversion_util",
[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;
231 llvm_ir::PrimitiveTypeToIrType(to_type, module)); in EmitIntegralToFloating()
236 llvm_ir::PrimitiveTypeToIrType(to_type, module)); in EmitIntegralToFloating()
269 llvm_ir::PrimitiveTypeToIrType(PRED, module_)); in EmitIntegerUnaryOp()
273 llvm_ir::PrimitiveTypeToIrType(to_type, module_), in EmitIntegerUnaryOp()
286 auto to_ir_component_type = llvm_ir::PrimitiveTypeToIrType( in EmitIntegerUnaryOp()
312 llvm_ir::PrimitiveTypeToIrType(to_type, module_)); in EmitIntegerUnaryOp()
326 llvm_ir::PrimitiveTypeToIrType(op->shape().element_type(), module_); in EmitIntegerUnaryOp()
[all …]
Delemental_ir_emitter.h39 absl::flat_hash_map<const HloInstruction*, llvm_ir::ElementGenerator>;
48 llvm_ir::ElementGenerator MakeElementGenerator(
60 virtual llvm_ir::IrArray::Index GetSourceIndexOfBitcast( in GetSourceIndexOfBitcast()
61 const llvm_ir::IrArray::Index& index, const HloInstruction* hlo) { in GetSourceIndexOfBitcast()
235 const llvm_ir::IrArray::Index& index);
240 const llvm_ir::IrArray::Index& index);
245 const llvm_ir::IrArray::Index& target_index);
250 const llvm_ir::IrArray::Index& index);
255 const llvm_ir::IrArray::Index& index);
260 const llvm_ir::IrArray::Index& index);
[all …]
/external/tensorflow/tensorflow/compiler/xla/service/cpu/tests/
Dcpu_noalias_test.cc72 llvm_ir::AliasAnalysis aa(*hlo_module, *status_or_buffer_assn.ValueOrDie(), in TEST_F()
92 llvm_ir::IrArray param_x_array(param_x_val, param_x_val->getValueType(), 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, concat1_val->getValueType(), in TEST_F()
107 llvm_ir::IrArray::Index zero_2d({zero, zero}, shape, zero->getType()); in TEST_F()
116 llvm_ir::IrArray concat2_array(concat2_val, concat2_val->getValueType(), in TEST_F()
119 llvm_ir::IrArray::Index zero_2d({zero, zero}, shape, zero->getType()); in TEST_F()
128 llvm_ir::IrArray add_array(concat2_val, concat2_val->getValueType(), shape); in TEST_F()
130 llvm_ir::IrArray::Index zero_2d({zero, zero}, shape, zero->getType()); in TEST_F()
149 RunFileCheck(llvm_ir::DumpModuleToString(ir_module), filecheck_pattern)); in TEST_F()
/external/tensorflow/tensorflow/compiler/xla/service/llvm_ir/
Dkernel_support_library.cc46 std::unique_ptr<llvm_ir::ForLoop> loop = llvm_ir::ForLoop::EmitForLoop( in ForWithStatus()
55 llvm_ir::SetToLastInsertPoint(loop->GetExitBasicBlock(), b_); in ForWithStatus()
64 llvm_ir::LlvmIfData if_data = in IfWithStatus()
65 llvm_ir::EmitIfThenElse(condition, name, b_, in IfWithStatus()
73 llvm_ir::SetToLastInsertPoint(if_data.after_block, b_); in IfWithStatus()
85 module->getFunction(llvm_ir::AsStringRef(kernel_name)); in EmitAndCallOutlinedKernel()
109 function = llvm_ir::CreateCpuFunction(function_type, in EmitAndCallOutlinedKernel()
140 b->CreateCall(function, llvm_ir::AsArrayRef(sanitized_args)); in EmitAndCallOutlinedKernel()
Dllvm_loop.h33 namespace llvm_ir {
84 UnrollMode unroll_mode = llvm_ir::UnrollMode::kDefaultUnroll,
197 UnrollMode unroll_mode = xla::llvm_ir::UnrollMode::kDefaultUnroll,
204 UnrollMode unroll_mode = xla::llvm_ir::UnrollMode::kDefaultUnroll,
212 UnrollMode unroll_mode = xla::llvm_ir::UnrollMode::kDefaultUnroll,
218 UnrollMode unroll_mode = xla::llvm_ir::UnrollMode::kDefaultUnroll,
252 const llvm_ir::IrArray& operand_array, int64_t 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
151 if (unroll_mode_ == xla::llvm_ir::UnrollMode::kNoUnroll) { in GetLoopMetadata()
162 if (unroll_mode_ == xla::llvm_ir::UnrollMode::kFullyUnroll) { in GetLoopMetadata()
170 return llvm_ir::IrName(prefix_, llvm_ir::IrName(name, suffix_)); in GetQualifiedName()
248 std::unique_ptr<llvm_ir::ForLoop> loop = AddLoop( in AddLoopsForShapeOnDimensions()
252 llvm_ir::IrName(suffix, absl::StrCat(dimension))); in AddLoopsForShapeOnDimensions()
259 const llvm_ir::IrArray& operand_array, int64_t dimension_to_skip, in EmitOperandArrayLoopNest()
Dtuple_ops.cc32 namespace llvm_ir { namespace
58 [](const llvm_ir::IrArray& buffer) { return buffer.GetBasePointer(); }); in EmitTuple()
59 llvm_ir::EmitTuple(tuple, buffer_ptrs, b); in EmitTuple()
78 llvm_ir::PrimitiveTypeToIrType(element_shape.element_type(), module); in EmitTupleAllocasAtFunctionEntry()
/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()

1234