Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

[XLA/GPU] rsqrt is cheap and should be fused. #40998

Merged
merged 2 commits into from
Jul 17, 2020
Merged

[XLA/GPU] rsqrt is cheap and should be fused. #40998

Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
31aeac6
[XLA/GPU] Declares that rsqrt is cheap and should be fused.
trentlo Jun 23, 2020
372cc81
[XLA] Make utility functions anonymous
trentlo Jul 9, 2020
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Jump to
Jump to file
Failed to load files.
Diff view
Diff view
23 changes: 19 additions & 4 deletions tensorflow/compiler/xla/service/gpu/instruction_fusion.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -29,12 +29,27 @@ limitations under the License.
namespace xla {
namespace gpu {

namespace {
bool ElementIsF32OrF16(const Shape& shape) {
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

static or put under anonymous namespace.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

My oversight. Thanks for the catch. Will update it soon,

PrimitiveType type = shape.element_type();
return type == F32 || type == F16;
}
} // namespace

/*static*/ bool GpuInstructionFusion::IsExpensive(
const HloInstruction& instruction) {
// We say that floating-point division is cheap on the GPU.
if (instruction.opcode() == HloOpcode::kDivide &&
ShapeUtil::ElementIsFloating(instruction.shape())) {
return false;

// We say that some floating-point math ops are cheap on the GPU. Unlike other
// intrinsics that can be expanded into many instructions, Div and Rsqrt are
// lowered into single hardware instructions.
switch (instruction.opcode()) {
case HloOpcode::kDivide:
case HloOpcode::kRsqrt:
if (ElementIsF32OrF16(instruction.shape())) {
return false;
}
default:
break;
}
return InstructionFusion::IsExpensive(instruction);
}
Expand Down