| 538 | } |
| 539 | |
| 540 | Status HloCostAnalysis::HandleConvolution(const HloInstruction* convolution) { |
| 541 | auto lhs = convolution->operand(0); |
| 542 | auto rhs = convolution->operand(1); |
| 543 | Window window = convolution->window(); |
| 544 | const auto& result_shape = convolution->shape(); |
| 545 | const Shape& lhs_shape = lhs->shape(); |
| 546 | const Shape& rhs_shape = rhs->shape(); |
| 547 | |
| 548 | const auto& dnums = convolution->convolution_dimension_numbers(); |
| 549 | |
| 550 | const int64 input_batch_dim = dnums.input_batch_dimension(); |
| 551 | const int64 input_feature_dim = dnums.input_feature_dimension(); |
| 552 | const int64 output_feature_dim = dnums.output_feature_dimension(); |
| 553 | const int64 input_feature = |
| 554 | ShapeUtil::GetDimension(lhs_shape, input_feature_dim); |
| 555 | const int64 output_feature = |
| 556 | ShapeUtil::GetDimension(result_shape, output_feature_dim); |
| 557 | const int64 batch = ShapeUtil::GetDimension(lhs_shape, input_batch_dim); |
| 558 | |
| 559 | DimensionVector kernel_limits; |
| 560 | DimensionVector output_limits; |
| 561 | DimensionVector input_limits; |
| 562 | if (window.dimensions().empty()) { |
| 563 | window = window_util::MakeWindow({1}); |
| 564 | kernel_limits.push_back(1); |
| 565 | output_limits.push_back(1); |
| 566 | input_limits.push_back(1); |
| 567 | } else { |
| 568 | for (int64 spatial_dimension = 0; |
| 569 | spatial_dimension < window.dimensions_size(); ++spatial_dimension) { |
| 570 | // Spatial dimension number for kernel (rhs). |
| 571 | const int64 kernel_spatial_dim = |
| 572 | dnums.kernel_spatial_dimensions(spatial_dimension); |
| 573 | const int64 kernel_limit = rhs_shape.dimensions(kernel_spatial_dim); |
| 574 | kernel_limits.push_back(kernel_limit); |
| 575 | |
| 576 | // Spatial dimension number for output. |
| 577 | const int64 output_spatial_dim = |
| 578 | dnums.output_spatial_dimensions(spatial_dimension); |
| 579 | const int64 output_limit = result_shape.dimensions(output_spatial_dim); |
| 580 | output_limits.push_back(output_limit); |
| 581 | |
| 582 | // Spatial dimension number for input (lhs). |
| 583 | const int64 input_spatial_dim = |
| 584 | dnums.input_spatial_dimensions(spatial_dimension); |
| 585 | const int64 input_limit = lhs_shape.dimensions(input_spatial_dim); |
| 586 | input_limits.push_back(input_limit); |
| 587 | } |
| 588 | } |
| 589 | |
| 590 | DimensionVector valid_position_counts; |
| 591 | |
| 592 | // Loop over each spatial dimension. |
| 593 | for (int64 spatial_dimension = 0; |
| 594 | spatial_dimension < window.dimensions_size(); ++spatial_dimension) { |
| 595 | const auto& window_dim = window.dimensions(spatial_dimension); |
| 596 | // These two conditions will create an N^2 iteration pattern with only N |
| 597 | // valid elements. This is a performance optimization and produces the same |
nothing calls this directly
no test coverage detected