tensorflow/compiler/mlir/tensorflow/transforms/tpu_space_to_depth_pass.cc
// Transforms Conv2DBackPropFilter for space to depth.
Transforms Conv2DBackPropFilter for space to depth.
Transforms Conv2DBackPropFilter for space to depth .
void HandleConv2DBackPropFilter(TF::Conv2DBackpropFilterOp backprop,
                                ArrayRef<int32_t> old_filter_shape,
                                ArrayRef<int32_t> new_filter_shape,
                                int64_t block_size) {
  OpBuilder builder(backprop);
  builder.setInsertionPoint(backprop);
  auto input = backprop.input();
  auto new_filter_sizes = builder.create<TF::ConstOp>(
      backprop.getLoc(),
      DenseIntElementsAttr::get(
          RankedTensorType::get({4}, builder.getIntegerType(32)),
          new_filter_shape));
  MLIRContext* context = backprop.getContext();
  SmallVector<int64_t, 4> values = {1, 1, 1, 1};
  auto attrs = llvm::map_range(values, [context](int64_t v) -> Attribute {
    return IntegerAttr::get(IntegerType::get(context, 64), APInt(64, v));
  });
  auto strides = ArrayAttr::get(context, llvm::to_vector<4>(attrs));
  SmallVector<int64_t, 4> new_shape(new_filter_shape.begin(),
                                    new_filter_shape.end());
  auto new_result_type =
      RankedTensorType::get(new_shape, getElementTypeOrSelf(input));
  auto loc = backprop.getLoc();
  auto new_backprop = builder.create<TF::Conv2DBackpropFilterOp>(
      loc, new_result_type, input, new_filter_sizes, backprop.out_backprop(),
      strides, backprop.use_cudnn_on_gpu(), backprop.padding(),
      backprop.explicit_paddings(), backprop.data_format(),
      backprop.dilations());
  SmallVector<int64_t, 6> first_reshape_shape = {
      new_filter_shape[0],
      new_filter_shape[1],
      block_size,
      block_size,
      new_filter_shape[2] / (block_size * block_size),
      new_filter_shape[3]};
  auto first_reshape_op =
      GetReshapeOpForConv2DFilter(first_reshape_shape, new_backprop, &builder);
  auto transpose_op = GetTransposeOpForConv2DFilter(&builder, first_reshape_op);
  SmallVector<int64_t, 4> last_reshape_shape = {
      new_filter_shape[0] * block_size, new_filter_shape[1] * block_size,
      new_filter_shape[2] / (block_size * block_size), new_filter_shape[3]};
  auto final_reshape_op =
      GetReshapeOpForConv2DFilter(last_reshape_shape, transpose_op, &builder);
  auto slice_op = GetSliceOpForConv2DBackPropFilter(old_filter_shape,
                                                    final_reshape_op, &builder);
  backprop.replaceAllUsesWith(slice_op.getResult());
}
void HandleConv2DBackPropFilter ( TF :: Conv2DBackpropFilterOp backprop , ArrayRef < int32_t > old_filter_shape , ArrayRef < int32_t > new_filter_shape , int64_t block_size ) { OpBuilder builder ( backprop ) ; builder . setInsertionPoint ( backprop ) ; auto input = backprop . input ( ) ; auto new_filter_sizes = builder . create < TF :: ConstOp > ( backprop . getLoc ( ) , DenseIntElementsAttr :: get ( RankedTensorType :: get ( { 4 } , builder . getIntegerType ( 32 ) ) , new_filter_shape ) ) ; MLIRContext * context = backprop . getContext ( ) ; SmallVector < int64_t , 4 > values = { 1 , 1 , 1 , 1 } ; auto attrs = llvm :: map_range ( values , [ context ] ( int64_t v ) -> Attribute { return IntegerAttr :: get ( IntegerType :: get ( context , 64 ) , APInt ( 64 , v ) ) ; } ) ; auto strides = ArrayAttr :: get ( context , llvm :: to_vector < 4 > ( attrs ) ) ; SmallVector < int64_t , 4 > new_shape ( new_filter_shape . begin ( ) , new_filter_shape . end ( ) ) ; auto new_result_type = RankedTensorType :: get ( new_shape , getElementTypeOrSelf ( input ) ) ; auto loc = backprop . getLoc ( ) ; auto new_backprop = builder . create < TF :: Conv2DBackpropFilterOp > ( loc , new_result_type , input , new_filter_sizes , backprop . out_backprop ( ) , strides , backprop . use_cudnn_on_gpu ( ) , backprop . padding ( ) , backprop . explicit_paddings ( ) , backprop . data_format ( ) , backprop . dilations ( ) ) ; SmallVector < int64_t , 6 > first_reshape_shape = { new_filter_shape [ 0 ] , new_filter_shape [ 1 ] , block_size , block_size , new_filter_shape [ 2 ] / ( block_size * block_size ) , new_filter_shape [ 3 ] } ; auto first_reshape_op = GetReshapeOpForConv2DFilter ( first_reshape_shape , new_backprop , & builder ) ; auto transpose_op = GetTransposeOpForConv2DFilter ( & builder , first_reshape_op ) ; SmallVector < int64_t , 4 > last_reshape_shape = { new_filter_shape [ 0 ] * block_size , new_filter_shape [ 1 ] * block_size , new_filter_shape [ 2 ] / ( block_size * block_size ) , new_filter_shape [ 3 ] } ; auto final_reshape_op = GetReshapeOpForConv2DFilter ( last_reshape_shape , transpose_op , & builder ) ; auto slice_op = GetSliceOpForConv2DBackPropFilter ( old_filter_shape , final_reshape_op , & builder ) ; backprop . replaceAllUsesWith ( slice_op . getResult ( ) ) ; }
Transforms Conv2DBackPropFilter for space to depth.
Transforms Conv2DBackPropFilter for space to depth .
// Get new filter size from new_filter_shape. // Set stride to [1, 1, 1, 1]. // new result type. // Build new BackPropFilterOp. // For example, if new filter shape is [4, 4, 12, 64], old filter shape // is [7, 7, 3, 64] with block_size 2. // Below transforms will be applied to the filter: // 1. Reshape to [4, 4, 2, 2, 3, 64]; // 2. Transpose to [4, 2, 4, 2, 3, 64]; // 3. Reshape to [8, 8, 3, 64]; // 4. Slice to [7, 7, 3, 64]. // Transpose. // Last Reshape op. // create slice op. // Update backprop's user with the slice op.