blob: c91f6343a69268ca687004dbe0ffbb863271a95c (
plain)
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
|
/* Copyright 2017 The TensorFlow Authors. All Rights Reserved.
Licensed under the Apache License, Version 2.0 (the "License");
you may not use this file except in compliance with the License.
You may obtain a copy of the License at
http://www.apache.org/licenses/LICENSE-2.0
Unless required by applicable law or agreed to in writing, software
distributed under the License is distributed on an "AS IS" BASIS,
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and
limitations under the License.
==============================================================================*/
#ifndef TENSORFLOW_COMPILER_XLA_SERVICE_GPU_INSTRUCTION_FUSION_H_
#define TENSORFLOW_COMPILER_XLA_SERVICE_GPU_INSTRUCTION_FUSION_H_
#include "tensorflow/compiler/xla/service/hlo_instruction.h"
#include "tensorflow/compiler/xla/service/instruction_fusion.h"
namespace xla {
namespace gpu {
class GpuInstructionFusion : public InstructionFusion {
public:
explicit GpuInstructionFusion(bool may_duplicate)
: InstructionFusion(GpuInstructionFusion::IsExpensive, may_duplicate) {}
// Maximum number of operands plus outputs allowed on a single fusion node.
// Exposed publicly mainly for tests.
static constexpr int64 kMaxOperandsAndOutputsPerFusion = 64;
// Determines whether the combination of `a` and `b` into a (possibly
// multi-output) fusion would be "too large" -- i.e., have more operands and
// outputs than is allowed.
//
// `ShouldFuse` and `ShouldFuseIntoMultiOutput` call this; it's public so that
// other fusion passes (e.g. GPU multi-output fusion) can also call this.
static bool FusionWouldBeTooLarge(const HloInstruction* a,
const HloInstruction* b);
static bool IsExpensive(const HloInstruction& instruction);
bool ShouldFuse(HloInstruction* consumer, int64 operand_index) override;
bool ShouldFuseIntoMultiOutput(HloInstruction* consumer,
int64 operand_index) override;
HloInstruction::FusionKind ChooseKind(
const HloInstruction* producer, const HloInstruction* consumer) override;
};
} // namespace gpu
} // namespace xla
#endif // TENSORFLOW_COMPILER_XLA_SERVICE_GPU_INSTRUCTION_FUSION_H_
|
