/external/tensorflow/tensorflow/compiler/xla/service/gpu/ |
H A D | 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_ [all...] |
H A D | 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...] |
H A D | 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_); 925 .CastToShape(reduced_input_shape, &ir_builder_), 927 .CastToShape(reduced_output_shape, &ir_builder_), 928 tile_size, num_rows, &ir_builder_); 994 llvm::Value* partial_reduction_result_address = ir_builder_ [all...] |
H A D | 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_ 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...] |
H A D | 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
|
H A D | 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());
|
H A D | 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_)
|
/external/tensorflow/tensorflow/compiler/xla/service/ |
H A D | 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_ [all...] |
H A D | 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/ |
H A D | 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_ [all...] |
H A D | 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_ 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...] |
H A D | 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::__anon25556::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_ 416 llvm::IRBuilder<>* ir_builder_; member in class:xla::cpu::__anon25556::RowMajorMatrixVectorProductEmitter 764 llvm_ir::ForLoopNest loop_nest(llvm_ir::IrName(&dot_), ir_builder_); local [all...] |
H A D | parallel_loop_emitter.cc | 37 llvm_ir::ForLoopNest loop_nest(loop_name, ir_builder_); 66 ir_builder_);
|
H A D | 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)));
|
H A D | 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_);
|
H A D | dot_op_emitter.h | 146 llvm::IRBuilder<>* ir_builder_; member in class:xla::cpu::DotOpEmitter
|
H A D | ir_function.h | 100 llvm::IRBuilder<>* ir_builder_; member in class:xla::cpu::IrFunction
|
/external/tensorflow/tensorflow/compiler/xla/service/llvm_ir/ |
H A D | 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_ [all...] |
H A D | 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);
|
H A D | 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
|
H A D | 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_);
|
H A D | fused_ir_emitter.h | 60 ir_builder_(elemental_emitter->ir_builder()), 99 llvm::IRBuilder<>* ir_builder_; member in class:xla::FusedIrEmitter
|
H A D | 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,
|
H A D | loop_emitter.h | 87 llvm::IRBuilder<>* ir_builder_; member in class:xla::llvm_ir::LoopEmitter
|
H A D | llvm_loop.h | 181 ir_builder_(ir_builder) {} 260 llvm::IRBuilder<>* ir_builder_; member in class:xla::llvm_ir::ForLoopNest
|