[arrayfire] 169/408: Removed uncessary corner sorting for SUSAN
Ghislain Vaillant
ghisvail-guest at moszumanska.debian.org
Mon Sep 21 19:11:48 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 2f5c6723671061eb5b7d5e1812f0d643f06d6a68
Author: pradeep <pradeep at arrayfire.com>
Date: Tue Jul 28 15:21:27 2015 -0400
Removed uncessary corner sorting for SUSAN
---
src/backend/cpu/susan.cpp | 41 +++----------------------------------
src/backend/cuda/kernel/susan.hpp | 30 ---------------------------
src/backend/cuda/susan.cu | 28 +++----------------------
src/backend/opencl/kernel/susan.cl | 18 ----------------
src/backend/opencl/kernel/susan.hpp | 37 ---------------------------------
src/backend/opencl/susan.cpp | 29 +++-----------------------
6 files changed, 9 insertions(+), 174 deletions(-)
diff --git a/src/backend/cpu/susan.cpp b/src/backend/cpu/susan.cpp
index c2d504d..770419c 100644
--- a/src/backend/cpu/susan.cpp
+++ b/src/backend/cpu/susan.cpp
@@ -10,7 +10,6 @@
#include <af/features.h>
#include <Array.hpp>
#include <cmath>
-#include <sort_index.hpp>
#include <math.hpp>
using af::features;
@@ -88,19 +87,6 @@ void non_maximal(float* x_out, float* y_out, float* resp_out,
}
}
-static void keep_corners(float* x_out, float* y_out, float* resp_out,
- const float* x_in, const float* y_in,
- const float* resp_in, const unsigned* resp_idx,
- const unsigned n_corners)
-{
- // Keep only the first n_feat features
- for (unsigned f = 0; f < n_corners; f++) {
- x_out[f] = x_in[resp_idx[f]];
- y_out[f] = y_in[resp_idx[f]];
- resp_out[f] = resp_in[f];
- }
-}
-
template<typename T>
unsigned susan(Array<float> &x_out, Array<float> &y_out, Array<float> &resp_out,
const Array<T> &in,
@@ -128,30 +114,9 @@ unsigned susan(Array<float> &x_out, Array<float> &y_out, Array<float> &resp_out,
if (corners_out == 0)
return 0;
- if (corners_found > corners_out) {
- Array<float> susan_responses = createDeviceDataArray<float>(dim4(corners_found), (void*)resp_corners);
- Array<float> susan_sorted = createEmptyArray<float>(dim4(corners_found));
- Array<unsigned> susan_idx = createEmptyArray<unsigned>(dim4(corners_found));
-
- // Sort susan responses
- sort_index<float, false>(susan_sorted, susan_idx, susan_responses, 0);
-
- x_out = createEmptyArray<float>(dim4(corners_out));
- y_out = createEmptyArray<float>(dim4(corners_out));
- resp_out = createEmptyArray<float>(dim4(corners_out));
-
- // Keep only the corners with higher SUSAN responses
- keep_corners(x_out.get(), y_out.get(), resp_out.get(),
- x_corners, y_corners, susan_sorted.get(), susan_idx.get(),
- corners_out);
-
- memFree(x_corners);
- memFree(y_corners);
- } else {
- x_out = createDeviceDataArray<float>(dim4(corners_out), (void*)x_corners);
- y_out = createDeviceDataArray<float>(dim4(corners_out), (void*)y_corners);
- resp_out = createDeviceDataArray<float>(dim4(corners_out), (void*)resp_corners);
- }
+ x_out = createDeviceDataArray<float>(dim4(corners_out), (void*)x_corners);
+ y_out = createDeviceDataArray<float>(dim4(corners_out), (void*)y_corners);
+ resp_out = createDeviceDataArray<float>(dim4(corners_out), (void*)resp_corners);
return corners_out;
}
diff --git a/src/backend/cuda/kernel/susan.hpp b/src/backend/cuda/kernel/susan.hpp
index addf87a..882adcf 100644
--- a/src/backend/cuda/kernel/susan.hpp
+++ b/src/backend/cuda/kernel/susan.hpp
@@ -176,36 +176,6 @@ void nonMaximal(float* x_out, float* y_out, float* resp_out,
memFree(d_corners_found);
}
-__global__
-void keepCornersKernel(float* x_out, float* y_out, float* resp_out,
- const float* x_in, const float* y_in,
- const float* resp_in, const unsigned* resp_idx,
- const unsigned n_corners)
-{
- const unsigned f = blockDim.x * blockIdx.x + threadIdx.x;
-
- // Keep only the first n_feat features
- if (f < n_corners) {
- x_out[f] = x_in[(unsigned)resp_idx[f]];
- y_out[f] = y_in[(unsigned)resp_idx[f]];
- resp_out[f] = resp_in[f];
- }
-}
-
-void keepCorners(float* x_out, float* y_out, float* resp_out,
- const float* x_in, const float* y_in,
- const float* resp_in, const unsigned* resp_idx,
- const unsigned n_corners)
-{
- dim3 threads(THREADS_PER_BLOCK, 1);
- dim3 blocks(divup(n_corners, threads.x), 1);
-
- keepCornersKernel<<<blocks, threads>>>(x_out, y_out, resp_out,
- x_in, y_in, resp_in, resp_idx, n_corners);
-
- POST_LAUNCH_CHECK();
-}
-
}
}
diff --git a/src/backend/cuda/susan.cu b/src/backend/cuda/susan.cu
index d11e75c..8474454 100644
--- a/src/backend/cuda/susan.cu
+++ b/src/backend/cuda/susan.cu
@@ -11,7 +11,6 @@
#include <Array.hpp>
#include <err_cuda.hpp>
#include <susan.hpp>
-#include <sort_index.hpp>
#include <kernel/susan.hpp>
using af::features;
@@ -46,30 +45,9 @@ unsigned susan(Array<float> &x_out, Array<float> &y_out, Array<float> &resp_out,
if (corners_out == 0)
return 0;
- if (corners_found > corners_out) {
- Array<float> susan_responses = createDeviceDataArray<float>(dim4(corners_found), (void*)resp_corners);
- Array<float> susan_sorted = createEmptyArray<float>(dim4(corners_found));
- Array<unsigned> susan_idx = createEmptyArray<unsigned>(dim4(corners_found));
-
- // Sort susan responses
- sort_index<float, false>(susan_sorted, susan_idx, susan_responses, 0);
-
- x_out = createEmptyArray<float>(dim4(corners_out));
- y_out = createEmptyArray<float>(dim4(corners_out));
- resp_out = createEmptyArray<float>(dim4(corners_out));
-
- // Keep only the corners with higher SUSAN responses
- kernel::keepCorners(x_out.get(), y_out.get(), resp_out.get(),
- x_corners, y_corners, susan_sorted.get(), susan_idx.get(),
- corners_out);
-
- memFree(x_corners);
- memFree(y_corners);
- } else {
- x_out = createDeviceDataArray<float>(dim4(corners_out), (void*)x_corners);
- y_out = createDeviceDataArray<float>(dim4(corners_out), (void*)y_corners);
- resp_out = createDeviceDataArray<float>(dim4(corners_out), (void*)resp_corners);
- }
+ x_out = createDeviceDataArray<float>(dim4(corners_out), (void*)x_corners);
+ y_out = createDeviceDataArray<float>(dim4(corners_out), (void*)y_corners);
+ resp_out = createDeviceDataArray<float>(dim4(corners_out), (void*)resp_corners);
return corners_out;
}
diff --git a/src/backend/opencl/kernel/susan.cl b/src/backend/opencl/kernel/susan.cl
index b70538e..69131fd 100644
--- a/src/backend/opencl/kernel/susan.cl
+++ b/src/backend/opencl/kernel/susan.cl
@@ -101,21 +101,3 @@ void non_maximal(global float* x_out, global float* y_out,
}
}
#endif
-
-#ifdef KEEP_CORNERS
-kernel
-void keep_corners(global float* x_out, global float* y_out, global float* resp_out,
- global const float* x_in, global const float* y_in,
- global const float* resp_in, global const unsigned* resp_idx,
- const unsigned n_corners)
-{
- const unsigned f = get_global_id(0);
-
- // Keep only the first n_feat features
- if (f < n_corners) {
- x_out[f] = x_in[(unsigned)resp_idx[f]];
- y_out[f] = y_in[(unsigned)resp_idx[f]];
- resp_out[f] = resp_in[f];
- }
-}
-#endif
diff --git a/src/backend/opencl/kernel/susan.hpp b/src/backend/opencl/kernel/susan.hpp
index ba3bcda..9551dbc 100644
--- a/src/backend/opencl/kernel/susan.hpp
+++ b/src/backend/opencl/kernel/susan.hpp
@@ -138,43 +138,6 @@ unsigned nonMaximal(cl::Buffer* x_out, cl::Buffer* y_out, cl::Buffer* resp_out,
return corners_found;
}
-void keepCorners(cl::Buffer* x_out, cl::Buffer* y_out, cl::Buffer* resp_out,
- const cl::Buffer* x_in, const cl::Buffer* y_in,
- const cl::Buffer* resp_in, const cl::Buffer* resp_idx,
- const unsigned n_corners)
-{
- try {
- static std::once_flag compileFlags[DeviceManager::MAX_DEVICES];
- static std::map<int, Program*> kcProg;
- static std::map<int, Kernel*> kcKernel;
-
- int device = getActiveDeviceId();
-
- std::call_once( compileFlags[device], [device] () {
- std::ostringstream options;
- options << " -D KEEP_CORNERS";
- cl::Program prog;
- buildProgram(prog, susan_cl, susan_cl_len, options.str());
- kcProg[device] = new Program(prog);
- kcKernel[device] = new Kernel(*kcProg[device], "keep_corners");
- });
-
- auto kcOp = make_kernel<Buffer, Buffer, Buffer,
- Buffer, Buffer,
- Buffer, Buffer,
- unsigned>(*kcKernel[device]);
-
- NDRange local(THREADS_PER_BLOCK, 1);
- NDRange global(divup(n_corners, local[0]), 1);
-
- kcOp(EnqueueArgs(getQueue(), global, local),
- *x_out, *y_out, *resp_out, *x_in, *y_in, *resp_in, *resp_idx, n_corners);
- } catch (cl::Error err) {
- CL_TO_AF_ERROR(err);
- throw;
- }
-}
-
}
}
diff --git a/src/backend/opencl/susan.cpp b/src/backend/opencl/susan.cpp
index 7d2db5a..71aca1b 100644
--- a/src/backend/opencl/susan.cpp
+++ b/src/backend/opencl/susan.cpp
@@ -13,7 +13,6 @@
#include <kernel/susan.hpp>
#include <cmath>
#include <algorithm>
-#include <sort_index.hpp>
using af::features;
@@ -55,31 +54,9 @@ unsigned susan(Array<float> &x_out, Array<float> &y_out, Array<float> &resp_out,
if (corners_out == 0)
return 0;
- if (corners_found > corners_out) {
- Array<float> susan_responses = createDeviceDataArray<float>(dim4(corners_found),
- (void*)((*resp_corners)()));
- Array<float> susan_sorted = createEmptyArray<float>(dim4(corners_found));
- Array<unsigned> susan_idx = createEmptyArray<unsigned>(dim4(corners_found));
-
- // Sort susan responses
- sort_index<float, false>(susan_sorted, susan_idx, susan_responses, 0);
-
- x_out = createEmptyArray<float>(dim4(corners_out));
- y_out = createEmptyArray<float>(dim4(corners_out));
- resp_out = createEmptyArray<float>(dim4(corners_out));
-
- // Keep only the corners with higher SUSAN responses
- kernel::keepCorners(x_out.get(), y_out.get(), resp_out.get(),
- x_corners, y_corners, susan_sorted.get(), susan_idx.get(),
- corners_out);
-
- bufferFree(x_corners);
- bufferFree(y_corners);
- } else {
- x_out = createDeviceDataArray<float>(dim4(corners_out), (void*)((*x_corners)()));
- y_out = createDeviceDataArray<float>(dim4(corners_out), (void*)((*y_corners)()));
- resp_out = createDeviceDataArray<float>(dim4(corners_out), (void*)((*resp_corners)()));
- }
+ x_out = createDeviceDataArray<float>(dim4(corners_out), (void*)((*x_corners)()));
+ y_out = createDeviceDataArray<float>(dim4(corners_out), (void*)((*y_corners)()));
+ resp_out = createDeviceDataArray<float>(dim4(corners_out), (void*)((*resp_corners)()));
return corners_out;
}
--
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