| /external/tensorflow/tensorflow/compiler/xla/service/gpu/ |
| parallel_loop_emitter.cc | 67 llvm::Intrinsic::nvvm_read_ptx_sreg_ctaid_x, {}, {}, ir_builder_);
71 ir_builder_->CreateZExt(block_id, ir_builder_->getInt64Ty(), "block_id");
78 llvm::Intrinsic::nvvm_read_ptx_sreg_tid_x, {}, {}, ir_builder_);
81 thread_id = ir_builder_->CreateZExt(thread_id, ir_builder_->getInt64Ty(),
84 llvm::Value* linear_index = ir_builder_->CreateAdd(
85 ir_builder_->CreateMul(
87 ir_builder_->getInt64(launch_dimensions_.threads_per_block()), "",
101 {ir_builder_->CreateICmpULT
[all...] |
| ir_emitter.cc | 59 ir_builder_(module_->getContext()),
61 &ir_emitter_context->buffer_assignment(), &ir_builder_, module_,
64 ir_builder_.setFastMathFlags(llvm_ir::GetFastMathFlags(
74 .EmitReadArrayElement(index, &ir_builder_);
78 *hlo, GpuElementalIrEmitter(hlo_module_config_, module_, &ir_builder_,
121 /*alignment=*/1, GetBasePointer(*operand), &ir_builder_, module_));
151 llvm_ir::EmitTuple(GetIrArray(*tuple, *tuple), base_ptrs, &ir_builder_,
173 ir_builder_.CreateCall(emitted_function, arguments);
193 llvm::Value* source = ir_builder_.CreateLoad(source_address, "source");
200 {output_address->getType()}, &ir_builder_);
504 llvm_ir::ForLoopNest loop_nest(IrName(dot), &ir_builder_); local
649 llvm_ir::ForLoopNest loops(IrName(reduce, "inner"), &ir_builder_); local
[all...] |
| ir_emitter_unnested.cc | 239 std::vector<llvm::Type*>(args.size(), ir_builder_.getInt8PtrTy()),
275 llvm::ConstantAsMetadata::get(ir_builder_.getInt32(1))}));
283 ir_builder_.SetInsertPoint(llvm::ReturnInst::Create(context, entry_bb));
515 &ir_builder_, GetNestedComputer());
577 &ir_builder_, GetNestedComputer());
595 launch_dimensions, &ir_builder_);
[all...] |
| elemental_ir_emitter.cc | 89 converted_operands[i] = ir_builder_->CreateFPCast(
90 converted_operands[i], ir_builder_->getFloatTy());
109 result = ir_builder_->CreateFPCast(result, ir_builder_->getHalfTy());
212 return ir_builder_->CreateFDiv(llvm::ConstantFP::get(llvm_ty, 1), sqrt);
289 ir_builder_->GetInsertBlock()->getModule()->getOrInsertFunction(
296 return ir_builder_->CreateCall(callee, llvm_ir::AsArrayRef(operands));
300 llvm::Value* block_id = ir_builder_->CreateIntCast(
302 {}, {}, ir_builder_),
303 ir_builder_->getIntNTy(128), /*isSigned=*/true, "block.id")
367 llvm_ir::ForLoopNest loops(IrName(hlo), ir_builder_); local
426 llvm_ir::ForLoopNest loops(IrName(hlo), ir_builder_); member in namespace:xla::gpu
[all...] |
| hlo_to_ir_bindings.h | 43 ir_builder_(ir_builder),
46 &ir_builder_->getContext()) {}
107 llvm::IRBuilder<>* ir_builder_; member in class:xla::gpu::HloToIrBindings
|
| ir_emitter_nested.cc | 72 argument_types.push_back(ir_builder_.getInt8PtrTy());
75 llvm::FunctionType::get(ir_builder_.getVoidTy(), argument_types, false);
98 ir_builder_.SetInsertPoint(
120 &ir_builder_)
|
| hlo_to_ir_bindings.cc | 42 llvm::Function* function = ir_builder_->GetInsertBlock()->getParent();
82 llvm::Value* base_ptr = ir_builder_->CreateInBoundsGEP(
83 temp_buffer_base_, ir_builder_->getInt64(offset));
112 ir_builder_->CreateAlloca(pointee_type), index);
118 ir_builder_->CreateInBoundsGEP(temp_buffer_base_,
119 ir_builder_->getInt64(offset)),
132 GetTypedIrValue(*gte->operand(0), {}, base_ptr), ir_builder_, module_);
136 EmitGetTupleElement(gte->operand(0), base_ptr), ir_builder_, module_);
152 ir_builder_->CreateBitCast(ir_value, pointee_type->getPointerTo());
|
| /external/tensorflow/tensorflow/compiler/xla/service/ |
| elemental_ir_emitter.cc | 227 return ir_builder_->CreateIntCast(
235 ir_builder_),
236 ir_builder_);
239 module_, ir_builder_);
247 ir_builder_->CreateSIToFP(operand_value, to_ir_component_type),
254 ir_builder_->CreateUIToFP(operand_value, to_ir_component_type),
271 return ir_builder_->CreateBitCast(
289 auto cmp = ir_builder_->CreateICmpSGE(operand_value, zero);
290 return ir_builder_->CreateSelect(cmp, operand_value,
291 ir_builder_->CreateNeg(operand_value))
[all...] |
| elemental_ir_emitter.h | 38 : ir_builder_(ir_builder),
57 llvm::IRBuilder<>* ir_builder() const { return ir_builder_; }
142 return ir_builder_->getIntN(128, 0);
145 llvm::IRBuilder<>* const ir_builder_; member in class:xla::ElementalIrEmitter::llvm
|
| /external/tensorflow/tensorflow/compiler/xla/service/cpu/ |
| elemental_ir_emitter.cc | 41 operand_value = ir_builder_->CreateFPCast(operand_value,
42 ir_builder_->getFloatTy());
62 llvm::Value* result = ir_builder_->CreateCall(function, operand_value);
64 result = ir_builder_->CreateFPCast(result, ir_builder_->getHalfTy());
80 lhs = ir_builder_->CreateFPCast(lhs, ir_builder_->getFloatTy());
81 rhs = ir_builder_->CreateFPCast(rhs, ir_builder_->getFloatTy());
101 llvm::Value* result = ir_builder_->CreateCall(function, {lhs, rhs})
[all...] |
| ir_emitter.cc | 91 ir_builder_(llvm_module->getContext()),
101 ir_builder_.setFastMathFlags(llvm_ir::GetFastMathFlags(
151 module_, &ir_builder_, num_dynamic_loop_bounds_));
159 ir_builder_.CreateBitCast(GetEmittedValueFor(bitcast->operand(0)),
309 GetEmittedValueFor(operand), &ir_builder_, module_);
323 GetEmittedValueFor(on_false), &ir_builder_, module_); local
368 llvm_ir::EmitTuple(infeed_array, tuple_element_addresses, &ir_builder_,
392 shape, &shape_length, &ir_builder_));
397 llvm::Type* int32_type = ir_builder_.getInt32Ty();
417 ir_builder_.getVoidTy()
491 llvm_ir::EmitTuple(GetIrArrayFor(tuple), base_ptrs, &ir_builder_, module_); local
559 &ir_builder_); local
666 llvm_ir::ForLoopNest source_loops(IrName(select_and_scatter), &ir_builder_); local
688 &ir_builder_); local
1020 llvm_ir::ForLoopNest loops(IrName(convolution, "inner"), &ir_builder_); local
[all...] |
| dot_op_emitter.cc | 159 ir_builder_(ir_builder),
160 ksl_(ir_builder_),
161 vsl_(scalar_type_, /*vector_size=*/tile_rows_, ir_builder_, "") {
172 return TileLoader(&vsl_, ir_builder_, /*matrix=*/lhs_,
206 llvm::IRBuilder<>* ir_builder_; member in class:xla::cpu::__anon38877::ColumnMajorMatrixVectorProductEmitter
235 EmitOuterLoopBody(ir_builder_->getInt64(column_limit), column_remainder,
267 llvm::Value* columns_llvm = ir_builder_->getInt64(columns);
278 /*end=*/ir_builder_->CreateAdd(columns_llvm, current_tile_col),
283 ir_builder_->CreateMul(col, ir_builder_->getInt64(m_))
416 llvm::IRBuilder<>* ir_builder_; member in class:xla::cpu::__anon38877::RowMajorMatrixVectorProductEmitter
764 llvm_ir::ForLoopNest loop_nest(llvm_ir::IrName(&dot_), ir_builder_); local
[all...] |
| parallel_loop_emitter.cc | 37 llvm_ir::ForLoopNest loop_nest(loop_name, ir_builder_);
66 ir_builder_);
|
| ir_function.cc | 54 : ir_builder_(ir_builder),
64 ir_builder_->CreateRetVoid();
177 ir_builder_->SetInsertPoint(llvm::BasicBlock::Create(
187 return ir_builder_->CreateLoad(
188 ir_builder_->CreateGEP(CHECK_NOTNULL(dynamic_loop_bounds_arg_),
189 ir_builder_->getInt64(offset), AsStringRef(name)));
|
| vector_support_library.cc | 30 ir_builder_(ir_builder),
33 primitive_type, ir_builder_->GetInsertBlock()->getModule());
85 return llvm_ir::EmitFloatMax(lhs, rhs, ir_builder_);
114 llvm_ir::EmitFloatMax(a, GetConstantFloat(type, low), ir_builder_),
115 GetConstantFloat(type, high), ir_builder_);
412 : ir_builder_(ir_builder) {
413 alloca_ = llvm_ir::EmitAllocaAtFunctionEntry(type, "", ir_builder_);
417 return ir_builder_->CreateLoad(alloca_);
421 ir_builder_->CreateStore(new_value, alloca_);
|
| dot_op_emitter.h | 146 llvm::IRBuilder<>* ir_builder_; member in class:xla::cpu::DotOpEmitter
|
| ir_function.h | 100 llvm::IRBuilder<>* ir_builder_; member in class:xla::cpu::IrFunction
|
| /external/tensorflow/tensorflow/compiler/xla/service/llvm_ir/ |
| kernel_support_library.cc | 26 If(ir_builder_->CreateICmpSLT(start, end), [&]() {
28 For(name, ir_builder_->CreateAdd(start, step), end, step,
40 for_body_generator(indvar, ir_builder_->getInt1(is_first_iteration));
44 name, start, end, step, ir_builder_,
47 ir_builder_->SetInsertPoint(&loop->GetBodyBasicBlock()->back());
49 /*is_first_iteration=*/ir_builder_->CreateICmpEQ(
51 llvm_ir::SetToLastInsertPoint(loop->GetExitBasicBlock(), ir_builder_);
59 llvm_ir::EmitIfThenElse(condition, "", ir_builder_);
60 ir_builder_->SetInsertPoint(&if_data.true_block->back());
62 ir_builder_->SetInsertPoint(&if_data.false_block->back())
[all...] |
| fused_ir_emitter.cc | 55 generated_value_bb == ir_builder_->GetInsertBlock()) {
63 << llvm_ir::AsString(ir_builder_->GetInsertBlock()->getName())
80 *ir_builder_->GetInsertBlock()->getModule(), initializer->getType(),
85 .EmitReadArrayElement(index, ir_builder_);
105 /*alignment=*/1, it->second, ir_builder_, module_);
113 .EmitReadArrayElement(index, ir_builder_);
122 .EmitReadArrayElement(index, ir_builder_);
142 ir_builder_->getContext(), operand_elemental_ir_types));
145 ret = ir_builder_->CreateInsertValue(ret, val_i, i);
|
| kernel_support_library.h | 38 : ir_builder_(ir_builder),
59 For(name, /*start=*/ir_builder_->getInt64(start),
60 /*end=*/ir_builder_->getInt64(end),
61 /*step=*/ir_builder_->getInt64(step), for_body_generator);
90 /*step=*/ir_builder_->getInt64(step), peel_first_iteration,
106 For(name, /*start=*/ir_builder_->getInt64(start),
107 /*end=*/ir_builder_->getInt64(end),
108 /*step=*/ir_builder_->getInt64(step), for_body_generator);
176 llvm::IRBuilder<>* ir_builder_; member in class:xla::KernelSupportLibrary
|
| loop_emitter.cc | 37 : body_emitter_(body_emitter), shape_(shape), ir_builder_(ir_builder) {}
52 ir_builder_(ir_builder) {}
83 ir_builder_(ir_builder) {
103 ForLoopNest loop_nest(loop_name, ir_builder_);
117 ir_builder_->SetInsertPoint(innermost_body_bb,
131 // Set the insertion point of ir_builder_ to the loop exit, so that
134 ir_builder_->SetInsertPoint(exit_bb_);
|
| fused_ir_emitter.h | 60 ir_builder_(elemental_emitter->ir_builder()),
99 llvm::IRBuilder<>* ir_builder_; member in class:xla::FusedIrEmitter
|
| loop_emitter.h | 87 llvm::IRBuilder<>* ir_builder_; member in class:xla::llvm_ir::LoopEmitter
|
| llvm_loop.cc | 183 return AddLoop(suffix, start_index, end_index, ir_builder_->getInt64(1),
195 ir_builder_->SetInsertPoint(&*inner_loop_body_bb_->getFirstInsertionPt());
200 loop->Emit(ir_builder_);
221 return AddLoop(suffix, ir_builder_->getInt64(start_index),
222 ir_builder_->getInt64(end_index), prevent_unrolling,
232 return AddLoop(suffix, ir_builder_->getInt64(start_index),
233 ir_builder_->getInt64(end_index),
234 ir_builder_->getInt64(stride), prevent_unrolling,
|
| llvm_loop.h | 181 ir_builder_(ir_builder) {}
260 llvm::IRBuilder<>* ir_builder_; member in class:xla::llvm_ir::ForLoopNest
|
