diff options
author | 2017-07-23 19:27:12 -0700 | |
---|---|---|
committer | 2017-07-23 19:36:11 -0700 | |
commit | 2661f6841d0ad9ec1381d177a1f9df02e73d001c (patch) | |
tree | 4a9a6df5ea7d198caf3f3f43a1425256e3526ca3 /tensorflow | |
parent | b4fa05012bd961a612639e11d26bf98aec1b19a4 (diff) |
[XLA] Add support for sin(x) transcendental.
PiperOrigin-RevId: 162889962
Diffstat (limited to 'tensorflow')
16 files changed, 59 insertions, 0 deletions
diff --git a/tensorflow/compiler/tests/unary_ops_test.py b/tensorflow/compiler/tests/unary_ops_test.py index 51d8786ce3..ce35eb9197 100644 --- a/tensorflow/compiler/tests/unary_ops_test.py +++ b/tensorflow/compiler/tests/unary_ops_test.py @@ -152,6 +152,16 @@ class UnaryOpsTest(XLATestCase): np.array([[1, 2]], dtype=dtype), expected=np.array([[0, 0.69314718]], dtype=dtype)) + self._assertOpOutputMatchesExpected( + math_ops.sin, + np.array([[1, 2]], dtype=dtype), + expected=np.array([[0.841478, 0.909302]], dtype=dtype)) + + self._assertOpOutputMatchesExpected( + math_ops.cos, + np.array([[1, 2]], dtype=dtype), + expected=np.array([[0.540297, -0.41614]], dtype=dtype)) + # TODO(b/34703906): improve log1p implementation and make tolerance # tighter. self._assertOpOutputMatchesExpected( diff --git a/tensorflow/compiler/tf2xla/kernels/unary_ops.cc b/tensorflow/compiler/tf2xla/kernels/unary_ops.cc index ab96d86ed2..626ddd17d3 100644 --- a/tensorflow/compiler/tf2xla/kernels/unary_ops.cc +++ b/tensorflow/compiler/tf2xla/kernels/unary_ops.cc @@ -45,6 +45,7 @@ namespace { XLAJIT_MAKE_UNARY(Abs, b->Abs(x)); XLAJIT_MAKE_UNARY(Ceil, b->Ceil(x)); XLAJIT_MAKE_UNARY(Cos, b->Cos(x)); +XLAJIT_MAKE_UNARY(Sin, b->Sin(x)); XLAJIT_MAKE_UNARY(Exp, b->Exp(x)); XLAJIT_MAKE_UNARY(Floor, b->Floor(x)); // Returns 0 if x is 0, -1 if x < 0 and 1 if x > 0. diff --git a/tensorflow/compiler/xla/client/computation_builder.cc b/tensorflow/compiler/xla/client/computation_builder.cc index 9810ccc042..212bcd27d2 100644 --- a/tensorflow/compiler/xla/client/computation_builder.cc +++ b/tensorflow/compiler/xla/client/computation_builder.cc @@ -976,6 +976,11 @@ ComputationDataHandle ComputationBuilder::Cos( return UnaryOp(UNOP_COS, operand); } +ComputationDataHandle ComputationBuilder::Sin( + const ComputationDataHandle& operand) { + return UnaryOp(UNOP_SIN, operand); +} + ComputationDataHandle ComputationBuilder::Tanh( const ComputationDataHandle& operand) { return UnaryOp(UNOP_TANH, operand); diff --git a/tensorflow/compiler/xla/client/computation_builder.h b/tensorflow/compiler/xla/client/computation_builder.h index b411346459..94602bd473 100644 --- a/tensorflow/compiler/xla/client/computation_builder.h +++ b/tensorflow/compiler/xla/client/computation_builder.h @@ -513,6 +513,9 @@ class ComputationBuilder { // Enqueues a cosine instruction onto the computation. ComputationDataHandle Cos(const ComputationDataHandle& operand); + // Enqueues a sine instruction onto the computation. + ComputationDataHandle Sin(const ComputationDataHandle& operand); + // Enqueues a tanh instruction onto the computation. ComputationDataHandle Tanh(const ComputationDataHandle& operand); diff --git a/tensorflow/compiler/xla/service/dfs_hlo_visitor.h b/tensorflow/compiler/xla/service/dfs_hlo_visitor.h index 0deedcf184..e6067ae9ea 100644 --- a/tensorflow/compiler/xla/service/dfs_hlo_visitor.h +++ b/tensorflow/compiler/xla/service/dfs_hlo_visitor.h @@ -145,6 +145,9 @@ class DfsHloVisitor { virtual Status HandleCos(HloInstruction* cos, HloInstruction* operand) { return HandleElementwiseUnary(cos, HloOpcode::kCos); } + virtual Status HandleSin(HloInstruction* sin, HloInstruction* operand) { + return HandleElementwiseUnary(sin, HloOpcode::kSin); + } virtual Status HandleTanh(HloInstruction* tanh, HloInstruction* operand) { return HandleElementwiseUnary(tanh, HloOpcode::kTanh); } diff --git a/tensorflow/compiler/xla/service/elemental_ir_emitter.cc b/tensorflow/compiler/xla/service/elemental_ir_emitter.cc index 9c4380b39d..81092e42d5 100644 --- a/tensorflow/compiler/xla/service/elemental_ir_emitter.cc +++ b/tensorflow/compiler/xla/service/elemental_ir_emitter.cc @@ -177,6 +177,10 @@ StatusOr<llvm::Value*> ElementalIrEmitter::EmitFloatUnaryOp( return llvm_ir::EmitCallToIntrinsic(llvm::Intrinsic::cos, {operand_value}, {operand_value->getType()}, ir_builder_); + case HloOpcode::kSin: + return llvm_ir::EmitCallToIntrinsic(llvm::Intrinsic::sin, {operand_value}, + {operand_value->getType()}, + ir_builder_); case HloOpcode::kFloor: return llvm_ir::EmitCallToIntrinsic( llvm::Intrinsic::floor, {operand_value}, {operand_value->getType()}, @@ -788,6 +792,7 @@ llvm_ir::ElementGenerator ElementalIrEmitter::MakeElementGenerator( case HloOpcode::kLog: case HloOpcode::kNegate: case HloOpcode::kSign: + case HloOpcode::kSin: case HloOpcode::kTanh: case HloOpcode::kLogicalNot: return [this, hlo, &operand_to_generator]( diff --git a/tensorflow/compiler/xla/service/gpu/elemental_ir_emitter.cc b/tensorflow/compiler/xla/service/gpu/elemental_ir_emitter.cc index c2dec7ed6a..c03213ab6d 100644 --- a/tensorflow/compiler/xla/service/gpu/elemental_ir_emitter.cc +++ b/tensorflow/compiler/xla/service/gpu/elemental_ir_emitter.cc @@ -211,6 +211,12 @@ StatusOr<llvm::Value*> GpuElementalIrEmitter::EmitFloatUnaryOp( case HloOpcode::kLog: return EmitLibdeviceMathCall("__nv_log", {operand_value}, {input_type}, output_type); + case HloOpcode::kCos: + return EmitLibdeviceMathCall("__nv_cos", {operand_value}, {input_type}, + output_type); + case HloOpcode::kSin: + return EmitLibdeviceMathCall("__nv_sin", {operand_value}, {input_type}, + output_type); case HloOpcode::kTanh: return EmitLibdeviceMathCall("__nv_tanh", {operand_value}, {input_type}, output_type); diff --git a/tensorflow/compiler/xla/service/hlo_graph_dumper.cc b/tensorflow/compiler/xla/service/hlo_graph_dumper.cc index 097a762015..acd26c4e31 100644 --- a/tensorflow/compiler/xla/service/hlo_graph_dumper.cc +++ b/tensorflow/compiler/xla/service/hlo_graph_dumper.cc @@ -360,6 +360,7 @@ string InstructionSequenceGraph( case HloOpcode::kRemainder: case HloOpcode::kSelect: case HloOpcode::kSign: + case HloOpcode::kSin: case HloOpcode::kSlice: case HloOpcode::kSort: case HloOpcode::kSubtract: diff --git a/tensorflow/compiler/xla/service/hlo_instruction.cc b/tensorflow/compiler/xla/service/hlo_instruction.cc index 16440f9a27..f52882cca5 100644 --- a/tensorflow/compiler/xla/service/hlo_instruction.cc +++ b/tensorflow/compiler/xla/service/hlo_instruction.cc @@ -130,6 +130,7 @@ HloInstruction::CreateGetTupleElement(const Shape& shape, case HloOpcode::kLogicalNot: case HloOpcode::kNegate: case HloOpcode::kSign: + case HloOpcode::kSin: case HloOpcode::kSort: case HloOpcode::kTanh: break; @@ -792,6 +793,7 @@ std::unique_ptr<HloInstruction> HloInstruction::CloneWithNewOperands( case HloOpcode::kLogicalNot: case HloOpcode::kNegate: case HloOpcode::kSign: + case HloOpcode::kSin: case HloOpcode::kSort: case HloOpcode::kTanh: CHECK_EQ(new_operands.size(), 1); @@ -1196,6 +1198,7 @@ bool HloInstruction::Identical( case HloOpcode::kRemainder: case HloOpcode::kSelect: case HloOpcode::kSign: + case HloOpcode::kSin: case HloOpcode::kSubtract: case HloOpcode::kTanh: case HloOpcode::kTuple: @@ -1897,6 +1900,8 @@ Status HloInstruction::Visit(DfsHloVisitor* visitor) { return visitor->HandleTanh(this, operands_[0]); case HloOpcode::kCos: return visitor->HandleCos(this, operands_[0]); + case HloOpcode::kSin: + return visitor->HandleSin(this, operands_[0]); case HloOpcode::kIsFinite: return visitor->HandleIsFinite(this, operands_[0]); case HloOpcode::kLogicalNot: @@ -2172,6 +2177,7 @@ bool HloInstruction::IsElementwise() const { case HloOpcode::kNegate: case HloOpcode::kReducePrecision: case HloOpcode::kSign: + case HloOpcode::kSin: case HloOpcode::kTanh: return true; diff --git a/tensorflow/compiler/xla/service/hlo_opcode.cc b/tensorflow/compiler/xla/service/hlo_opcode.cc index ac637dfd3e..3888f757ad 100644 --- a/tensorflow/compiler/xla/service/hlo_opcode.cc +++ b/tensorflow/compiler/xla/service/hlo_opcode.cc @@ -145,6 +145,8 @@ string HloOpcodeString(HloOpcode opcode) { return "send"; case HloOpcode::kSign: return "sign"; + case HloOpcode::kSin: + return "sine"; case HloOpcode::kSlice: return "slice"; case HloOpcode::kSort: diff --git a/tensorflow/compiler/xla/service/hlo_opcode.h b/tensorflow/compiler/xla/service/hlo_opcode.h index 7086dfff42..358e611d57 100644 --- a/tensorflow/compiler/xla/service/hlo_opcode.h +++ b/tensorflow/compiler/xla/service/hlo_opcode.h @@ -87,6 +87,7 @@ enum class HloOpcode { kSelectAndScatter, kSend, kSign, + kSin, kSlice, kSort, kSubtract, diff --git a/tensorflow/compiler/xla/service/instruction_fusion.cc b/tensorflow/compiler/xla/service/instruction_fusion.cc index e7e2e4f9d3..482ab9b94a 100644 --- a/tensorflow/compiler/xla/service/instruction_fusion.cc +++ b/tensorflow/compiler/xla/service/instruction_fusion.cc @@ -69,6 +69,7 @@ namespace xla { case HloOpcode::kReverse: case HloOpcode::kSelect: case HloOpcode::kSign: + case HloOpcode::kSin: case HloOpcode::kSlice: case HloOpcode::kSubtract: case HloOpcode::kTranspose: diff --git a/tensorflow/compiler/xla/service/shape_inference.cc b/tensorflow/compiler/xla/service/shape_inference.cc index 69fee6a621..40206145c8 100644 --- a/tensorflow/compiler/xla/service/shape_inference.cc +++ b/tensorflow/compiler/xla/service/shape_inference.cc @@ -185,6 +185,7 @@ StatusOr<Shape> InferWindowOutputShape(const Shape& base_shape, case UNOP_FLOOR: case UNOP_CEIL: case UNOP_COS: + case UNOP_SIN: case UNOP_EXP: case UNOP_LOG: case UNOP_TANH: diff --git a/tensorflow/compiler/xla/service/user_computation.cc b/tensorflow/compiler/xla/service/user_computation.cc index d9079d70df..3ab780e7d0 100644 --- a/tensorflow/compiler/xla/service/user_computation.cc +++ b/tensorflow/compiler/xla/service/user_computation.cc @@ -64,6 +64,8 @@ HloOpcode UnaryOperationToHloOpcode(UnaryOperation unop) { return HloOpcode::kNegate; case UNOP_SIGN: return HloOpcode::kSign; + case UNOP_SIN: + return HloOpcode::kSin; case UNOP_SORT: return HloOpcode::kSort; case UNOP_TANH: diff --git a/tensorflow/compiler/xla/tests/array_elementwise_ops_test.cc b/tensorflow/compiler/xla/tests/array_elementwise_ops_test.cc index fa6c97bcb8..fb913e200f 100644 --- a/tensorflow/compiler/xla/tests/array_elementwise_ops_test.cc +++ b/tensorflow/compiler/xla/tests/array_elementwise_ops_test.cc @@ -1542,6 +1542,15 @@ XLA_TEST_F(ArrayElementwiseOpTest, CosF32s) { error_spec_); } +XLA_TEST_F(ArrayElementwiseOpTest, SinF32s) { + ComputationBuilder builder(client_, TestName()); + auto a = builder.ConstantR1<float>({3.14159f, 0.0f, 1.570796f, -0.78539f}); + auto result = builder.Sin(a); + + ComputeAndCompareR1<float>(&builder, {0.0f, 0.0f, 1.0f, -0.707107f}, {}, + error_spec_); +} + TEST_F(ArrayElementwiseOpTest, TanhF32s) { ComputationBuilder builder(client_, TestName()); auto a = builder.ConstantR1<float>({-2.5f, 3.14f, 2.25f}); diff --git a/tensorflow/compiler/xla/xla_data.proto b/tensorflow/compiler/xla/xla_data.proto index baecfda1c8..166f31c9fb 100644 --- a/tensorflow/compiler/xla/xla_data.proto +++ b/tensorflow/compiler/xla/xla_data.proto @@ -637,6 +637,9 @@ enum UnaryOperation { // Elementwise, computes the cosine of x. UNOP_COS = 12; + + // Elementwise, computes the sine of x. + UNOP_SIN = 13; } message UnaryOpRequest { |