get rid of kernel build failure

* fix out-of-boundary access
opencv · Jan 25, 2021 · 7fda4cb · 7fda4cb
1 parent d733a80
commit 7fda4cb
Show file tree

Hide file tree

Showing 2 changed files with 8 additions and 6 deletions.
diff --git a/modules/rgbd/src/opencl/tsdf.cl b/modules/rgbd/src/opencl/tsdf.cl
@@ -4,8 +4,7 @@
 
 // This code is also subject to the license terms in the LICENSE_KinectFusion.md file found in this module's directory
 
-typedef __INT8_TYPE__ int8_t;
-
+typedef char int8_t;
 typedef int8_t TsdfType;
 typedef uchar WeightType;
 
@@ -28,14 +27,17 @@ static inline float tsdfToFloat(TsdfType num)
 }
 
 __kernel void preCalculationPixNorm (__global float * pixNorms,
+                                     int pix_step, int pix_offset,
+                                     int pix_rows, int pix_cols,
                                      const __global float * xx,
                                      const __global float * yy,
                                      int width)
 {    
     int i = get_global_id(0);
     int j = get_global_id(1);
     int idx = i*width + j;
-    pixNorms[idx] = sqrt(xx[j] * xx[j] + yy[i] * yy[i] + 1.0f);
+    if(i < pix_rows && j < pix_cols)
+        pixNorms[idx] = sqrt(xx[j] * xx[j] + yy[i] * yy[i] + 1.0f);
 }
 
 __kernel void integrate(__global const char * depthptr,
@@ -85,7 +87,7 @@ __kernel void integrate(__global const char * depthptr,
     int volYidx = x*volDims.x + y*volDims.y;
 
     int startZ, endZ;
-    if(fabs(zStep.z) > 1e-5)
+    if(fabs(zStep.z) > 1e-5f)
     {
         int baseZ = convert_int(-basePt.z / zStep.z);
         if(zStep.z > 0)

diff --git a/modules/rgbd/src/tsdf.cpp b/modules/rgbd/src/tsdf.cpp
@@ -856,11 +856,11 @@ static cv::UMat preCalculationPixNormGPU(int depth_rows, int depth_cols, Vec2f f
         throw std::runtime_error("Failed to create kernel: " + errorStr);
 
     AccessFlag af = ACCESS_READ;
-    UMat pixNorm = _pixNorm.getUMat(af);
+    UMat pixNorm = _pixNorm.getUMat(ACCESS_RW);
     UMat xx = x.getUMat(af);
     UMat yy = y.getUMat(af);
 
-    kk.args(ocl::KernelArg::PtrReadWrite(pixNorm),
+    kk.args(ocl::KernelArg::ReadWrite(pixNorm),
         ocl::KernelArg::PtrReadOnly(xx),
         ocl::KernelArg::PtrReadOnly(yy),
         depth_cols);