Remove redundant code

1. remove threshold_opencv_kernel
2. typo ASCENDC_KERNELS_H
3. add ALIGN_UP macro

Author: MengqingCao

modules/cannops/ascendc_kernels/threshold_opencv_kernel.cpp

@@ -377,11 +377,3 @@ extern "C" __global__ __aicore__ void threshold_opencv(GM_ADDR tilingGM, GM_ADDR
     // Clear tiling GM cache manually. (cce compiler bug)
     dcci(tilingGM, 1);
 }
-
-#ifndef __CCE_KT_TEST__
-void threshold_opencv_kernel(uint32_t blockDim, void* l2ctrl, void* stream, uint8_t* tiling,
-                             uint8_t* x, uint8_t* y)
-{
-    threshold_opencv<<<blockDim, l2ctrl, stream>>>(tiling, x, y);
-}
-#endif

modules/cannops/include/opencv2/ascendc_kernels.hpp

@@ -1,5 +1,5 @@
 #ifndef ASCENDC_KERNELS_H
-#define KERNEL_TILINASCENDC_KERNELS_HG_H
+#define ASCENDC_KERNELS_H
 
 #include "../../ascendc_kernels/kernel_tiling_types.h"
 #include "aclrtlaunch_threshold_opencv.h"

modules/cannops/src/element_operations.cpp

@@ -444,7 +444,7 @@ double threshold(const AscendMat& src, AscendMat& dst, double thresh, double max
         tiling.thresh = thresh;
         // AscendMat memory will be align to 32B, it's safe to set totalLengh a little bigger.
         size_t totalBytes = src.rows * src.cols * src.channels();
-        tiling.totalLength = ((totalBytes + 32) & ~31);
+        tiling.totalLength = ALIGN_UP(totalBytes, 32);
         tiling.threshType = type;
         tiling.dtype = src.depth();
 

modules/cannops/src/precomp.hpp

@@ -11,5 +11,6 @@
 #include "opencv2/cann_interface.hpp"
 #include "opencv2/cann_private.hpp"
 #include "opencv2/ascendc_kernels.hpp"
+#define ALIGN_UP(num, align) (((num) + (align) - 1) & ~((align) - 1))
 
 #endif /* __OPENCV_PRECOMP_H__ */