Home
last modified time | relevance | path

Searched refs:GetBasePointer (Results 1 – 9 of 9) sorted by relevance

/external/tensorflow/tensorflow/compiler/xla/service/llvm_ir/
Dtuple_ops.cc43 b->CreateLoad(pred.GetBasePointer(), "load_predicate_value"); in EmitTupleSelect()
65 b->CreateInBoundsGEP(select.GetBasePointer(), element_index); in EmitTupleSelect()
78 b->CreateInBoundsGEP(tuple.GetBasePointer(), in EmitTuple()
90 [](const llvm_ir::IrArray& buffer) { return buffer.GetBasePointer(); }); in EmitTuple()
Dfused_ir_emitter.h112 return GetIrArrayForFusedParameter(parameter_number).GetBasePointer(); in GetBasePointerForFusedParameter()
Dir_array.h216 llvm::Value* GetBasePointer() const { return base_ptr_; } in GetBasePointer() function
/external/tensorflow/tensorflow/compiler/xla/service/cpu/
Ddot_op_emitter.cc229 llvm::Value* lhs = lhs_array_.GetBasePointer(); in EmitTiledLlvmIrGemm()
230 llvm::Value* rhs = rhs_array_.GetBasePointer(); in EmitTiledLlvmIrGemm()
231 llvm::Value* target = target_array_.GetBasePointer(); in EmitTiledLlvmIrGemm()
313 llvm::Value* result_op = target_array_.GetBasePointer(); in EmitTiledLlvmIrGemv()
315 swap_operands ? rhs_array_.GetBasePointer() : lhs_array_.GetBasePointer(); in EmitTiledLlvmIrGemv()
317 swap_operands ? lhs_array_.GetBasePointer() : rhs_array_.GetBasePointer(); in EmitTiledLlvmIrGemv()
339 /*addend=*/addend_array_ ? addend_array_->GetBasePointer() : nullptr, in EmitTiledLlvmIrGemv()
349 /*addend=*/addend_array_ ? addend_array_->GetBasePointer() : nullptr, in EmitTiledLlvmIrGemv()
675 b_->CreateBitCast(target_array_.GetBasePointer(), float_ptr_type), in EmitCallToRuntime()
676 b_->CreateBitCast(lhs->GetBasePointer(), float_ptr_type), in EmitCallToRuntime()
[all …]
Dir_emitter.cc937 << llvm_ir::DumpToString(*lhs_array.GetBasePointer()); in HandleDot()
939 << llvm_ir::DumpToString(*rhs_array.GetBasePointer()); in HandleDot()
941 << llvm_ir::DumpToString(*target_array.GetBasePointer()); in HandleDot()
2559 GetIrArrayFor(branch_index).GetBasePointer(), "load_predicate_value"); in HandleConditional()
2594 GetIrArrayFor(branch_index).GetBasePointer(), "load_branch_index_value"); in HandleConditional()
2975 tuple_operand_ptrs.push_back(output_arrays[i].GetBasePointer()); in EmitTargetElementLoop()
/external/tensorflow/tensorflow/compiler/xla/service/gpu/
Dir_emitter.cc92 bindings_.BindHloToIrValue(*bitcast, GetBasePointer(*operand)); in HandleBitcast()
104 bindings_.BindHloToIrValue(*add_dependency, GetBasePointer(*operand)); in HandleAddDependency()
118 /*alignment=*/1, GetBasePointer(*operand), &b_)); in HandleGetTupleElement()
145 base_ptrs.push_back(GetBasePointer(*operand)); in HandleTuple()
436 GetBasePointer(*on_true), GetBasePointer(*on_false), in HandleTupleSelect()
680 Store(Load(GetBasePointer(*init_value)), accumulator_addr); in HandleReduce()
738 operand_addresses.push_back(GetBasePointer(*operand)); in HandleCall()
741 GetBasePointer(*call)); in HandleCall()
Dir_emitter.h128 llvm::Value* GetBasePointer(const HloInstruction& inst) const { in GetBasePointer() function
129 return bindings_.GetBasePointer(inst); in GetBasePointer()
Dhlo_to_ir_bindings.h73 llvm::Value* GetBasePointer(const HloInstruction& hlo,
Dhlo_to_ir_bindings.cc216 llvm::Value* base_ptr = GetBasePointer(hlo, shape_index); in GetIrArray()