Skip to content

Fixes initial values for erode and dilate #1547

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
pavanky merged 2 commits into from
Aug 17, 2016
Merged

Fixes initial values for erode and dilate #1547

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
f27c01e
fix initial value for morph
Aug 16, 2016
a301136
Fixing initial values in morph functions for all backends
pavanky Aug 17, 2016
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
10 changes: 8 additions & 2 deletions src/backend/cpu/kernel/morph.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -8,8 +8,10 @@
********************************************************/

#pragma once
#include <limits>
#include <Array.hpp>
#include <utility.hpp>
#include <ops.hpp>

namespace cpu
{
Expand All @@ -30,14 +32,16 @@ void morph(Array<T> out, Array<T> const in, Array<T> const mask)
const dim_t R0 = window[0]/2;
const dim_t R1 = window[1]/2;

T init = IsDilation ? Binary<T, af_max_t>().init() : Binary<T, af_min_t>().init();

for(dim_t b3=0; b3<dims[3]; ++b3) {
for(dim_t b2=0; b2<dims[2]; ++b2) {
// either channels or batch is handled by outer most loop
for(dim_t j=0; j<dims[1]; ++j) {
// j steps along 2nd dimension
for(dim_t i=0; i<dims[0]; ++i) {
// i steps along 1st dimension
T filterResult = inData[ getIdx(istrides, i, j) ];
T filterResult = init;

// wj,wi steps along 2nd & 1st dimensions of filter window respectively
for(dim_t wj=0; wj<window[1]; wj++) {
Expand Down Expand Up @@ -88,6 +92,8 @@ void morph3d(Array<T> out, Array<T> const in, Array<T> const mask)
const T* inData = in.get();
const T* filter = mask.get();

T init = IsDilation ? Binary<T, af_max_t>().init() : Binary<T, af_min_t>().init();

for(dim_t batchId=0; batchId<bCount; ++batchId) {
// either channels or batch is handled by outer most loop
for(dim_t k=0; k<dims[2]; ++k) {
Expand All @@ -96,7 +102,7 @@ void morph3d(Array<T> out, Array<T> const in, Array<T> const mask)
// j steps along 2nd dimension
for(dim_t i=0; i<dims[0]; ++i) {
// i steps along 1st dimension
T filterResult = inData[ getIdx(istrides, i, j, k) ];
T filterResult = init;

// wk, wj,wi steps along 2nd & 1st dimensions of filter window respectively
for(dim_t wk=0; wk<window[2]; wk++) {
Expand Down
6 changes: 4 additions & 2 deletions src/backend/cuda/kernel/morph.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -7,11 +7,13 @@
* http://arrayfire.com/licenses/BSD-3-Clause
********************************************************/

#include <limits>
#include <backend.hpp>
#include <dispatch.hpp>
#include <Param.hpp>
#include <debug_cuda.hpp>
#include <math.hpp>
#include <ops.hpp>
#include "shared.hpp"

namespace cuda
Expand Down Expand Up @@ -102,7 +104,7 @@ static __global__ void morphKernel(Param<T> out, CParam<T> in,
__syncthreads();

const T * d_filt = (const T *)cFilter;
T acc = shrdMem[ lIdx(i, j, shrdLen, 1) ];
T acc = isDilation ? Binary<T, af_max_t>().init() : Binary<T, af_min_t>().init();
#pragma unroll
for(int wj=0; wj<windLen; ++wj) {
int joff = wj*windLen;
Expand Down Expand Up @@ -196,7 +198,7 @@ static __global__ void morph3DKernel(Param<T> out, CParam<T> in, int nBBS)
int k = lz + halo;

const T * d_filt = (const T *)cFilter;
T acc = shrdMem[ lIdx3D(i, j, k, shrdArea, shrdLen, 1) ];
T acc = isDilation ? Binary<T, af_max_t>().init() : Binary<T, af_min_t>().init();
#pragma unroll
for(int wk=0; wk<windLen; ++wk) {
int koff = wk*se_area;
Expand Down
4 changes: 2 additions & 2 deletions src/backend/opencl/kernel/morph.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -69,7 +69,7 @@ void morph(__global T * out,
int j = ly + halo;
barrier(CLK_LOCAL_MEM_FENCE);

T acc = localMem[ lIdx(i, j, shrdLen, 1) ];
T acc = init;
#pragma unroll
for(int wj=0; wj<windLen; ++wj) {
int joff = wj*windLen;
Expand Down Expand Up @@ -169,7 +169,7 @@ void morph3d(__global T * out,
int j = ly + halo;
int k = lz + halo;

T acc = localMem[ lIdx3D(i, j, k, shrdArea, shrdLen, 1) ];
T acc = init;
#pragma unroll
for(int wk=0; wk<windLen; ++wk) {
int koff = wk*se_area;
Expand Down
60 changes: 34 additions & 26 deletions src/backend/opencl/kernel/morph.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,6 +18,8 @@
#include <Param.hpp>
#include <debug_opencl.hpp>
#include <memory.hpp>
#include <ops.hpp>
#include <type_util.hpp>

using cl::Buffer;
using cl::Program;
Expand Down Expand Up @@ -54,19 +56,22 @@ void morph(Param out,
int device = getActiveDeviceId();

std::call_once( compileFlags[device], [device] () {
std::ostringstream options;
options << " -D T=" << dtype_traits<T>::getName()
<< " -D isDilation="<< isDilation
<< " -D windLen=" << windLen;
if (std::is_same<T, double>::value ||
std::is_same<T, cdouble>::value) {
options << " -D USE_DOUBLE";
}
Program prog;
buildProgram(prog, morph_cl, morph_cl_len, options.str());
morProgs[device] = new Program(prog);
morKernels[device] = new Kernel(*morProgs[device], "morph");
});
ToNum<T> toNum;
T init = isDilation ? Binary<T, af_max_t>().init() : Binary<T, af_min_t>().init();
std::ostringstream options;
options << " -D T=" << dtype_traits<T>::getName()
<< " -D isDilation="<< isDilation
<< " -D init=" << toNum(init)
<< " -D windLen=" << windLen;
if (std::is_same<T, double>::value ||
std::is_same<T, cdouble>::value) {
options << " -D USE_DOUBLE";
}
Program prog;
buildProgram(prog, morph_cl, morph_cl_len, options.str());
morProgs[device] = new Program(prog);
morKernels[device] = new Kernel(*morProgs[device], "morph");
});

auto morphOp = KernelFunctor<Buffer, KParam,
Buffer, KParam,
Expand Down Expand Up @@ -119,19 +124,22 @@ void morph3d(Param out,
int device = getActiveDeviceId();

std::call_once( compileFlags[device], [device] () {
std::ostringstream options;
options << " -D T=" << dtype_traits<T>::getName()
<< " -D isDilation="<< isDilation
<< " -D windLen=" << windLen;
if (std::is_same<T, double>::value ||
std::is_same<T, cdouble>::value) {
options << " -D USE_DOUBLE";
}
Program prog;
buildProgram(prog, morph_cl, morph_cl_len, options.str());
morProgs[device] = new Program(prog);
morKernels[device] = new Kernel(*morProgs[device], "morph3d");
});
ToNum<T> toNum;
T init = isDilation ? Binary<T, af_max_t>().init() : Binary<T, af_min_t>().init();
std::ostringstream options;
options << " -D T=" << dtype_traits<T>::getName()
<< " -D isDilation="<< isDilation
<< " -D init=" << toNum(init)
<< " -D windLen=" << windLen;
if (std::is_same<T, double>::value ||
std::is_same<T, cdouble>::value) {
options << " -D USE_DOUBLE";
}
Program prog;
buildProgram(prog, morph_cl, morph_cl_len, options.str());
morProgs[device] = new Program(prog);
morKernels[device] = new Kernel(*morProgs[device], "morph3d");
});

auto morphOp = KernelFunctor<Buffer, KParam,
Buffer, KParam,
Expand Down