From f27c01e3fa7d303e17fe68aebe06b95ad9e70ad7 Mon Sep 17 00:00:00 2001
From: Vardan Akopian <vakopian@sustain.net>
Date: Tue, 16 Aug 2016 15:08:06 -0700
Subject: [PATCH 1/2] fix initial value for morph

---
 src/backend/cpu/kernel/morph.hpp  | 7 +++++--
 src/backend/cuda/kernel/morph.hpp | 7 +++++--
 2 files changed, 10 insertions(+), 4 deletions(-)

diff --git a/src/backend/cpu/kernel/morph.hpp b/src/backend/cpu/kernel/morph.hpp
index d990bb873b..0589dd222f 100644
--- a/src/backend/cpu/kernel/morph.hpp
+++ b/src/backend/cpu/kernel/morph.hpp
@@ -8,6 +8,7 @@
  ********************************************************/
 
 #pragma once
+#include <limits>
 #include <Array.hpp>
 #include <utility.hpp>
 
@@ -37,7 +38,8 @@ void morph(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) ];
+                    T filterResult = isDilation ? (std::is_integral<T>::value ? std::numeric_limits<T>::lowest() : -std::numeric_limits<T>::infinity())
+			: (std::is_integral<T>::value ? std::numeric_limits<T>::max() : std::numeric_limits<T>::infinity());
 
                     // wj,wi steps along 2nd & 1st dimensions of filter window respectively
                     for(dim_t wj=0; wj<window[1]; wj++) {
@@ -96,7 +98,8 @@ 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 = isDilation ? (std::is_integral<T>::value ? std::numeric_limits<T>::lowest() : -std::numeric_limits<T>::infinity())
+			: (std::is_integral<T>::value ? std::numeric_limits<T>::max() : std::numeric_limits<T>::infinity());
 
                     // wk, wj,wi steps along 2nd & 1st dimensions of filter window respectively
                     for(dim_t wk=0; wk<window[2]; wk++) {
diff --git a/src/backend/cuda/kernel/morph.hpp b/src/backend/cuda/kernel/morph.hpp
index f92ed01c26..697f8760ec 100644
--- a/src/backend/cuda/kernel/morph.hpp
+++ b/src/backend/cuda/kernel/morph.hpp
@@ -7,6 +7,7 @@
  * http://arrayfire.com/licenses/BSD-3-Clause
  ********************************************************/
 
+#include <limits>
 #include <backend.hpp>
 #include <dispatch.hpp>
 #include <Param.hpp>
@@ -102,7 +103,8 @@ 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 ? (std::is_integral<T>::value ? std::numeric_limits<T>::lowest() : -std::numeric_limits<T>::infinity())
+	: (std::is_integral<T>::value ? std::numeric_limits<T>::max() : std::numeric_limits<T>::infinity());
 #pragma unroll
     for(int wj=0; wj<windLen; ++wj) {
         int joff   = wj*windLen;
@@ -196,7 +198,8 @@ 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 ? (std::is_integral<T>::value ? std::numeric_limits<T>::lowest() : -std::numeric_limits<T>::infinity())
+	: (std::is_integral<T>::value ? std::numeric_limits<T>::max() : std::numeric_limits<T>::infinity());
 #pragma unroll
     for(int wk=0; wk<windLen; ++wk) {
         int koff   = wk*se_area;

From a301136043e3bbee209def3d837f84210bcd55b4 Mon Sep 17 00:00:00 2001
From: Pavan Yalamanchili <pavan@arrayfire.com>
Date: Wed, 17 Aug 2016 13:59:40 -0400
Subject: [PATCH 2/2] Fixing initial values in morph functions for all backends

---
 src/backend/cpu/kernel/morph.hpp    | 11 ++++--
 src/backend/cuda/kernel/morph.hpp   |  7 ++--
 src/backend/opencl/kernel/morph.cl  |  4 +-
 src/backend/opencl/kernel/morph.hpp | 60 ++++++++++++++++-------------
 4 files changed, 46 insertions(+), 36 deletions(-)

diff --git a/src/backend/cpu/kernel/morph.hpp b/src/backend/cpu/kernel/morph.hpp
index 0589dd222f..4cec3b363a 100644
--- a/src/backend/cpu/kernel/morph.hpp
+++ b/src/backend/cpu/kernel/morph.hpp
@@ -11,6 +11,7 @@
 #include <limits>
 #include <Array.hpp>
 #include <utility.hpp>
+#include <ops.hpp>
 
 namespace cpu
 {
@@ -31,6 +32,8 @@ 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
@@ -38,8 +41,7 @@ void morph(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 = isDilation ? (std::is_integral<T>::value ? std::numeric_limits<T>::lowest() : -std::numeric_limits<T>::infinity())
-			: (std::is_integral<T>::value ? std::numeric_limits<T>::max() : std::numeric_limits<T>::infinity());
+                    T filterResult = init;
 
                     // wj,wi steps along 2nd & 1st dimensions of filter window respectively
                     for(dim_t wj=0; wj<window[1]; wj++) {
@@ -90,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) {
@@ -98,8 +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 = isDilation ? (std::is_integral<T>::value ? std::numeric_limits<T>::lowest() : -std::numeric_limits<T>::infinity())
-			: (std::is_integral<T>::value ? std::numeric_limits<T>::max() : std::numeric_limits<T>::infinity());
+                    T filterResult = init;
 
                     // wk, wj,wi steps along 2nd & 1st dimensions of filter window respectively
                     for(dim_t wk=0; wk<window[2]; wk++) {
diff --git a/src/backend/cuda/kernel/morph.hpp b/src/backend/cuda/kernel/morph.hpp
index 697f8760ec..beef1dbbb9 100644
--- a/src/backend/cuda/kernel/morph.hpp
+++ b/src/backend/cuda/kernel/morph.hpp
@@ -13,6 +13,7 @@
 #include <Param.hpp>
 #include <debug_cuda.hpp>
 #include <math.hpp>
+#include <ops.hpp>
 #include "shared.hpp"
 
 namespace cuda
@@ -103,8 +104,7 @@ static __global__ void morphKernel(Param<T> out, CParam<T> in,
     __syncthreads();
 
     const T * d_filt = (const T *)cFilter;
-    T acc = isDilation ? (std::is_integral<T>::value ? std::numeric_limits<T>::lowest() : -std::numeric_limits<T>::infinity())
-	: (std::is_integral<T>::value ? std::numeric_limits<T>::max() : std::numeric_limits<T>::infinity());
+    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;
@@ -198,8 +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 = isDilation ? (std::is_integral<T>::value ? std::numeric_limits<T>::lowest() : -std::numeric_limits<T>::infinity())
-	: (std::is_integral<T>::value ? std::numeric_limits<T>::max() : std::numeric_limits<T>::infinity());
+    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;
diff --git a/src/backend/opencl/kernel/morph.cl b/src/backend/opencl/kernel/morph.cl
index cbbbb84707..4b5a24f9e2 100644
--- a/src/backend/opencl/kernel/morph.cl
+++ b/src/backend/opencl/kernel/morph.cl
@@ -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;
@@ -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;
diff --git a/src/backend/opencl/kernel/morph.hpp b/src/backend/opencl/kernel/morph.hpp
index bde71bd335..fa7b81618a 100644
--- a/src/backend/opencl/kernel/morph.hpp
+++ b/src/backend/opencl/kernel/morph.hpp
@@ -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;
@@ -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,
@@ -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,