[arrayfire] 31/408: Added AF_INTERP_LOWER and implementation for resize
Ghislain Vaillant
ghisvail-guest at moszumanska.debian.org
Mon Sep 21 19:11:10 UTC 2015
This is an automated email from the git hooks/post-receive script.
ghisvail-guest pushed a commit to branch debian/sid
in repository arrayfire.
commit 2175ccbaade33d60fdfb3a607ed47f02f69cc563
Author: Shehzan Mohammed <shehzan at arrayfire.com>
Date: Tue Jun 23 18:27:39 2015 -0400
Added AF_INTERP_LOWER and implementation for resize
---
include/af/defines.h | 3 ++-
src/api/c/resize.cpp | 5 ++++-
src/backend/cpu/resize.cpp | 30 ++++++++++++++++++++++++++++++
src/backend/cuda/kernel/resize.hpp | 25 +++++++++++++++++++++++++
src/backend/cuda/resize.cu | 3 +++
src/backend/opencl/kernel/resize.cl | 21 +++++++++++++++++++++
src/backend/opencl/kernel/resize.hpp | 3 ++-
src/backend/opencl/resize.cpp | 3 +++
8 files changed, 90 insertions(+), 3 deletions(-)
diff --git a/include/af/defines.h b/include/af/defines.h
index f04da32..db22b5f 100644
--- a/include/af/defines.h
+++ b/include/af/defines.h
@@ -175,7 +175,8 @@ typedef enum {
AF_INTERP_NEAREST, ///< Nearest Interpolation
AF_INTERP_LINEAR, ///< Linear Interpolation
AF_INTERP_BILINEAR, ///< Bilinear Interpolation
- AF_INTERP_CUBIC ///< Cubic Interpolation
+ AF_INTERP_CUBIC, ///< Cubic Interpolation
+ AF_INTERP_LOWER ///< Floor Indexed
} af_interp_type;
typedef enum {
diff --git a/src/api/c/resize.cpp b/src/api/c/resize.cpp
index 4036f9c..419af85 100644
--- a/src/api/c/resize.cpp
+++ b/src/api/c/resize.cpp
@@ -32,7 +32,10 @@ af_err af_resize(af_array *out, const af_array in, const dim_t odim0, const dim_
ArrayInfo info = getInfo(in);
af_dtype type = info.getType();
- ARG_ASSERT(4, (method == AF_INTERP_BILINEAR || method == AF_INTERP_NEAREST));
+ ARG_ASSERT(4, (method == AF_INTERP_BILINEAR ||
+ method == AF_INTERP_NEAREST ||
+ method == AF_INTERP_LOWER
+ ));
DIM_ASSERT(2, odim0 > 0);
DIM_ASSERT(3, odim1 > 0);
diff --git a/src/backend/cpu/resize.cpp b/src/backend/cpu/resize.cpp
index c3b18d0..a4ba43f 100644
--- a/src/backend/cpu/resize.cpp
+++ b/src/backend/cpu/resize.cpp
@@ -129,6 +129,33 @@ namespace cpu
}
};
+ template<typename T>
+ struct resize_op<T, AF_INTERP_LOWER>
+ {
+ void operator()(T *outPtr, const T *inPtr, const af::dim4 &odims, const af::dim4 &idims,
+ const af::dim4 &ostrides, const af::dim4 &istrides,
+ const dim_t x, const dim_t y)
+ {
+ // Compute Indices
+ dim_t i_x = floor((float)x / (odims[0] / (float)idims[0]));
+ dim_t i_y = floor((float)y / (odims[1] / (float)idims[1]));
+
+ if (i_x >= idims[0]) i_x = idims[0] - 1;
+ if (i_y >= idims[1]) i_y = idims[1] - 1;
+
+ dim_t i_off = i_y * istrides[1] + i_x;
+ dim_t o_off = y * ostrides[1] + x;
+ // Copy values from all channels
+ for(dim_t w = 0; w < odims[3]; w++) {
+ dim_t wost = w * ostrides[3];
+ dim_t wist = w * istrides[3];
+ for(dim_t z = 0; z < odims[2]; z++) {
+ outPtr[o_off + z * ostrides[2] + wost] = inPtr[i_off + z * istrides[2] + wist];
+ }
+ }
+ }
+ };
+
template<typename T, af_interp_type method>
void resize_(T *outPtr, const T *inPtr, const af::dim4 &odims, const af::dim4 &idims,
const af::dim4 &ostrides, const af::dim4 &istrides)
@@ -165,6 +192,9 @@ namespace cpu
case AF_INTERP_BILINEAR:
resize_<T, AF_INTERP_BILINEAR>(outPtr, inPtr, odims, idims, ostrides, istrides);
break;
+ case AF_INTERP_LOWER:
+ resize_<T, AF_INTERP_LOWER>(outPtr, inPtr, odims, idims, ostrides, istrides);
+ break;
default:
break;
}
diff --git a/src/backend/cuda/kernel/resize.hpp b/src/backend/cuda/kernel/resize.hpp
index 524ee63..d61c20e 100644
--- a/src/backend/cuda/kernel/resize.hpp
+++ b/src/backend/cuda/kernel/resize.hpp
@@ -120,6 +120,29 @@ namespace cuda
}
///////////////////////////////////////////////////////////////////////////
+ // lower resampling
+ ///////////////////////////////////////////////////////////////////////////
+ template<typename T>
+ __host__ __device__
+ void resize_l(Param<T> out, CParam<T> in,
+ const int o_off, const int i_off,
+ const int blockIdx_x, const int blockIdx_y,
+ const float xf, const float yf)
+ {
+ const int ox = threadIdx.x + blockIdx_x * blockDim.x;
+ const int oy = threadIdx.y + blockIdx_y * blockDim.y;
+
+ int ix = (ox * xf);
+ int iy = (oy * yf);
+
+ if (ox >= out.dims[0] || oy >= out.dims[1]) { return; }
+ if (ix >= in.dims[0]) { ix = in.dims[0] - 1; }
+ if (iy >= in.dims[1]) { iy = in.dims[1] - 1; }
+
+ out.ptr[o_off + ox + oy * out.strides[1]] = in.ptr[i_off + ix + iy * in.strides[1]];
+ }
+
+ ///////////////////////////////////////////////////////////////////////////
// Resize Kernel
///////////////////////////////////////////////////////////////////////////
template<typename T, af_interp_type method>
@@ -140,6 +163,8 @@ namespace cuda
resize_n(out, in, o_off, i_off, blockIdx_x, blockIdx_y, xf, yf);
} else if(method == AF_INTERP_BILINEAR) {
resize_b(out, in, o_off, i_off, blockIdx_x, blockIdx_y, xf, yf);
+ } else if(method == AF_INTERP_LOWER) {
+ resize_l(out, in, o_off, i_off, blockIdx_x, blockIdx_y, xf, yf);
}
}
diff --git a/src/backend/cuda/resize.cu b/src/backend/cuda/resize.cu
index dd54426..dcec972 100644
--- a/src/backend/cuda/resize.cu
+++ b/src/backend/cuda/resize.cu
@@ -31,6 +31,9 @@ namespace cuda
case AF_INTERP_BILINEAR:
kernel::resize<T, AF_INTERP_BILINEAR>(out, in);
break;
+ case AF_INTERP_LOWER:
+ kernel::resize<T, AF_INTERP_LOWER>(out, in);
+ break;
default:
break;
}
diff --git a/src/backend/opencl/kernel/resize.cl b/src/backend/opencl/kernel/resize.cl
index db886bd..f4eaf6f 100644
--- a/src/backend/opencl/kernel/resize.cl
+++ b/src/backend/opencl/kernel/resize.cl
@@ -23,6 +23,7 @@
#define NEAREST resize_n_
#define BILINEAR resize_b_
+#define LOWER resize_l_
////////////////////////////////////////////////////////////////////////////////////
// nearest-neighbor resampling
@@ -86,6 +87,26 @@ void resize_b_(__global T* d_out, const KParam out,
}
////////////////////////////////////////////////////////////////////////////////////
+// lower resampling
+void resize_l_(__global T* d_out, const KParam out,
+ __global const T* d_in, const KParam in,
+ const int blockIdx_x, const int blockIdx_y,
+ const float xf, const float yf)
+{
+ int const ox = get_local_id(0) + blockIdx_x * get_local_size(0);
+ int const oy = get_local_id(1) + blockIdx_y * get_local_size(1);
+
+ int ix = (ox * xf);
+ int iy = (oy * yf);
+
+ if (ox >= out.dims[0] || oy >= out.dims[1]) { return; }
+ if (ix >= in.dims[0]) { ix = in.dims[0] - 1; }
+ if (iy >= in.dims[1]) { iy = in.dims[1] - 1; }
+
+ d_out[ox + oy * out.strides[1]] = d_in[ix + iy * in.strides[1]];
+}
+
+////////////////////////////////////////////////////////////////////////////////////
// Wrapper Kernel
__kernel
void resize_kernel(__global T *d_out, const KParam out,
diff --git a/src/backend/opencl/kernel/resize.hpp b/src/backend/opencl/kernel/resize.hpp
index 5c485ca..53a952b 100644
--- a/src/backend/opencl/kernel/resize.hpp
+++ b/src/backend/opencl/kernel/resize.hpp
@@ -62,8 +62,9 @@ namespace opencl
options << " -D WT=" << dtype_traits<wtype_t<BT>>::getName();
switch(method) {
- case AF_INTERP_NEAREST: options <<" -D INTERP=NEAREST" ; break;
+ case AF_INTERP_NEAREST: options <<" -D INTERP=NEAREST" ; break;
case AF_INTERP_BILINEAR: options <<" -D INTERP=BILINEAR"; break;
+ case AF_INTERP_LOWER: options <<" -D INTERP=LOWER" ; break;
default: break;
}
diff --git a/src/backend/opencl/resize.cpp b/src/backend/opencl/resize.cpp
index d5f358f..10f2735 100644
--- a/src/backend/opencl/resize.cpp
+++ b/src/backend/opencl/resize.cpp
@@ -32,6 +32,9 @@ namespace opencl
case AF_INTERP_BILINEAR:
kernel::resize<T, AF_INTERP_BILINEAR>(out, in);
break;
+ case AF_INTERP_LOWER:
+ kernel::resize<T, AF_INTERP_LOWER>(out, in);
+ break;
default:
break;
}
--
Alioth's /usr/local/bin/git-commit-notice on /srv/git.debian.org/git/debian-science/packages/arrayfire.git
More information about the debian-science-commits
mailing list