/external/tensorflow/tensorflow/compiler/xla/tests/ |
D | convolution_test.cc | 422 ConvolutionDimensionNumbers dnums; in XLA_TEST_F() local 423 dnums.set_input_batch_dimension(0); in XLA_TEST_F() 424 dnums.set_output_batch_dimension(0); in XLA_TEST_F() 425 dnums.add_input_spatial_dimensions(1); in XLA_TEST_F() 426 dnums.add_output_spatial_dimensions(1); in XLA_TEST_F() 427 dnums.add_input_spatial_dimensions(2); in XLA_TEST_F() 428 dnums.add_output_spatial_dimensions(2); in XLA_TEST_F() 429 dnums.add_input_spatial_dimensions(3); in XLA_TEST_F() 430 dnums.add_output_spatial_dimensions(3); in XLA_TEST_F() 431 dnums.set_input_feature_dimension(4); in XLA_TEST_F() [all …]
|
D | convolution_variants_test.cc | 980 ConvolutionDimensionNumbers dnums; in XLA_TEST_F() local 982 dnums.set_input_batch_dimension(0); in XLA_TEST_F() 983 dnums.set_output_batch_dimension(0); in XLA_TEST_F() 984 dnums.add_input_spatial_dimensions(1); in XLA_TEST_F() 985 dnums.add_output_spatial_dimensions(1); in XLA_TEST_F() 986 dnums.add_input_spatial_dimensions(2); in XLA_TEST_F() 987 dnums.add_output_spatial_dimensions(2); in XLA_TEST_F() 988 dnums.set_input_feature_dimension(3); in XLA_TEST_F() 989 dnums.set_output_feature_dimension(3); in XLA_TEST_F() 992 dnums.add_kernel_spatial_dimensions(0); in XLA_TEST_F() [all …]
|
D | dot_operation_test.cc | 614 DotDimensionNumbers dnums; in XLA_TYPED_TEST() local 615 dnums.add_lhs_contracting_dimensions(2); in XLA_TYPED_TEST() 616 dnums.add_rhs_contracting_dimensions(1); in XLA_TYPED_TEST() 617 dnums.add_lhs_batch_dimensions(0); in XLA_TYPED_TEST() 618 dnums.add_rhs_batch_dimensions(0); in XLA_TYPED_TEST() 620 DotGeneral(x, y, dnums); in XLA_TYPED_TEST() 651 DotDimensionNumbers dnums; in XLA_TYPED_TEST() local 652 dnums.add_lhs_contracting_dimensions(1); in XLA_TYPED_TEST() 653 dnums.add_rhs_contracting_dimensions(1); in XLA_TYPED_TEST() 654 dnums.add_lhs_batch_dimensions(0); in XLA_TYPED_TEST() [all …]
|
/external/tensorflow/tensorflow/compiler/xla/service/gpu/ |
D | stream_executor_util.cc | 38 StreamExecutorConvLayoutsToXlaLayouts(const ConvolutionDimensionNumbers& dnums, in StreamExecutorConvLayoutsToXlaLayouts() argument 44 input_layout.push_back(dnums.input_batch_dimension()); in StreamExecutorConvLayoutsToXlaLayouts() 45 input_layout.push_back(dnums.input_feature_dimension()); in StreamExecutorConvLayoutsToXlaLayouts() 47 dnums.input_spatial_dimensions().begin(), in StreamExecutorConvLayoutsToXlaLayouts() 48 dnums.input_spatial_dimensions().end()); in StreamExecutorConvLayoutsToXlaLayouts() 51 input_layout.push_back(dnums.input_batch_dimension()); in StreamExecutorConvLayoutsToXlaLayouts() 53 dnums.input_spatial_dimensions().begin(), in StreamExecutorConvLayoutsToXlaLayouts() 54 dnums.input_spatial_dimensions().end()); in StreamExecutorConvLayoutsToXlaLayouts() 55 input_layout.push_back(dnums.input_feature_dimension()); in StreamExecutorConvLayoutsToXlaLayouts() 60 ConvolutionDimensionNumbersToString(dnums)); in StreamExecutorConvLayoutsToXlaLayouts() [all …]
|
D | cudnn_conv_rewriter.cc | 42 const ConvolutionDimensionNumbers& dnums, in CreateCudnnConv() argument 61 custom_call->set_convolution_dimension_numbers(dnums); in CreateCudnnConv() 68 const ConvolutionDimensionNumbers& dnums = in CanImplementAsCudnnForwardConv() local 70 if (dnums.input_spatial_dimensions_size() > 3) { in CanImplementAsCudnnForwardConv() 82 if (dnums.input_spatial_dimensions_size() == 2 in CanImplementAsCudnnForwardConv() 273 ConvolutionDimensionNumbers dnums = conv->convolution_dimension_numbers(); in MatchBackwardInput() local 293 absl::c_is_permutation(dnums.kernel_spatial_dimensions(), in MatchBackwardInput() 327 const auto& input_spatial_dims = dnums.input_spatial_dimensions(); in MatchBackwardInput() 328 const auto& output_spatial_dims = dnums.output_spatial_dimensions(); in MatchBackwardInput() 433 dnums.set_kernel_input_feature_dimension( in MatchBackwardInput() [all …]
|
D | cudnn_conv_runner.cc | 70 const ConvolutionDimensionNumbers* dnums; member 125 const ConvolutionDimensionNumbers& dnums = *params.dnums; in RunCudnnConvImpl() local 137 VLOG(3) << "Dim nums: { " << dnums.ShortDebugString() << " }"; in RunCudnnConvImpl() 156 CHECK_EQ(num_dimensions, dnums.input_spatial_dimensions_size()); in RunCudnnConvImpl() 157 CHECK_EQ(num_dimensions, dnums.kernel_spatial_dimensions_size()); in RunCudnnConvImpl() 158 CHECK_EQ(num_dimensions, dnums.output_spatial_dimensions_size()); in RunCudnnConvImpl() 178 dnums, input_shape.layout(), filter_shape.layout(), in RunCudnnConvImpl() 184 input_shape.dimensions(dnums.input_feature_dimension())) in RunCudnnConvImpl() 185 .set_count(input_shape.dimensions(dnums.input_batch_dimension())); in RunCudnnConvImpl() 190 input_shape.dimensions(dnums.input_spatial_dimensions(dim))); in RunCudnnConvImpl() [all …]
|
D | cudnn_conv_pad_for_tensor_cores.cc | 123 const auto& dnums = conv->convolution_dimension_numbers(); in PadForTensorCores() local 167 new_input_shape->dimensions(dnums.input_feature_dimension()); in PadForTensorCores() 169 new_output_shape->dimensions(dnums.output_feature_dimension()); in PadForTensorCores() 171 new_input_shape->set_dimensions(dnums.input_feature_dimension(), 4); in PadForTensorCores() 172 new_filter_shape->set_dimensions(dnums.kernel_input_feature_dimension(), 4); in PadForTensorCores() 177 pad_dim(new_input_shape, dnums.input_feature_dimension()); in PadForTensorCores() 178 pad_dim(new_filter_shape, dnums.kernel_input_feature_dimension()); in PadForTensorCores() 179 pad_dim(new_filter_shape, dnums.kernel_output_feature_dimension()); in PadForTensorCores() 180 pad_dim(new_output_shape, dnums.output_feature_dimension()); in PadForTensorCores()
|
D | stream_executor_util.h | 36 StreamExecutorConvLayoutsToXlaLayouts(const ConvolutionDimensionNumbers& dnums, 44 XlaConvLayoutsToStreamExecutorLayouts(const ConvolutionDimensionNumbers& dnums,
|
/external/tensorflow/tensorflow/compiler/xla/service/cpu/ |
D | conv_canonicalization_test.cc | 70 ConvolutionDimensionNumbers dnums; in TEST_F() local 71 dnums.set_input_batch_dimension(1); in TEST_F() 72 dnums.set_output_batch_dimension(1); in TEST_F() 73 dnums.add_input_spatial_dimensions(2); in TEST_F() 74 dnums.add_output_spatial_dimensions(2); in TEST_F() 75 dnums.add_input_spatial_dimensions(3); in TEST_F() 76 dnums.add_output_spatial_dimensions(3); in TEST_F() 77 dnums.set_input_feature_dimension(0); in TEST_F() 78 dnums.set_output_feature_dimension(0); in TEST_F() 79 dnums.add_kernel_spatial_dimensions(2); in TEST_F() [all …]
|
D | ir_emission_utils.cc | 83 const ConvolutionDimensionNumbers& dnums = in PotentiallyImplementedAsEigenConvolution() local 87 const int64 num_spatial_dims = dnums.output_spatial_dimensions_size(); in PotentiallyImplementedAsEigenConvolution() 93 if (dnums.input_spatial_dimensions(i) != i + 1) { in PotentiallyImplementedAsEigenConvolution() 96 if (dnums.kernel_spatial_dimensions(i) != i) { in PotentiallyImplementedAsEigenConvolution() 99 if (dnums.output_spatial_dimensions(i) != i + 1) { in PotentiallyImplementedAsEigenConvolution() 104 return dnums.input_batch_dimension() == 0 && in PotentiallyImplementedAsEigenConvolution() 105 dnums.input_feature_dimension() == input_shape.dimensions_size() - 1 && in PotentiallyImplementedAsEigenConvolution() 106 dnums.output_batch_dimension() == 0 && in PotentiallyImplementedAsEigenConvolution() 107 dnums.output_feature_dimension() == in PotentiallyImplementedAsEigenConvolution() 109 dnums.kernel_input_feature_dimension() == in PotentiallyImplementedAsEigenConvolution() [all …]
|
D | conv_canonicalization.cc | 38 const ConvolutionDimensionNumbers& dnums = in Run() local 40 auto input_batch_dim = dnums.input_batch_dimension(); in Run() 41 auto input_feature_dim = dnums.input_feature_dimension(); in Run() 42 auto kernel_input_feature_dim = dnums.kernel_input_feature_dimension(); in Run() 43 auto kernel_output_feature_dim = dnums.kernel_output_feature_dimension(); in Run() 45 const int64 num_spatial_dims = dnums.output_spatial_dimensions_size(); in Run() 64 new_input_dim_order[i + 1] = dnums.input_spatial_dimensions(i); in Run() 66 input->shape().dimensions(dnums.input_spatial_dimensions(i)); in Run() 83 new_kernel_dim_order[i] = dnums.kernel_spatial_dimensions(i); in Run() 85 kernel->shape().dimensions(dnums.kernel_spatial_dimensions(i)); in Run() [all …]
|
/external/tensorflow/tensorflow/compiler/xla/service/ |
D | transpose_folding_test.cc | 230 auto dnums = XlaBuilder::CreateDefaultConvDimensionNumbers(); in TEST_F() local 240 transpose_y->shape().dimensions(dnums.kernel_spatial_dimensions(i))); in TEST_F() 244 /*batch_group_count=*/1, window, dnums); in TEST_F() 248 /*feature_group_count=*/1, /*batch_group_count=*/1, window, dnums, in TEST_F() 266 EXPECT_EQ(dnums.kernel_input_feature_dimension(), in TEST_F() 269 EXPECT_EQ(dnums.kernel_output_feature_dimension(), in TEST_F() 286 auto dnums = XlaBuilder::CreateDefaultConvDimensionNumbers(); in TEST_F() local 296 transpose_y->shape().dimensions(dnums.kernel_spatial_dimensions(i))); in TEST_F() 300 /*batch_group_count=*/1, window, dnums); in TEST_F() 304 /*feature_group_count=*/1, /*batch_group_count=*/1, window, dnums, in TEST_F() [all …]
|
D | dynamic_padder_test.cc | 118 auto dnums = XlaBuilder::CreateDefaultConvDimensionNumbers(0); in TEST_F() local 120 dnums.set_kernel_input_feature_dimension(0); in TEST_F() 121 dnums.set_kernel_output_feature_dimension(1); in TEST_F() 122 dnums.set_input_batch_dimension(0); in TEST_F() 123 dnums.set_output_batch_dimension(1); in TEST_F() 124 dnums.set_output_feature_dimension(0); in TEST_F() 130 /*batch_group_count=*/1, window, dnums, in TEST_F() 161 auto dnums = XlaBuilder::CreateDefaultConvDimensionNumbers(0); in TEST_F() local 163 dnums.set_kernel_input_feature_dimension(0); in TEST_F() 164 dnums.set_kernel_output_feature_dimension(1); in TEST_F() [all …]
|
D | dot_decomposer.cc | 36 const DotDimensionNumbers& dnums = dot->dot_dimension_numbers(); in DecomposeBatchDot() local 44 CHECK_EQ(dnums.lhs_batch_dimensions_size(), in DecomposeBatchDot() 45 dnums.rhs_batch_dimensions_size()); in DecomposeBatchDot() 46 const int64 num_batch_dims = dnums.lhs_batch_dimensions_size(); in DecomposeBatchDot() 51 CHECK_EQ(lhs_shape.dimensions(dnums.lhs_batch_dimensions(i)), in DecomposeBatchDot() 52 rhs_shape.dimensions(dnums.rhs_batch_dimensions(i))); in DecomposeBatchDot() 53 batch_size *= lhs_shape.dimensions(dnums.lhs_batch_dimensions(i)); in DecomposeBatchDot() 57 CHECK_EQ(1, dnums.lhs_contracting_dimensions_size()); in DecomposeBatchDot() 58 const int64 lhs_contracting_dim_number = dnums.lhs_contracting_dimensions(0); in DecomposeBatchDot() 61 CHECK_EQ(1, dnums.rhs_contracting_dimensions_size()); in DecomposeBatchDot() [all …]
|
D | hlo_evaluator_test.cc | 914 ConvolutionDimensionNumbers dnums; in TEST_P() local 915 dnums.set_input_batch_dimension(0); in TEST_P() 916 dnums.set_output_batch_dimension(0); in TEST_P() 917 dnums.set_input_feature_dimension(1); in TEST_P() 918 dnums.set_output_feature_dimension(1); in TEST_P() 919 dnums.add_input_spatial_dimensions(2); in TEST_P() 920 dnums.add_output_spatial_dimensions(2); in TEST_P() 922 dnums.set_kernel_output_feature_dimension(0); in TEST_P() 923 dnums.set_kernel_input_feature_dimension(1); in TEST_P() 924 dnums.add_kernel_spatial_dimensions(2); in TEST_P() [all …]
|
D | shape_inference_test.cc | 391 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 …]
|
D | transpose_folding.cc | 127 const ConvolutionDimensionNumbers& dnums = in FoldTransposeIntoConvolution() local 129 ConvolutionDimensionNumbers new_dnums = dnums; in FoldTransposeIntoConvolution() 142 transpose_dimensions[dnums.input_batch_dimension()]); in FoldTransposeIntoConvolution() 144 transpose_dimensions[dnums.input_feature_dimension()]); in FoldTransposeIntoConvolution() 165 transpose_dimensions[dnums.kernel_input_feature_dimension()]); in FoldTransposeIntoConvolution() 167 transpose_dimensions[dnums.kernel_output_feature_dimension()]); in FoldTransposeIntoConvolution()
|
D | hlo_cost_analysis_test.cc | 167 DotDimensionNumbers dnums; in TEST_F() local 168 dnums.add_lhs_contracting_dimensions(1); in TEST_F() 169 dnums.add_lhs_contracting_dimensions(2); in TEST_F() 170 dnums.add_rhs_contracting_dimensions(0); in TEST_F() 171 dnums.add_rhs_contracting_dimensions(1); in TEST_F() 172 DotGeneral(lhs, rhs, dnums); in TEST_F() 196 DotDimensionNumbers dnums; in TEST_F() local 197 dnums.add_lhs_contracting_dimensions(1); in TEST_F() 198 dnums.add_lhs_batch_dimensions(2); in TEST_F() 199 dnums.add_rhs_contracting_dimensions(0); in TEST_F() [all …]
|
D | shape_inference.cc | 1574 const ConvolutionDimensionNumbers& dnums) { in InferConvolveShape() argument 1602 if (dnums.input_spatial_dimensions_size() != in InferConvolveShape() 1603 dnums.kernel_spatial_dimensions_size()) { in InferConvolveShape() 1607 dnums.DebugString()); in InferConvolveShape() 1610 if (dnums.input_spatial_dimensions_size() != in InferConvolveShape() 1611 dnums.output_spatial_dimensions_size()) { in InferConvolveShape() 1615 dnums.DebugString()); in InferConvolveShape() 1618 const int num_spatial_dims = dnums.input_spatial_dimensions_size(); in InferConvolveShape() 1623 window.DebugString(), dnums.DebugString()); in InferConvolveShape() 1643 input_dnums[0] = dnums.input_batch_dimension(); in InferConvolveShape() [all …]
|
/external/tensorflow/tensorflow/compiler/tf2xla/kernels/ |
D | conv_op_helpers.cc | 351 xla::ConvolutionDimensionNumbers dnums; in MakeXlaBackpropInputConvOp() local 352 dnums.set_input_batch_dimension(batch_dim); in MakeXlaBackpropInputConvOp() 353 dnums.set_output_batch_dimension(batch_dim); in MakeXlaBackpropInputConvOp() 354 dnums.set_input_feature_dimension(feature_dim); in MakeXlaBackpropInputConvOp() 355 dnums.set_output_feature_dimension(feature_dim); in MakeXlaBackpropInputConvOp() 359 dnums.set_kernel_input_feature_dimension(attrs.num_spatial_dims + 1); in MakeXlaBackpropInputConvOp() 360 dnums.set_kernel_output_feature_dimension(attrs.num_spatial_dims); in MakeXlaBackpropInputConvOp() 369 dnums.add_input_spatial_dimensions(dim); in MakeXlaBackpropInputConvOp() 370 dnums.add_kernel_spatial_dimensions(i); in MakeXlaBackpropInputConvOp() 371 dnums.add_output_spatial_dimensions(dim); in MakeXlaBackpropInputConvOp() [all …]
|
D | reverse_sequence_op.cc | 100 xla::GatherDimensionNumbers dnums; in Compile() local 101 dnums.set_index_vector_dim(2); in Compile() 104 dnums.add_start_index_map(batch_dim_); in Compile() 105 dnums.add_start_index_map(seq_dim_); in Compile() 111 dnums.add_offset_dims(i); in Compile() 113 dnums.add_collapsed_slice_dims(i); in Compile() 122 xla::Gather(input, start_indices, dnums, slice_sizes)); in Compile()
|
/external/tensorflow/tensorflow/compiler/tests/ |
D | xla_ops_test.py | 111 dnums = xla_data_pb2.ConvolutionDimensionNumbers() 113 dnums.input_batch_dimension = 0 114 dnums.input_feature_dimension = 1 115 dnums.output_batch_dimension = 0 116 dnums.output_feature_dimension = 1 117 dnums.kernel_output_feature_dimension = 0 118 dnums.kernel_input_feature_dimension = 1 119 dnums.input_spatial_dimensions.extend(range(2, 2 + num_spatial_dims)) 120 dnums.kernel_spatial_dimensions.extend(range(2, 2 + num_spatial_dims)) 121 dnums.output_spatial_dimensions.extend(range(2, 2 + num_spatial_dims)) [all …]
|
/external/tensorflow/tensorflow/compiler/xla/client/ |
D | xla_builder_test.cc | 685 ConvolutionDimensionNumbers dnums; in TEST_F() local 686 dnums.set_input_batch_dimension(0); in TEST_F() 687 dnums.set_output_batch_dimension(0); in TEST_F() 688 dnums.add_input_spatial_dimensions(1); in TEST_F() 689 dnums.add_output_spatial_dimensions(1); in TEST_F() 690 dnums.add_input_spatial_dimensions(2); in TEST_F() 691 dnums.add_output_spatial_dimensions(2); in TEST_F() 692 dnums.set_input_feature_dimension(3); in TEST_F() 693 dnums.set_output_feature_dimension(3); in TEST_F() 694 dnums.add_kernel_spatial_dimensions(0); in TEST_F() [all …]
|
/external/tensorflow/tensorflow/compiler/xla/ |
D | reference_util.cc | 74 const ConvolutionDimensionNumbers& dnums) { in ConvArray3DGeneralDimensionsDilated() argument 75 CHECK_EQ(dnums.input_spatial_dimensions_size(), 1); in ConvArray3DGeneralDimensionsDilated() 76 CHECK_EQ(dnums.kernel_spatial_dimensions_size(), 1); in ConvArray3DGeneralDimensionsDilated() 77 CHECK_EQ(dnums.output_spatial_dimensions_size(), 1); in ConvArray3DGeneralDimensionsDilated() 92 ConvolutionDimensionNumbers dnums2d = dnums; in ConvArray3DGeneralDimensionsDilated() 479 ConvolutionDimensionNumbers dnums) { in ConvArray4DGeneralDimensionsDilated() argument 487 if (dnums.kernel_spatial_dimensions(0) > dnums.kernel_spatial_dimensions(1)) { in ConvArray4DGeneralDimensionsDilated() 496 lhs_literal.shape().dimensions(dnums.input_spatial_dimensions(0)); in ConvArray4DGeneralDimensionsDilated() 498 lhs_literal.shape().dimensions(dnums.input_spatial_dimensions(1)); in ConvArray4DGeneralDimensionsDilated() 500 rhs_literal.shape().dimensions(dnums.kernel_spatial_dimensions(0)); in ConvArray4DGeneralDimensionsDilated() [all …]
|
/external/tensorflow/tensorflow/compiler/xla/client/lib/ |
D | matrix.cc | 195 DotDimensionNumbers dnums; in Einsum() local 209 dnums.add_lhs_batch_dimensions(i); in Einsum() 210 dnums.add_rhs_batch_dimensions(rhs_dimension_number(dim_name)); in Einsum() 212 dnums.add_lhs_contracting_dimensions(i); in Einsum() 213 dnums.add_rhs_contracting_dimensions(rhs_dimension_number(dim_name)); in Einsum() 233 for (auto d : dnums.lhs_batch_dimensions()) { in Einsum() 251 return Transpose(DotGeneral(x, y, dnums, &precision_proto), transpose_dims); in Einsum()
|