message: Run CublasGemmPadForTensorCores before GemmRewriter.
  GemmRewriter will rewrite dots into custom calls. This prevents
  CublasGemmPadForTensorCores from working.
diff: @@ -148,18 +148,23 @@ Status NVPTXCompiler::OptimizeHloConvolutionCanonicalization(
  Status NVPTXCompiler::OptimizeHloPostLayoutAssignment(
  HloModule* hlo_module, se::StreamExecutor* stream_exec,
  se::DeviceMemoryAllocator* device_allocator) {
  - TF_RETURN_IF_ERROR(GpuCompiler::OptimizeHloPostLayoutAssignment(
  - hlo_module, stream_exec, device_allocator));
  -
  - HloPassPipeline pipeline("nvptx post-layout_assignment");
  + HloPassPipeline pre_pipeline("nvptx post-layout_assignment part 1");
  // Pad the dimensions of matrices in dot operations to multiples of 8.
  + // This needs to run before GemmRewriter, which is part of
  + // OptimizeHloPostLayoutAssignment().
  if (IsVoltaOrLater(*stream_exec)) {
  - pipeline.AddPass<CublasGemmPadForTensorCores>();
  + pre_pipeline.AddPass<CublasGemmPadForTensorCores>();
  }
  + TF_RETURN_IF_ERROR(pre_pipeline.Run(hlo_module).status());
  +
  + TF_RETURN_IF_ERROR(GpuCompiler::OptimizeHloPostLayoutAssignment(
  + hlo_module, stream_exec, device_allocator));
  +
  + HloPassPipeline post_pipeline("nvptx post-layout_assignment part 2");
  // Find the fastest algorithm for GEMMs.
  - pipeline.AddPass<GemmAlgorithmPicker>(stream_exec, device_allocator);
  - TF_RETURN_IF_ERROR(pipeline.Run(hlo_module).status());
  + post_pipeline.AddPass<GemmAlgorithmPicker>(stream_exec, device_allocator);
  + TF_RETURN_IF_ERROR(post_pipeline.Run(hlo_module).status());
  return Status::OK();
  }
  