| 977 | } |
| 978 | |
| 979 | Status IrEmitter::HandleDot(HloInstruction* dot) { |
| 980 | auto lhs = dot->operand(0); |
| 981 | auto rhs = dot->operand(1); |
| 982 | TF_RETURN_IF_ERROR(ElementTypesSameAndSupported( |
| 983 | /*instruction=*/*dot, /*operands=*/{lhs, rhs}, |
| 984 | /*supported_types=*/{S32, F16, F32, F64, C64, C128})); |
| 985 | const DotDimensionNumbers& dnums = dot->dot_dimension_numbers(); |
| 986 | |
| 987 | if (dnums.lhs_contracting_dimensions_size() != 1) { |
| 988 | // This is disallowed by ShapeInference today. |
| 989 | return Unimplemented( |
| 990 | "Dot with multiple contracting dimensions not implemented."); |
| 991 | } |
| 992 | |
| 993 | llvm_ir::IrArray lhs_array(GetIrArrayFor(lhs)); |
| 994 | llvm_ir::IrArray rhs_array(GetIrArrayFor(rhs)); |
| 995 | |
| 996 | TF_RETURN_IF_ERROR(EmitTargetAddressForOp(dot)); |
| 997 | llvm_ir::IrArray target_array = GetIrArrayFor(dot); |
| 998 | |
| 999 | VLOG(2) << "HandleDot: "; |
| 1000 | VLOG(2) << " lhs operand: " |
| 1001 | << llvm_ir::DumpToString(*lhs_array.GetBasePointer()); |
| 1002 | VLOG(2) << " rhs operand: " |
| 1003 | << llvm_ir::DumpToString(*rhs_array.GetBasePointer()); |
| 1004 | VLOG(2) << " target: " |
| 1005 | << llvm_ir::DumpToString(*target_array.GetBasePointer()); |
| 1006 | |
| 1007 | // Dot operation is complicated so we delegate to a helper class. |
| 1008 | return EmitDotOperation(*dot, target_array, lhs_array, rhs_array, |
| 1009 | /*addend_array=*/nullptr, |
| 1010 | GetExecutableRunOptionsArgument(), &b_, |
| 1011 | hlo_module_config_, target_machine_features_); |
| 1012 | } |
| 1013 | |
| 1014 | StatusOr<llvm::Value*> IrEmitter::EmitElementalConvolution( |
| 1015 | const HloConvolutionInstruction* convolution, |
nothing calls this directly
no test coverage detected