Merge pull request #1392 from hjchen2/ocr_attention

Add less than implamentation and unit test

src/common/types.cpp

@@ -75,6 +75,7 @@ const char *G_OP_TYPE_TOP_K = "top_k";
 const char *G_OP_TYPE_CAST = "cast";
 const char *G_OP_TYPE_LOG = "log";
 const char *G_OP_TYPE_LOD_RESET = "lod_reset";
+const char *G_OP_TYPE_LESS_THAN = "less_than";
 
 const char *G_OP_TYPE_QUANTIZE = "quantize";
 const char *G_OP_TYPE_DEQUANTIZE = "dequantize";
@@ -175,5 +176,6 @@ std::unordered_map<
         {G_OP_TYPE_SEQUENCE_SOFTMAX, {{"X"}, {"Out"}}},
         {G_OP_TYPE_NORM, {{"X"}, {"Out", "Norm"}}},
         {G_OP_TYPE_LOG, {{"X"}, {"Out"}}},
-        {G_OP_TYPE_LOD_RESET, {{"X", "Y"}, {"Out"}}}};
+        {G_OP_TYPE_LOD_RESET, {{"X", "Y"}, {"Out"}}},
+        {G_OP_TYPE_LESS_THAN, {{"X", "Y"}, {"Out"}}}};
 }  // namespace paddle_mobile

src/common/types.h

@@ -158,6 +158,7 @@ extern const char *G_OP_TYPE_TOP_K;
 extern const char *G_OP_TYPE_CAST;
 extern const char *G_OP_TYPE_LOG;
 extern const char *G_OP_TYPE_LOD_RESET;
+extern const char *G_OP_TYPE_LESS_THAN;
 
 extern const char *G_OP_TYPE_QUANTIZE;
 extern const char *G_OP_TYPE_DEQUANTIZE;

src/framework/load_ops.h

@@ -273,3 +273,9 @@ LOAD_OP1(sequence_pool, CPU);
 #ifdef LOG_OP
 LOAD_OP1(log, CPU);
 #endif
+#ifdef LOD_RESET_OP
+LOAD_OP1(lod_reset, CPU);
+#endif
+#ifdef LESS_THAN_OP
+LOAD_OP1(less_than, CPU);
+#endif

src/operators/activation_op.cpp

@@ -17,7 +17,7 @@ limitations under the License. */
 namespace paddle_mobile {
 namespace operators {
 
-#define DEFINE_INFERSHAPE(OpName)                           \
+#define DEFINE_ACTIVATION_INFERSHAPE(OpName)                \
   template <typename Dtype, typename T>                     \
   void OpName##Op<Dtype, T>::InferShape() const {           \
     const auto &input_dims = this->param_.InputX()->dims(); \
@@ -25,20 +25,20 @@ namespace operators {
   }
 
 #ifdef RELU_OP
-DEFINE_INFERSHAPE(Relu);
-DEFINE_INFERSHAPE(Relu6);
+DEFINE_ACTIVATION_INFERSHAPE(Relu);
+DEFINE_ACTIVATION_INFERSHAPE(Relu6);
 #endif  // RELU_OP
 
 #ifdef SIGMOID_OP
-DEFINE_INFERSHAPE(Sigmoid);
+DEFINE_ACTIVATION_INFERSHAPE(Sigmoid);
 #endif  // SIGMOID_OP
 
 #ifdef TANH_OP
-DEFINE_INFERSHAPE(Tanh);
+DEFINE_ACTIVATION_INFERSHAPE(Tanh);
 #endif  // TANH_OP
 
 #ifdef LOG_OP
-DEFINE_INFERSHAPE(Log);
+DEFINE_ACTIVATION_INFERSHAPE(Log);
 #endif  // LOG_OP
 
 }  // namespace operators

src/operators/compare_op.cpp

+/* Copyright (c) 2018 PaddlePaddle 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. */
+
+#include "operators/compare_op.h"
+
+namespace paddle_mobile {
+namespace operators {
+
+#ifdef LESS_THAN_OP
+template <typename Dtype, typename T>
+void LessThanOp<Dtype, T>::InferShape() const {
+  const auto &input_dims = this->param_.input_x_->dims();
+  this->param_.output_->Resize(input_dims);
+}
+#endif  // LESS_THAN_OP
+
+}  // namespace operators
+}  // namespace paddle_mobile
+
+namespace ops = paddle_mobile::operators;
+#ifdef LESS_THAN_OP
+REGISTER_OPERATOR_CPU(less_than, ops::LessThanOp);
+#endif  // LESS_THAN_OP

src/operators/compare_op.h

+/* Copyright (c) 2018 PaddlePaddle 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. */
+
+#pragma once
+
+#include <string>
+#include "framework/operator.h"
+#include "operators/kernel/compare_kernel.h"
+#include "operators/op_param.h"
+
+namespace paddle_mobile {
+namespace operators {
+
+#ifdef LESS_THAN_OP
+DECLARE_OPERATOR(LessThan, CompareParam, LessThanKernel);
+#endif  // LESS_THAN_OP
+
+}  // namespace operators
+}  // namespace paddle_mobile

src/operators/kernel/arm/compare_kernel.cpp

+/* Copyright (c) 2018 PaddlePaddle 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. */
+
+#include "operators/kernel/compare_kernel.h"
+#if defined(__ARM_NEON__) || defined(__ARM_NEON)
+#include <arm_neon.h>
+#endif
+
+namespace paddle_mobile {
+namespace operators {
+
+typedef enum {
+  LESS_THAN = 0,
+  LESS_EQUAL = 1,
+  GREATER_THAN = 2,
+  GREATER_EQUAL = 3,
+  EQUAL = 4,
+  NOT_EQUAL = 5,
+} CompareType;
+
+#if defined(__ARM_NEON__) || defined(__ARM_NEON)
+template <CompareType Comp = LESS_THAN>
+inline uint32x4_t vcmpq_f32(const float32x4_t x, const float32x4_t y) {
+  return vcleq_f32(x, y);
+}
+#endif
+
+template <CompareType Comp = LESS_THAN>
+inline uint8_t Compare(const float x, const float y) {
+  return static_cast<uint8_t>(x < y);
+}
+
+template <CompareType Comp = LESS_THAN>
+inline uint8_t Compare(const int64_t x, const int64_t y) {
+  return static_cast<uint8_t>(x < y);
+}
+
+template <typename Dtype, CompareType Comp>
+struct CompareCompute {
+  void operator()(const Tensor *X, const Tensor *Y, const int Axis,
+                  Tensor *Out) {}
+};
+
+template <CompareType Comp>
+struct CompareCompute<float, Comp> {
+  void operator()(const Tensor *X, const Tensor *Y, const int Axis,
+                  Tensor *Out) {
+    const float *x = X->data<float>();
+    const float *y = Y->data<float>();
+    uint8_t *output = reinterpret_cast<uint8_t *>(Out->mutable_data<bool>());
+    const auto &x_dims = X->dims();
+    const auto &y_dims = Y->dims();
+    /// axis = -1 represent the last dimensions.
+    int axis = (Axis == -1 ? x_dims.size() - y_dims.size() : Axis);
+    int batch = 1;
+    int channels = 1;
+    int elementwise_num = 1;
+    for (int i = 0; i < axis; ++i) {
+      batch *= x_dims[i];
+    }
+    for (int i = 0; i < y_dims.size(); ++i) {
+      channels *= y_dims[i];
+    }
+    for (int i = y_dims.size() + axis; i < x_dims.size(); ++i) {
+      elementwise_num *= x_dims[i];
+    }
+    // if elementwise_num == 1, compare rowwise
+    if (elementwise_num == 1) {
+      int remain_start = 0;
+#if defined(__ARM_NEON__) || defined(__ARM_NEON)
+      remain_start = channels & 0xfff8;
+      uint8x8_t __mask = vdup_n_u8(0x1);
+      for (int i = 0; i < batch; ++i) {
+        for (int j = 0; j < channels - 7; j += 8) {
+          int x_offset = i * channels + j;
+          float32x4_t __x0 = vld1q_f32(x + x_offset);
+          float32x4_t __x1 = vld1q_f32(x + x_offset + 4);
+          float32x4_t __y0 = vld1q_f32(y + j);
+          float32x4_t __y1 = vld1q_f32(y + j + 4);
+          uint32x4_t __cmp0 = vcmpq_f32<Comp>(__x0, __y0);
+          uint32x4_t __cmp1 = vcmpq_f32<Comp>(__x1, __y1);
+          uint16x4_t __ncmp0 = vmovn_u32(__cmp0);
+          uint16x4_t __ncmp1 = vmovn_u32(__cmp1);
+          uint16x8_t __ncmp = vcombine_u16(__ncmp0, __ncmp1);
+          uint8x8_t __nncmp = vmovn_u16(__ncmp);
+          __nncmp = vand_u8(__nncmp, __mask);
+          vst1_u8(output + x_offset, __nncmp);
+        }
+      }
+#endif  // __ARM_NEON__
+      for (int i = 0; i < batch; ++i) {
+        for (int j = remain_start; j < channels; ++j) {
+          int x_offset = i * channels + j;
+          output[x_offset] = Compare<Comp>(x[x_offset], y[j]);
+        }
+      }
+    } else {
+      for (int i = 0; i < batch; ++i) {
+        for (int j = 0; j < channels; ++j) {
+          int x_offset = (i * channels + j) * elementwise_num;
+          int y_offset = j * elementwise_num;
+          int remain_start = 0;
+#if defined(__ARM_NEON__) || defined(__ARM_NEON)
+          remain_start = elementwise_num & 0xfff8;
+          uint8x8_t __mask = vdup_n_u8(0x1);
+          for (int k = 0; k < elementwise_num - 7; k += 8) {
+            float32x4_t __x0 = vld1q_f32(x + x_offset);
+            float32x4_t __x1 = vld1q_f32(x + x_offset + 4);
+            float32x4_t __y0 = vld1q_f32(y + y_offset);
+            uint32x4_t __cmp0 = vcmpq_f32<Comp>(__x0, __y0);
+            uint32x4_t __cmp1 = vcmpq_f32<Comp>(__x1, __y0);
+            uint16x4_t __ncmp0 = vmovn_u32(__cmp0);
+            uint16x4_t __ncmp1 = vmovn_u32(__cmp1);
+            uint16x8_t __ncmp = vcombine_u16(__ncmp0, __ncmp1);
+            uint8x8_t __nncmp = vmovn_u16(__ncmp);
+            __nncmp = vand_u8(__nncmp, __mask);
+            vst1_u8(output + x_offset, __nncmp);
+            x_offset += 8;
+            y_offset += 8;
+          }
+#endif  // __ARM_NEON__
+          for (int k = remain_start; k < elementwise_num; ++k) {
+            output[x_offset + k] = Compare<Comp>(x[x_offset + k], y[y_offset]);
+          }
+        }
+      }
+    }
+  }
+};
+
+template <CompareType Comp>
+struct CompareCompute<int64_t, Comp> {
+  void operator()(const Tensor *X, const Tensor *Y, const int Axis,
+                  Tensor *Out) {
+    const int64_t *x = X->data<int64_t>();
+    const int64_t *y = Y->data<int64_t>();
+    uint8_t *output = reinterpret_cast<uint8_t *>(Out->mutable_data<bool>());
+    const auto &x_dims = X->dims();
+    const auto &y_dims = Y->dims();
+    /// axis = -1 represent the last dimensions.
+    int axis = (Axis == -1 ? x_dims.size() - y_dims.size() : Axis);
+    int batch = 1;
+    int channels = 1;
+    int elementwise_num = 1;
+    for (int i = 0; i < axis; ++i) {
+      batch *= x_dims[i];
+    }
+    for (int i = 0; i < y_dims.size(); ++i) {
+      channels *= y_dims[i];
+    }
+    for (int i = y_dims.size() + axis; i < x_dims.size(); ++i) {
+      elementwise_num *= x_dims[i];
+    }
+    // if elementwise_num == 1, compare rowwise
+    if (elementwise_num == 1) {
+      for (int i = 0; i < batch; ++i) {
+        for (int j = 0; j < channels; ++j) {
+          int x_offset = i * channels + j;
+          output[x_offset] = Compare<Comp>(x[x_offset], y[j]);
+        }
+      }
+    } else {
+      for (int i = 0; i < batch; ++i) {
+        for (int j = 0; j < channels; ++j) {
+          int x_offset = (i * channels + j) * elementwise_num;
+          int y_offset = j * elementwise_num;
+          for (int k = 0; k < elementwise_num; ++k) {
+            output[x_offset + k] = Compare<Comp>(x[x_offset + k], y[y_offset]);
+          }
+        }
+      }
+    }
+  }
+};
+
+#ifdef LESS_THAN_OP
+template <>
+bool LessThanKernel<CPU, float>::Init(CompareParam<CPU> *param) {
+  return true;
+}
+
+template <>
+void LessThanKernel<CPU, float>::Compute(const CompareParam<CPU> &param) {
+  if (param.input_x_->type() == typeid(int64_t)) {
+    CompareCompute<int64_t, LESS_THAN>()(param.input_x_, param.input_y_,
+                                         param.axis_, param.output_);
+  } else if (param.input_x_->type() == typeid(float)) {
+    CompareCompute<float, LESS_THAN>()(param.input_x_, param.input_y_,
+                                       param.axis_, param.output_);
+  } else {
+    PADDLE_MOBILE_THROW_EXCEPTION(
+        "LessThan only support int64_t and float data type.");
+  }
+}
+#endif  // LESS_THAN_OP
+
+}  // namespace operators
+}  // namespace paddle_mobile

src/operators/kernel/compare_kernel.h

+/* Copyright (c) 2018 PaddlePaddle 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. */
+
+#pragma once
+
+#include "framework/operator.h"
+#include "operators/op_param.h"
+
+namespace paddle_mobile {
+namespace operators {
+
+#ifdef LESS_THAN_OP
+DECLARE_KERNEL(LessThan, CompareParam);
+#endif  // LESS_THAN_OP
+
+}  // namespace operators
+}  // namespace paddle_mobile

src/operators/kernel/kernels.h

@@ -21,15 +21,15 @@ namespace paddle_mobile {
 namespace operators {
 
 #ifdef TOP_K_OP
-DECLARE_KERNEL(TopK, TopKParam)
+DECLARE_KERNEL(TopK, TopKParam);
 #endif  // TOP_K_OP
 
 #ifdef CAST_OP
-DECLARE_KERNEL(Cast, CastParam)
+DECLARE_KERNEL(Cast, CastParam);
 #endif  // CAST_OP
 
 #ifdef LOD_RESET_OP
-DECLARE_KERNEL(LodReset, LodResetParam)
+DECLARE_KERNEL(LodReset, LodResetParam);
 #endif  // LOD_RESET_OP
 
 }  // namespace operators

src/operators/lod_reset_op.cpp

@@ -21,7 +21,7 @@ namespace operators {
 
 template <typename Dtype, typename T>
 void LodResetOp<Dtype, T>::InferShape() const {
-  const auto &input_dims = this->param_.input_->dims();
+  const auto &input_dims = this->param_.input_x_->dims();
   this->param_.output_->Resize(input_dims);
 }
 

src/operators/op_param.h

@@ -2856,5 +2856,28 @@ class LodResetParam : public OpParam {
 };
 #endif  // LOD_RESET_OP
 
+#ifdef LESS_THAN_OP
+template <typename Dtype>
+class CompareParam : public OpParam {
+  typedef typename DtypeTensorTrait<Dtype>::gtype GType;
+  typedef typename DtypeTensorTrait<Dtype>::rtype RType;
+
+ public:
+  CompareParam(const VariableNameMap &inputs, const VariableNameMap &outputs,
+               const AttributeMap &attrs, const Scope &scope) {
+    input_x_ = InputXFrom<GType>(inputs, scope);
+    input_y_ = InputYFrom<GType>(inputs, scope);
+    output_ = OutFrom<GType>(outputs, scope);
+    axis_ = OpParam::GetAttr<int>("axis", attrs);
+  }
+
+ public:
+  GType *input_x_;
+  GType *input_y_;
+  GType *output_;
+  int axis_;
+};
+#endif  // LESS_THAN_OP
+
 }  // namespace operators
 }  // namespace paddle_mobile

test/CMakeLists.txt

@@ -250,6 +250,9 @@ if (NOT FOUND_MATCH)
     ADD_EXECUTABLE(test-cast-op operators/test_cast_op.cpp test_helper.h test_include.h)
     target_link_libraries(test-cast-op paddle-mobile)
 
+    ADD_EXECUTABLE(test-less-than-op operators/test_less_than_op.cpp test_helper.h test_include.h)
+    target_link_libraries(test-less-than-op paddle-mobile)
+
     # gen test
     ADD_EXECUTABLE(test-fc-op operators/test_fusion_fc_op.cpp test_helper.h test_include.h)
     target_link_libraries(test-fc-op paddle-mobile)

test/operators/test_less_than_op.cpp

+/* Copyright (c) 2018 PaddlePaddle 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. */
+
+#include <cmath>
+#include <iostream>
+#include "../test_include.h"
+#include "operators/compare_op.h"
+
+namespace paddle_mobile {
+
+template <typename T>
+void LessThan(const framework::Tensor *X, const framework::Tensor *Y,
+              const int Axis, framework::Tensor *Out) {
+  const T *x = X->data<T>();
+  const T *y = Y->data<T>();
+  bool *output = Out->mutable_data<bool>();
+  const auto &x_dims = X->dims();
+  const auto &y_dims = Y->dims();
+  /// axis = -1 represent the last dimensions.
+  int axis = (Axis == -1 ? x_dims.size() - y_dims.size() : Axis);
+  int batch = 1;
+  int channels = 1;
+  int elementwise_num = 1;
+  for (int i = 0; i < axis; ++i) {
+    batch *= x_dims[i];
+  }
+  for (int i = 0; i < y_dims.size(); ++i) {
+    channels *= y_dims[i];
+  }
+  for (int i = y_dims.size() + axis; i < x_dims.size(); ++i) {
+    elementwise_num *= x_dims[i];
+  }
+  // less than
+  for (int i = 0; i < batch; ++i) {
+    for (int j = 0; j < channels; ++j) {
+      int x_offset = (i * channels + j) * elementwise_num;
+      int y_offset = j * elementwise_num;
+      for (int k = 0; k < elementwise_num; ++k) {
+        output[x_offset + k] = (x[x_offset + k] < y[y_offset]);
+      }
+    }
+  }
+}
+
+template <typename T>
+int TestLessThanOp(const std::vector<int> &x_shape,
+                   const std::vector<int> &y_shape, const int axis) {
+  framework::DDim xdims = framework::make_ddim(x_shape);
+  framework::DDim ydims = framework::make_ddim(y_shape);
+  VariableNameMap inputs;
+  VariableNameMap outputs;
+  auto scope = std::make_shared<framework::Scope>();
+  inputs["X"] = std::vector<std::string>({"inputx"});
+  inputs["Y"] = std::vector<std::string>({"inputy"});
+  outputs["Out"] = std::vector<std::string>({"output"});
+
+  auto inputx_var = scope.get()->Var("inputx");
+  auto inputx = inputx_var->template GetMutable<framework::LoDTensor>();
+  SetupTensor<T>(inputx, xdims, static_cast<T>(-100), static_cast<T>(100));
+  auto inputy_var = scope.get()->Var("inputy");
+  auto inputy = inputy_var->template GetMutable<framework::LoDTensor>();
+  SetupTensor<T>(inputy, ydims, static_cast<T>(-100), static_cast<T>(100));
+
+  auto output_var = scope.get()->Var("output");
+
+  framework::AttributeMap attrs;
+  attrs["axis"].Set<int>(axis);
+  auto *op = new operators::LessThanOp<CPU, float>("less_than", inputs, outputs,
+                                                   attrs, scope);
+  op->InferShape();
+  op->Init();
+  op->Run();
+
+  auto output = output_var->template Get<framework::LoDTensor>();
+
+  framework::Tensor output_cmp;
+  bool *output_cmp_data = output_cmp.mutable_data<bool>(output->dims());
+  LessThan<T>(inputx, inputy, axis, &output_cmp);
+
+  const bool *output_data = output->data<bool>();
+  for (int i = 0; i < output->numel(); ++i) {
+    if (output_data[i] != output_cmp_data[i]) {
+      LOG(kLOG_INFO) << "output_data[" << i << "] = " << output_data[i]
+                     << ", output_cmp_data[" << i
+                     << "] = " << output_cmp_data[i];
+      delete op;
+      exit(1);
+    }
+  }
+  delete op;
+  return 0;
+}
+
+}  // namespace paddle_mobile
+
+int main() {
+  paddle_mobile::TestLessThanOp<float>({1, 2, 3}, {1, 2, 3}, 0);
+  paddle_mobile::TestLessThanOp<float>({10, 2, 1}, {10, 2, 1}, 0);
+
+  paddle_mobile::TestLessThanOp<float>({2, 10, 1}, {1, 10, 1}, 1);
+  paddle_mobile::TestLessThanOp<float>({10, 2, 1}, {1, 2, 1}, 1);
+
+  paddle_mobile::TestLessThanOp<int64_t>({1, 2, 3}, {1, 2, 3}, 0);
+  paddle_mobile::TestLessThanOp<int64_t>({10, 2, 1}, {10, 2, 1}, 0);
+
+  paddle_mobile::TestLessThanOp<int64_t>({2, 10, 1}, {1, 10, 1}, 1);
+  paddle_mobile::TestLessThanOp<int64_t>({10, 2, 1}, {1, 2, 1}, 1);
+
+  std::cout << "test less_than op pass." << std::endl;
+  return 0;
+}

tools/build.sh

@@ -173,7 +173,7 @@ build_error() {
 }
 
 if [ $# -lt 1 ]; then
-	echo "error: target missing!"
+    echo "error: target missing!"
     echo "available targets: ios|android"
     echo "sample usage: ./build.sh android"
 else

tools/op.cmake

@@ -278,6 +278,8 @@ if(NOT FOUND_MATCH)
   set(SEQUENCE_SOFTMAX_OP ON)
   set(LOG_OP ON)
   set(TANH_OP ON)
+  set(LOD_RESET_OP ON)
+  set(LESS_THAN_OP ON)
 endif()
 
   # option(BATCHNORM_OP "" ON)
@@ -517,6 +519,12 @@ endif()
 if (LOG_OP)
   add_definitions(-DLOG_OP)
 endif()
+if (LOD_RESET_OP)
+  add_definitions(-DLOD_RESET_OP)
+endif()
+if (LESS_THAN_OP)
+  add_definitions(-DLESS_THAN_OP)
+endif()
 
 if (TANH_OP)
   add_definitions(-DTANH_OP)