tensorflow/compiler/xla/service/gpu/tests/swap_conv_operands_test.cc
// Here, we swap the operands of a convolution to avoid the performance penalty
// associated with convolutions with large padding. This tests that the operands
// are swapped in this case, and that the emitted convolution is sucessfully
// lowered to a cuDNN custom-call.
Here, we swap the operands of a convolution to avoid the performance penalty
associated with convolutions with large padding. This tests that the operands
are swapped in this case, and that the emitted convolution is sucessfully
lowered to a cuDNN custom-call.
Here we swap the operands of a convolution to avoid the performance penalty associated with convolutions with large padding . This tests that the operands are swapped in this case and that the emitted convolution is sucessfully lowered to a cuDNN custom - call .
TEST_F(SwapConvOperandsTest, LargePadding) {
  const char* hlo_text = R"(
HloModule swap_conv
ENTRY swap_conv {
  input = f32[512,128,3,3]{3,2,1,0} parameter(0)
  filter = f32[1,30,30,512]{3,2,1,0}  parameter(1)
  convolution = f32[1,32,32,128]{3,2,1,0} convolution(input, filter), window={size=30x30 pad=29_29x29_29}, dim_labels=fb01_o01i->f01b
  ROOT tuple = (f32[1,32,32,128]{3,2,1,0}) tuple(convolution)
}
)";
  MatchOptimizedHloWithShapes(hlo_text,
                              R"(
// CHECK-LABEL: ENTRY
// CHECK:   [[FILTER:%[^ ]+]] = f32[1,30,30,512]{2,1,3,0}
// CHECK:   [[INPUT:%[^ ]+]] = f32[512,128,3,3]{3,2,0,1}
// CHECK:   %cudnn-conv = (f32[1,32,32,128]{2,1,3,0}, u8[{{[0-9]*}}]{0}) custom-call(f32[1,30,30,512]{2,1,3,0} [[FILTER]], f32[512,128,3,3]{3,2,0,1} [[INPUT]]), window={size=3x3 pad=2_2x2_2 rhs_reversal=1x1}, dim_labels=b01f_io01->b01f
      )");
  EXPECT_TRUE(RunAndCompare(hlo_text, ErrorSpec{1e-5, 1e-5}));
}
TEST_F ( SwapConvOperandsTest , LargePadding ) { const char * hlo_text = R" ( 
HloModule swap_conv

ENTRY swap_conv {
  input = f32[512,128,3,3]{3,2,1,0} parameter(0)
  filter = f32[1,30,30,512]{3,2,1,0}  parameter(1)
  convolution = f32[1,32,32,128]{3,2,1,0} convolution(input, filter), window={size=30x30 pad=29_29x29_29}, dim_labels=fb01_o01i->f01b
  ROOT tuple = (f32[1,32,32,128]{3,2,1,0}) tuple(convolution)
}
 ) " ; MatchOptimizedHloWithShapes ( hlo_text , R" ( 
// CHECK-LABEL: ENTRY
// CHECK:   [[FILTER:%[^ ]+]] = f32[1,30,30,512]{2,1,3,0}
// CHECK:   [[INPUT:%[^ ]+]] = f32[512,128,3,3]{3,2,0,1}
// CHECK:   %cudnn-conv = (f32[1,32,32,128]{2,1,3,0}, u8[{{[0-9]*}}]{0}) custom-call(f32[1,30,30,512]{2,1,3,0} [[FILTER]], f32[512,128,3,3]{3,2,0,1} [[INPUT]]), window={size=3x3 pad=2_2x2_2 rhs_reversal=1x1}, dim_labels=b01f_io01->b01f
       ) " ) ; EXPECT_TRUE ( RunAndCompare ( hlo_text , ErrorSpec { 1e-5 , 1e-5 } ) ) ; }
Here, we swap the operands of a convolution to avoid the performance penalty
associated with convolutions with large padding.
Here we swap the operands of a convolution to avoid the performance penalty associated with convolutions with large padding .