Home
last modified time | relevance | path

Searched refs:dnums (Results 1 – 25 of 32) sorted by relevance

12

/external/tensorflow/tensorflow/compiler/xla/service/cpu/
Dconv_canonicalization_test.cc69 ConvolutionDimensionNumbers dnums; in TEST_F() local
70 dnums.set_input_batch_dimension(1); in TEST_F()
71 dnums.set_output_batch_dimension(1); in TEST_F()
72 dnums.add_input_spatial_dimensions(2); in TEST_F()
73 dnums.add_output_spatial_dimensions(2); in TEST_F()
74 dnums.add_input_spatial_dimensions(3); in TEST_F()
75 dnums.add_output_spatial_dimensions(3); in TEST_F()
76 dnums.set_input_feature_dimension(0); in TEST_F()
77 dnums.set_output_feature_dimension(0); in TEST_F()
78 dnums.add_kernel_spatial_dimensions(2); in TEST_F()
[all …]
Dir_emission_utils.cc51 const ConvolutionDimensionNumbers& dnums = in PotentiallyImplementedAsEigenConvolution() local
55 const int64 num_spatial_dims = dnums.output_spatial_dimensions_size(); in PotentiallyImplementedAsEigenConvolution()
61 if (dnums.input_spatial_dimensions(i) != i + 1) { in PotentiallyImplementedAsEigenConvolution()
64 if (dnums.kernel_spatial_dimensions(i) != i) { in PotentiallyImplementedAsEigenConvolution()
67 if (dnums.output_spatial_dimensions(i) != i + 1) { in PotentiallyImplementedAsEigenConvolution()
73 return dnums.input_batch_dimension() == 0 && in PotentiallyImplementedAsEigenConvolution()
74 dnums.input_feature_dimension() == input_shape.dimensions_size() - 1 && in PotentiallyImplementedAsEigenConvolution()
75 dnums.output_batch_dimension() == 0 && in PotentiallyImplementedAsEigenConvolution()
76 dnums.output_feature_dimension() == in PotentiallyImplementedAsEigenConvolution()
78 dnums.kernel_input_feature_dimension() == in PotentiallyImplementedAsEigenConvolution()
[all …]
Dconv_canonicalization.cc37 const ConvolutionDimensionNumbers& dnums = in Run() local
39 auto input_batch_dim = dnums.input_batch_dimension(); in Run()
40 auto input_feature_dim = dnums.input_feature_dimension(); in Run()
41 auto kernel_input_feature_dim = dnums.kernel_input_feature_dimension(); in Run()
42 auto kernel_output_feature_dim = dnums.kernel_output_feature_dimension(); in Run()
44 const int64 num_spatial_dims = dnums.output_spatial_dimensions_size(); in Run()
63 new_input_dim_order[i + 1] = dnums.input_spatial_dimensions(i); in Run()
65 input->shape().dimensions(dnums.input_spatial_dimensions(i)); in Run()
82 new_kernel_dim_order[i] = dnums.kernel_spatial_dimensions(i); in Run()
84 kernel->shape().dimensions(dnums.kernel_spatial_dimensions(i)); in Run()
[all …]
Dir_emitter.cc805 const DotDimensionNumbers& dnums = dot->dot_dimension_numbers(); in HandleDot() local
806 if (dnums.lhs_batch_dimensions_size() > 0 || in HandleDot()
807 dnums.rhs_batch_dimensions_size() > 0) { in HandleDot()
811 if (dnums.lhs_contracting_dimensions_size() != 1) { in HandleDot()
817 if (dnums.lhs_contracting_dimensions(0) != in HandleDot()
819 dnums.rhs_contracting_dimensions(0) != 0) { in HandleDot()
854 const ConvolutionDimensionNumbers& dnums = in HandleConvolution() local
874 const ConvolutionDimensionNumbers& dnums = in HandleConvolution() local
879 int64 input_batch = input_shape.dimensions(dnums.input_batch_dimension()); in HandleConvolution()
881 input_shape.dimensions(dnums.input_spatial_dimensions(0)); in HandleConvolution()
[all …]
/external/tensorflow/tensorflow/compiler/xla/tests/
Dconvolution_test.cc430 ConvolutionDimensionNumbers dnums; in XLA_TEST_F() local
431 dnums.set_input_batch_dimension(0); in XLA_TEST_F()
432 dnums.set_output_batch_dimension(0); in XLA_TEST_F()
433 dnums.add_input_spatial_dimensions(1); in XLA_TEST_F()
434 dnums.add_output_spatial_dimensions(1); in XLA_TEST_F()
435 dnums.add_input_spatial_dimensions(2); in XLA_TEST_F()
436 dnums.add_output_spatial_dimensions(2); in XLA_TEST_F()
437 dnums.add_input_spatial_dimensions(3); in XLA_TEST_F()
438 dnums.add_output_spatial_dimensions(3); in XLA_TEST_F()
439 dnums.set_input_feature_dimension(4); in XLA_TEST_F()
[all …]
Dconvolution_variants_test.cc981 ConvolutionDimensionNumbers dnums; in XLA_TEST_F() local
983 dnums.set_input_batch_dimension(0); in XLA_TEST_F()
984 dnums.set_output_batch_dimension(0); in XLA_TEST_F()
985 dnums.add_input_spatial_dimensions(1); in XLA_TEST_F()
986 dnums.add_output_spatial_dimensions(1); in XLA_TEST_F()
987 dnums.add_input_spatial_dimensions(2); in XLA_TEST_F()
988 dnums.add_output_spatial_dimensions(2); in XLA_TEST_F()
989 dnums.set_input_feature_dimension(3); in XLA_TEST_F()
990 dnums.set_output_feature_dimension(3); in XLA_TEST_F()
993 dnums.add_kernel_spatial_dimensions(0); in XLA_TEST_F()
[all …]
Ddot_operation_test.cc534 DotDimensionNumbers dnums; in XLA_TEST_F() local
535 dnums.add_lhs_contracting_dimensions(2); in XLA_TEST_F()
536 dnums.add_rhs_contracting_dimensions(1); in XLA_TEST_F()
537 dnums.add_lhs_batch_dimensions(0); in XLA_TEST_F()
538 dnums.add_rhs_batch_dimensions(0); in XLA_TEST_F()
540 auto out = builder.DotGeneral(x, y, dnums); in XLA_TEST_F()
/external/tensorflow/tensorflow/compiler/xla/service/
Dtranspose_folding_test.cc225 auto dnums = ComputationBuilder::CreateDefaultConvDimensionNumbers(); in TEST_F() local
235 transpose_y->shape().dimensions(dnums.kernel_spatial_dimensions(i))); in TEST_F()
238 x->shape(), transpose_y->shape(), window, dnums); in TEST_F()
241 conv_shape.ValueOrDie(), x, transpose_y, window, dnums)); in TEST_F()
258 EXPECT_EQ(dnums.kernel_input_feature_dimension(), in TEST_F()
261 EXPECT_EQ(dnums.kernel_output_feature_dimension(), in TEST_F()
278 auto dnums = ComputationBuilder::CreateDefaultConvDimensionNumbers(); in TEST_F() local
288 transpose_y->shape().dimensions(dnums.kernel_spatial_dimensions(i))); in TEST_F()
291 x->shape(), transpose_y->shape(), window, dnums); in TEST_F()
294 conv_shape.ValueOrDie(), x, transpose_y, window, dnums)); in TEST_F()
[all …]
Ddot_decomposer.cc34 const DotDimensionNumbers& dnums = dot->dot_dimension_numbers(); in DecomposeBatchDot() local
42 CHECK_EQ(dnums.lhs_batch_dimensions_size(), in DecomposeBatchDot()
43 dnums.rhs_batch_dimensions_size()); in DecomposeBatchDot()
44 const int64 num_batch_dims = dnums.lhs_batch_dimensions_size(); in DecomposeBatchDot()
49 CHECK_EQ(lhs_shape.dimensions(dnums.lhs_batch_dimensions(i)), in DecomposeBatchDot()
50 rhs_shape.dimensions(dnums.rhs_batch_dimensions(i))); in DecomposeBatchDot()
51 batch_size *= lhs_shape.dimensions(dnums.lhs_batch_dimensions(i)); in DecomposeBatchDot()
55 CHECK_EQ(1, dnums.lhs_contracting_dimensions_size()); in DecomposeBatchDot()
56 const int64 lhs_contracting_dim_number = dnums.lhs_contracting_dimensions(0); in DecomposeBatchDot()
59 CHECK_EQ(1, dnums.rhs_contracting_dimensions_size()); in DecomposeBatchDot()
[all …]
Dhlo_evaluator_test.cc765 ConvolutionDimensionNumbers dnums; in TEST_P() local
766 dnums.set_input_batch_dimension(0); in TEST_P()
767 dnums.set_output_batch_dimension(0); in TEST_P()
768 dnums.set_input_feature_dimension(1); in TEST_P()
769 dnums.set_output_feature_dimension(1); in TEST_P()
770 dnums.add_input_spatial_dimensions(2); in TEST_P()
771 dnums.add_output_spatial_dimensions(2); in TEST_P()
773 dnums.set_kernel_output_feature_dimension(0); in TEST_P()
774 dnums.set_kernel_input_feature_dimension(1); in TEST_P()
775 dnums.add_kernel_spatial_dimensions(2); in TEST_P()
[all …]
Dshape_inference_test.cc391 ConvolutionDimensionNumbers dnums; in TEST_F() local
395 dnums.set_input_batch_dimension(0); in TEST_F()
396 dnums.set_output_batch_dimension(0); in TEST_F()
397 dnums.set_input_feature_dimension(1); in TEST_F()
398 dnums.set_output_feature_dimension(1); in TEST_F()
399 dnums.add_input_spatial_dimensions(2); in TEST_F()
400 dnums.add_output_spatial_dimensions(2); in TEST_F()
401 dnums.add_input_spatial_dimensions(3); in TEST_F()
402 dnums.add_output_spatial_dimensions(3); in TEST_F()
406 dnums.set_kernel_input_feature_dimension(2); in TEST_F()
[all …]
Dtranspose_folding.cc108 const ConvolutionDimensionNumbers& dnums = in FoldTransposeIntoConvolution() local
110 ConvolutionDimensionNumbers new_dnums = dnums; in FoldTransposeIntoConvolution()
124 transpose_dimensions[dnums.input_batch_dimension()]); in FoldTransposeIntoConvolution()
126 transpose_dimensions[dnums.input_feature_dimension()]); in FoldTransposeIntoConvolution()
148 transpose_dimensions[dnums.kernel_input_feature_dimension()]); in FoldTransposeIntoConvolution()
150 transpose_dimensions[dnums.kernel_output_feature_dimension()]); in FoldTransposeIntoConvolution()
Dshape_inference.cc1593 const ConvolutionDimensionNumbers& dnums) { in InferConvolveShape() argument
1603 if (dnums.input_spatial_dimensions_size() != in InferConvolveShape()
1604 dnums.kernel_spatial_dimensions_size()) { in InferConvolveShape()
1611 const int num_spatial_dims = dnums.input_spatial_dimensions_size(); in InferConvolveShape()
1616 window.DebugString().c_str(), dnums.DebugString().c_str()); in InferConvolveShape()
1638 input_dnums[0] = dnums.input_batch_dimension(); in InferConvolveShape()
1639 input_dnums[1] = dnums.input_feature_dimension(); in InferConvolveShape()
1640 std::copy(dnums.input_spatial_dimensions().begin(), in InferConvolveShape()
1641 dnums.input_spatial_dimensions().end(), input_dnums.begin() + 2); in InferConvolveShape()
1645 window_dnums[0] = dnums.kernel_input_feature_dimension(); in InferConvolveShape()
[all …]
Dhlo_evaluator.cc902 const auto& dnums = conv->convolution_dimension_numbers(); in HandleConvolution() local
903 const int64 num_spatial_dims = dnums.output_spatial_dimensions_size(); in HandleConvolution()
904 CHECK_EQ(num_spatial_dims, dnums.input_spatial_dimensions_size()); in HandleConvolution()
905 CHECK_EQ(num_spatial_dims, dnums.kernel_spatial_dimensions_size()); in HandleConvolution()
917 window, dnums)); in HandleConvolution()
927 const int64 input_batch_dim = dnums.input_batch_dimension(); in HandleConvolution()
928 const int64 input_z_dim = dnums.input_feature_dimension(); in HandleConvolution()
930 const int64 kernel_input_z_dim = dnums.kernel_input_feature_dimension(); in HandleConvolution()
931 const int64 kernel_output_z_dim = dnums.kernel_output_feature_dimension(); in HandleConvolution()
933 const int64 output_batch_dim = dnums.output_batch_dimension(); in HandleConvolution()
[all …]
Dalgebraic_simplifier_test.cc889 ConvolutionDimensionNumbers dnums; in TEST_F() local
890 dnums.set_input_batch_dimension(0); in TEST_F()
891 dnums.add_input_spatial_dimensions(1); in TEST_F()
892 dnums.set_input_feature_dimension(2); in TEST_F()
894 dnums.set_output_batch_dimension(0); in TEST_F()
895 dnums.add_output_spatial_dimensions(1); in TEST_F()
896 dnums.set_output_feature_dimension(2); in TEST_F()
898 dnums.add_kernel_spatial_dimensions(0); in TEST_F()
899 dnums.set_kernel_input_feature_dimension(1); in TEST_F()
900 dnums.set_kernel_output_feature_dimension(2); in TEST_F()
[all …]
Dalgebraic_simplifier.cc782 const DotDimensionNumbers& dnums = dot->dot_dimension_numbers(); in OptimizeDotOfConcat() local
783 if (dnums.lhs_contracting_dimensions_size() != 1 || in OptimizeDotOfConcat()
784 dnums.lhs_batch_dimensions_size() != 0) { in OptimizeDotOfConcat()
788 const int64 lhs_contracting_dim = dnums.lhs_contracting_dimensions(0); in OptimizeDotOfConcat()
789 const int64 rhs_contracting_dim = dnums.rhs_contracting_dimensions(0); in OptimizeDotOfConcat()
1884 const ConvolutionDimensionNumbers& dnums = in HandleConvolution() local
1894 for (int64 i = 0; i < dnums.kernel_spatial_dimensions_size(); ++i) { in HandleConvolution()
1895 if (filter_shape.dimensions(dnums.kernel_spatial_dimensions(i)) != 1) { in HandleConvolution()
1924 dnums.input_feature_dimension() || in HandleConvolution()
1926 dnums.output_feature_dimension() || in HandleConvolution()
[all …]
/external/tensorflow/tensorflow/compiler/xla/service/gpu/
Dcudnn_convolution_runner.cc79 const ConvolutionDimensionNumbers& dnums, AlgorithmConfig algorithm, in RunCudnnConvolution() argument
89 VLOG(3) << "Dim nums: { " << dnums.ShortDebugString() << " }"; in RunCudnnConvolution()
109 CHECK_EQ(num_dimensions, dnums.input_spatial_dimensions_size()); in RunCudnnConvolution()
110 CHECK_EQ(num_dimensions, dnums.kernel_spatial_dimensions_size()); in RunCudnnConvolution()
111 CHECK_EQ(num_dimensions, dnums.output_spatial_dimensions_size()); in RunCudnnConvolution()
121 input_shape.dimensions(dnums.input_feature_dimension())) in RunCudnnConvolution()
122 .set_count(input_shape.dimensions(dnums.input_batch_dimension())); in RunCudnnConvolution()
127 input_shape.dimensions(dnums.input_spatial_dimensions(dim))); in RunCudnnConvolution()
133 filter_shape.dimensions(dnums.kernel_input_feature_dimension())) in RunCudnnConvolution()
135 filter_shape.dimensions(dnums.kernel_output_feature_dimension())); in RunCudnnConvolution()
[all …]
Dcudnn_convolution_rewriter.cc38 const ConvolutionDimensionNumbers& dnums = in CanImplementAsCudnnForwardConv() local
40 if (dnums.input_spatial_dimensions_size() > 3) { in CanImplementAsCudnnForwardConv()
226 ConvolutionDimensionNumbers dnums = conv->convolution_dimension_numbers(); in MatchBackwardInput() local
227 const auto& kernel_spatial_dims = dnums.kernel_spatial_dimensions(); in MatchBackwardInput()
274 const auto& input_spatial_dims = dnums.input_spatial_dimensions(); in MatchBackwardInput()
275 const auto& output_spatial_dims = dnums.output_spatial_dimensions(); in MatchBackwardInput()
387 dnums.set_kernel_input_feature_dimension( in MatchBackwardInput()
389 dnums.set_kernel_output_feature_dimension( in MatchBackwardInput()
392 return std::make_tuple(true, new_window, dnums); in MatchBackwardInput()
402 ConvolutionDimensionNumbers dnums; in RunOnInstruction() local
[all …]
Dir_emission_utils.cc134 const ConvolutionDimensionNumbers& dnums) { in CreateCudnnConv() argument
157 custom_call->set_convolution_dimension_numbers(dnums); in CreateCudnnConv()
163 const Window& window, const ConvolutionDimensionNumbers& dnums) { in CreateCudnnConvForward() argument
165 window, dnums); in CreateCudnnConvForward()
170 const Window& window, const ConvolutionDimensionNumbers& dnums) { in CreateCudnnConvBackwardInput() argument
172 reverse_filter, window, dnums); in CreateCudnnConvBackwardInput()
177 const Window& window, const ConvolutionDimensionNumbers& dnums) { in CreateCudnnConvBackwardFilter() argument
179 output, window, dnums); in CreateCudnnConvBackwardFilter()
Dcudnn_convolution_algorithm_picker.cc101 const ConvolutionDimensionNumbers& dnums) { in ShouldIncludeWinogradNonfusedAlgo() argument
102 int64 batch = input_shape.dimensions(dnums.input_batch_dimension()); in ShouldIncludeWinogradNonfusedAlgo()
103 int64 in_depths = input_shape.dimensions(dnums.input_feature_dimension()); in ShouldIncludeWinogradNonfusedAlgo()
104 int64 in_rows = input_shape.dimensions(dnums.input_spatial_dimensions(0)); in ShouldIncludeWinogradNonfusedAlgo()
106 dnums.input_spatial_dimensions_size() == 1 in ShouldIncludeWinogradNonfusedAlgo()
108 : input_shape.dimensions(dnums.input_spatial_dimensions(1)); in ShouldIncludeWinogradNonfusedAlgo()
109 int64 out_depths = output_shape.dimensions(dnums.output_feature_dimension()); in ShouldIncludeWinogradNonfusedAlgo()
170 const ConvolutionDimensionNumbers& dnums, HloInstruction* instr) { in PickBestAlgorithm() argument
213 ShouldIncludeWinogradNonfusedAlgo(input_shape, output_shape, dnums); in PickBestAlgorithm()
228 dnums, AlgorithmConfig(alg), &stream, &profile_result) in PickBestAlgorithm()
Dir_emission_utils.h108 const Window& window, const ConvolutionDimensionNumbers& dnums);
111 const Window& window, const ConvolutionDimensionNumbers& dnums);
114 const Window& window, const ConvolutionDimensionNumbers& dnums);
Dcudnn_convolution_runner.h79 const ConvolutionDimensionNumbers& dnums,
90 const Window& window, const ConvolutionDimensionNumbers& dnums,
/external/tensorflow/tensorflow/compiler/tf2xla/kernels/
Dconv_ops.cc398 xla::ConvolutionDimensionNumbers dnums; in Compile() local
399 dnums.set_input_batch_dimension(batch_dim); in Compile()
400 dnums.set_output_batch_dimension(batch_dim); in Compile()
401 dnums.set_input_feature_dimension(feature_dim); in Compile()
402 dnums.set_output_feature_dimension(feature_dim); in Compile()
406 dnums.set_kernel_input_feature_dimension(num_spatial_dims_ + 1); in Compile()
407 dnums.set_kernel_output_feature_dimension(num_spatial_dims_); in Compile()
416 dnums.add_input_spatial_dimensions(dim); in Compile()
417 dnums.add_kernel_spatial_dimensions(i); in Compile()
418 dnums.add_output_spatial_dimensions(dim); in Compile()
[all …]
/external/tensorflow/tensorflow/compiler/xla/
Dreference_util.cc104 const ConvolutionDimensionNumbers& dnums) { in ConvArray3DGeneralDimensionsDilated() argument
105 CHECK_EQ(dnums.input_spatial_dimensions_size(), 1); in ConvArray3DGeneralDimensionsDilated()
106 CHECK_EQ(dnums.kernel_spatial_dimensions_size(), 1); in ConvArray3DGeneralDimensionsDilated()
107 CHECK_EQ(dnums.output_spatial_dimensions_size(), 1); in ConvArray3DGeneralDimensionsDilated()
124 ConvolutionDimensionNumbers dnums2d = dnums; in ConvArray3DGeneralDimensionsDilated()
511 ConvolutionDimensionNumbers dnums) { in ConvArray4DGeneralDimensionsDilated() argument
519 if (dnums.kernel_spatial_dimensions(0) > dnums.kernel_spatial_dimensions(1)) { in ConvArray4DGeneralDimensionsDilated()
528 lhs_literal->shape().dimensions(dnums.input_spatial_dimensions(0)); in ConvArray4DGeneralDimensionsDilated()
530 lhs_literal->shape().dimensions(dnums.input_spatial_dimensions(1)); in ConvArray4DGeneralDimensionsDilated()
532 rhs_literal->shape().dimensions(dnums.kernel_spatial_dimensions(0)); in ConvArray4DGeneralDimensionsDilated()
[all …]
/external/tensorflow/tensorflow/compiler/xla/tools/parser/
Dhlo_parser.cc167 bool ParseConvolutionDimensionNumbers(ConvolutionDimensionNumbers* dnums);
686 optional<ConvolutionDimensionNumbers> dnums; in ParseInstruction() local
689 AttrTy::kConvolutionDimensionNumbers, &dnums}; in ParseInstruction()
698 shape, /*lhs=*/operands[0], /*rhs=*/operands[1], *window, *dnums)); in ParseInstruction()
2075 ConvolutionDimensionNumbers* dnums) { in ParseConvolutionDimensionNumbers() argument
2115 dnums->add_input_spatial_dimensions(-1); in ParseConvolutionDimensionNumbers()
2120 dnums->set_input_batch_dimension(i); in ParseConvolutionDimensionNumbers()
2122 dnums->set_input_feature_dimension(i); in ParseConvolutionDimensionNumbers()
2124 dnums->set_input_spatial_dimensions(c - '0', i); in ParseConvolutionDimensionNumbers()
2139 dnums->add_kernel_spatial_dimensions(-1); in ParseConvolutionDimensionNumbers()
[all …]

12