src/opencl/openclwrapper.cpp
static cl_int pixErodeCL_55(l_int32 wpl, l_int32 h) {
  size_t globalThreads[2];
  cl_mem pixtemp;
  cl_int status;
  int gsize;
  l_uint32 fwmask, lwmask;
  size_t localThreads[2];

  lwmask = lmask32[31 - 2];
  fwmask = rmask32[31 - 2];

  // Horizontal pass
  gsize = (wpl * h + GROUPSIZE_HMORX - 1) / GROUPSIZE_HMORX * GROUPSIZE_HMORX;
  globalThreads[0] = gsize;
  globalThreads[1] = GROUPSIZE_HMORY;
  localThreads[0] = GROUPSIZE_HMORX;
  localThreads[1] = GROUPSIZE_HMORY;

  rEnv.mpkKernel = clCreateKernel(rEnv.mpkProgram, "morphoErodeHor_5x5", &status);
  CHECK_OPENCL(status, "clCreateKernel morphoErodeHor_5x5");

  status = clSetKernelArg(rEnv.mpkKernel, 0, sizeof(cl_mem), &pixsCLBuffer);
  status = clSetKernelArg(rEnv.mpkKernel, 1, sizeof(cl_mem), &pixdCLBuffer);
  status = clSetKernelArg(rEnv.mpkKernel, 2, sizeof(wpl), &wpl);
  status = clSetKernelArg(rEnv.mpkKernel, 3, sizeof(h), &h);

  status = clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, rEnv.mpkKernel, 2, nullptr, globalThreads,
                                  localThreads, 0, nullptr, nullptr);

  // Swap source and dest buffers
  pixtemp = pixsCLBuffer;
  pixsCLBuffer = pixdCLBuffer;
  pixdCLBuffer = pixtemp;

  // Vertical
  gsize = (wpl + GROUPSIZE_X - 1) / GROUPSIZE_X * GROUPSIZE_X;
  globalThreads[0] = gsize;
  gsize = (h + GROUPSIZE_Y - 1) / GROUPSIZE_Y * GROUPSIZE_Y;
  globalThreads[1] = gsize;
  localThreads[0] = GROUPSIZE_X;
  localThreads[1] = GROUPSIZE_Y;

  rEnv.mpkKernel = clCreateKernel(rEnv.mpkProgram, "morphoErodeVer_5x5", &status);
  CHECK_OPENCL(status, "clCreateKernel morphoErodeVer_5x5");

  status = clSetKernelArg(rEnv.mpkKernel, 0, sizeof(cl_mem), &pixsCLBuffer);
  status = clSetKernelArg(rEnv.mpkKernel, 1, sizeof(cl_mem), &pixdCLBuffer);
  status = clSetKernelArg(rEnv.mpkKernel, 2, sizeof(wpl), &wpl);
  status = clSetKernelArg(rEnv.mpkKernel, 3, sizeof(h), &h);
  status = clSetKernelArg(rEnv.mpkKernel, 4, sizeof(fwmask), &fwmask);
  status = clSetKernelArg(rEnv.mpkKernel, 5, sizeof(lwmask), &lwmask);
  status = clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, rEnv.mpkKernel, 2, nullptr, globalThreads,
                                  localThreads, 0, nullptr, nullptr);

  return status;
}
// Morphology Erode operation for 5x5 structuring element. Invokes the relevant
// OpenCL kernels
Morphology Erode operation for 5x5 structuring element. Invokes the relevant
OpenCL kernels
Morphology Erode operation for 5x5 structuring element . Invokes the relevant OpenCL kernels
static cl_int pixErodeCL_55(l_int32 wpl, l_int32 h) {
  size_t globalThreads[2];
  cl_mem pixtemp;
  cl_int status;
  int gsize;
  l_uint32 fwmask, lwmask;
  size_t localThreads[2];
  lwmask = lmask32[31 - 2];
  fwmask = rmask32[31 - 2];
  gsize = (wpl * h + GROUPSIZE_HMORX - 1) / GROUPSIZE_HMORX * GROUPSIZE_HMORX;
  globalThreads[0] = gsize;
  globalThreads[1] = GROUPSIZE_HMORY;
  localThreads[0] = GROUPSIZE_HMORX;
  localThreads[1] = GROUPSIZE_HMORY;
  rEnv.mpkKernel = clCreateKernel(rEnv.mpkProgram, "morphoErodeHor_5x5", &status);
  CHECK_OPENCL(status, "clCreateKernel morphoErodeHor_5x5");
  status = clSetKernelArg(rEnv.mpkKernel, 0, sizeof(cl_mem), &pixsCLBuffer);
  status = clSetKernelArg(rEnv.mpkKernel, 1, sizeof(cl_mem), &pixdCLBuffer);
  status = clSetKernelArg(rEnv.mpkKernel, 2, sizeof(wpl), &wpl);
  status = clSetKernelArg(rEnv.mpkKernel, 3, sizeof(h), &h);
  status = clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, rEnv.mpkKernel, 2, nullptr, globalThreads,
                                  localThreads, 0, nullptr, nullptr);
  pixtemp = pixsCLBuffer;
  pixsCLBuffer = pixdCLBuffer;
  pixdCLBuffer = pixtemp;
  gsize = (wpl + GROUPSIZE_X - 1) / GROUPSIZE_X * GROUPSIZE_X;
  globalThreads[0] = gsize;
  gsize = (h + GROUPSIZE_Y - 1) / GROUPSIZE_Y * GROUPSIZE_Y;
  globalThreads[1] = gsize;
  localThreads[0] = GROUPSIZE_X;
  localThreads[1] = GROUPSIZE_Y;
  rEnv.mpkKernel = clCreateKernel(rEnv.mpkProgram, "morphoErodeVer_5x5", &status);
  CHECK_OPENCL(status, "clCreateKernel morphoErodeVer_5x5");
  status = clSetKernelArg(rEnv.mpkKernel, 0, sizeof(cl_mem), &pixsCLBuffer);
  status = clSetKernelArg(rEnv.mpkKernel, 1, sizeof(cl_mem), &pixdCLBuffer);
  status = clSetKernelArg(rEnv.mpkKernel, 2, sizeof(wpl), &wpl);
  status = clSetKernelArg(rEnv.mpkKernel, 3, sizeof(h), &h);
  status = clSetKernelArg(rEnv.mpkKernel, 4, sizeof(fwmask), &fwmask);
  status = clSetKernelArg(rEnv.mpkKernel, 5, sizeof(lwmask), &lwmask);
  status = clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, rEnv.mpkKernel, 2, nullptr, globalThreads,
                                  localThreads, 0, nullptr, nullptr);
  return status;
}
static cl_int pixErodeCL_55 ( l_int32 wpl , l_int32 h ) { size_t globalThreads [ 2 ] ; cl_mem pixtemp ; cl_int status ; int gsize ; l_uint32 fwmask , lwmask ; size_t localThreads [ 2 ] ; lwmask = lmask32 [ 31 - 2 ] ; fwmask = rmask32 [ 31 - 2 ] ; gsize = ( wpl * h + GROUPSIZE_HMORX - 1 ) / GROUPSIZE_HMORX * GROUPSIZE_HMORX ; globalThreads [ 0 ] = gsize ; globalThreads [ 1 ] = GROUPSIZE_HMORY ; localThreads [ 0 ] = GROUPSIZE_HMORX ; localThreads [ 1 ] = GROUPSIZE_HMORY ; rEnv . mpkKernel = clCreateKernel ( rEnv . mpkProgram , " " , & status ) ; CHECK_OPENCL ( status , " " ) ; status = clSetKernelArg ( rEnv . mpkKernel , 0 , sizeof ( cl_mem ) , & pixsCLBuffer ) ; status = clSetKernelArg ( rEnv . mpkKernel , 1 , sizeof ( cl_mem ) , & pixdCLBuffer ) ; status = clSetKernelArg ( rEnv . mpkKernel , 2 , sizeof ( wpl ) , & wpl ) ; status = clSetKernelArg ( rEnv . mpkKernel , 3 , sizeof ( h ) , & h ) ; status = clEnqueueNDRangeKernel ( rEnv . mpkCmdQueue , rEnv . mpkKernel , 2 , nullptr , globalThreads , localThreads , 0 , nullptr , nullptr ) ; pixtemp = pixsCLBuffer ; pixsCLBuffer = pixdCLBuffer ; pixdCLBuffer = pixtemp ; gsize = ( wpl + GROUPSIZE_X - 1 ) / GROUPSIZE_X * GROUPSIZE_X ; globalThreads [ 0 ] = gsize ; gsize = ( h + GROUPSIZE_Y - 1 ) / GROUPSIZE_Y * GROUPSIZE_Y ; globalThreads [ 1 ] = gsize ; localThreads [ 0 ] = GROUPSIZE_X ; localThreads [ 1 ] = GROUPSIZE_Y ; rEnv . mpkKernel = clCreateKernel ( rEnv . mpkProgram , " " , & status ) ; CHECK_OPENCL ( status , " " ) ; status = clSetKernelArg ( rEnv . mpkKernel , 0 , sizeof ( cl_mem ) , & pixsCLBuffer ) ; status = clSetKernelArg ( rEnv . mpkKernel , 1 , sizeof ( cl_mem ) , & pixdCLBuffer ) ; status = clSetKernelArg ( rEnv . mpkKernel , 2 , sizeof ( wpl ) , & wpl ) ; status = clSetKernelArg ( rEnv . mpkKernel , 3 , sizeof ( h ) , & h ) ; status = clSetKernelArg ( rEnv . mpkKernel , 4 , sizeof ( fwmask ) , & fwmask ) ; status = clSetKernelArg ( rEnv . mpkKernel , 5 , sizeof ( lwmask ) , & lwmask ) ; status = clEnqueueNDRangeKernel ( rEnv . mpkCmdQueue , rEnv . mpkKernel , 2 , nullptr , globalThreads , localThreads , 0 , nullptr , nullptr ) ; return status ; }
Morphology Erode operation for 5x5 structuring element.
Morphology Erode operation for 5x5 structuring element .
// Horizontal pass // Swap source and dest buffers // Vertical