Skip to content

Fixes sub array (opencl, cuda, cpu, oneapi) for orb #3670

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

Open
willyborn wants to merge 5 commits into
base: master
Choose a base branch
from
Open

Fixes sub array (opencl, cuda, cpu, oneapi) for orb #3670

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
28bd6b0
Adds test helpers for temporary array formats (JIT, SUB, ...)
willyborn Jun 10, 2025
a70ab50
Fixes sub-array (cuda, opencl) support for harris
willyborn Jun 27, 2025
072085b
Fixes sub-array (cpu, cuda, opencl) support for fast
willyborn Jun 27, 2025
a3a91e1
Fixes sub-array (cpu, cuda, opencl, oneapi) support for convolve
willyborn Jun 27, 2025
73b9dc9
Fixes sub-array (cpu, cuda, opencl, oneapi) support for orb
willyborn Jun 27, 2025
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
Prev Previous commit
Next Next commit
Fixes sub-array (cuda, opencl) support for harris
  • Loading branch information
@willyborn
willyborn committed Aug 5, 2025
commit a70ab50e51d7d18eb9da1e39a3268a58cf351134
32 changes: 19 additions & 13 deletions src/backend/cuda/kernel/harris.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -174,7 +174,7 @@ void harris(unsigned* corners_out, float** x_out, float** y_out,
filter.strides[k] = filter.dims[k - 1] * filter.strides[k - 1];
}

int filter_elem = filter.strides[3] * filter.dims[3];
int filter_elem = filter.dims[0] * filter.dims[1];
auto filter_alloc = memAlloc<convAccT>(filter_elem);
filter.ptr = filter_alloc.get();
CUDA_CHECK(cudaMemcpyAsync(filter.ptr, h_filter.data(),
Expand All @@ -184,9 +184,11 @@ void harris(unsigned* corners_out, float** x_out, float** y_out,
const unsigned border_len = filter_len / 2 + 1;

Param<T> ix, iy;
for (dim_t i = 0; i < 4; i++) {
ix.dims[0] = iy.dims[0] = in.dims[0];
ix.strides[0] = iy.strides[0] = 1;
for (dim_t i = 1; i < 4; i++) {
ix.dims[i] = iy.dims[i] = in.dims[i];
ix.strides[i] = iy.strides[i] = in.strides[i];
ix.strides[i] = iy.strides[i] = ix.strides[i - 1] * ix.dims[i - 1];
}
auto ix_alloc = memAlloc<T>(ix.dims[3] * ix.strides[3]);
auto iy_alloc = memAlloc<T>(iy.dims[3] * iy.strides[3]);
Expand All @@ -198,12 +200,16 @@ void harris(unsigned* corners_out, float** x_out, float** y_out,

Param<T> ixx, ixy, iyy;
Param<T> ixx_tmp, ixy_tmp, iyy_tmp;
for (dim_t i = 0; i < 4; i++) {
ixx.dims[i] = ixy.dims[i] = iyy.dims[i] = in.dims[i];
ixx_tmp.dims[i] = ixy_tmp.dims[i] = iyy_tmp.dims[i] = in.dims[i];
ixx.strides[i] = ixy.strides[i] = iyy.strides[i] = in.strides[i];
ixx_tmp.strides[i] = ixy_tmp.strides[i] = iyy_tmp.strides[i] =
in.strides[i];
ixx.dims[0] = ixy.dims[0] = iyy.dims[0] = ixx_tmp.dims[0] =
ixy_tmp.dims[0] = iyy_tmp.dims[0] = in.dims[0];
ixx.strides[0] = ixy.strides[0] = iyy.strides[0] = ixx_tmp.strides[0] =
ixy_tmp.strides[0] = iyy_tmp.strides[0] = 1;
for (dim_t i = 1; i < 4; i++) {
ixx.dims[i] = ixy.dims[i] = iyy.dims[i] = ixx_tmp.dims[i] =
ixy_tmp.dims[i] = iyy_tmp.dims[i] = in.dims[i];
ixx.strides[i] = ixy.strides[i] = iyy.strides[i] = ixx_tmp.strides[i] =
ixy_tmp.strides[i] = iyy_tmp.strides[i] =
ixx.strides[i - 1] * ixx.dims[i - 1];
}
auto ixx_alloc = memAlloc<T>(ixx.dims[3] * ixx.strides[3]);
auto ixy_alloc = memAlloc<T>(ixy.dims[3] * ixy.strides[3]);
Expand All @@ -214,9 +220,9 @@ void harris(unsigned* corners_out, float** x_out, float** y_out,

// Compute second-order derivatives
dim3 threads(THREADS_PER_BLOCK, 1);
dim3 blocks(divup(in.dims[3] * in.strides[3], threads.x), 1);
dim3 blocks(divup(in.dims[0] * in.dims[1], threads.x), 1);
CUDA_LAUNCH((second_order_deriv<T>), blocks, threads, ixx.ptr, ixy.ptr,
iyy.ptr, in.dims[3] * in.strides[3], ix.ptr, iy.ptr);
iyy.ptr, in.dims[0] * in.dims[1], ix.ptr, iy.ptr);

auto ixx_tmp_alloc = memAlloc<T>(ixx_tmp.dims[3] * ixx_tmp.strides[3]);
auto ixy_tmp_alloc = memAlloc<T>(ixy_tmp.dims[3] * ixy_tmp.strides[3]);
Expand All @@ -235,7 +241,7 @@ void harris(unsigned* corners_out, float** x_out, float** y_out,

// Number of corners is not known a priori, limit maximum number of corners
// according to image dimensions
unsigned corner_lim = in.dims[3] * in.strides[3] * 0.2f;
unsigned corner_lim = in.dims[0] * in.dims[1] * 0.2f;

auto d_corners_found = memAlloc<unsigned>(1);
CUDA_CHECK(cudaMemsetAsync(d_corners_found.get(), 0, sizeof(unsigned),
Expand All @@ -245,7 +251,7 @@ void harris(unsigned* corners_out, float** x_out, float** y_out,
auto d_y_corners = memAlloc<float>(corner_lim);
auto d_resp_corners = memAlloc<float>(corner_lim);

auto d_responses = memAlloc<T>(in.dims[3] * in.strides[3]);
auto d_responses = memAlloc<T>(in.dims[0] * in.dims[1]);

// Calculate Harris responses for all pixels
threads = dim3(BLOCK_SIZE, BLOCK_SIZE);
Expand Down
12 changes: 6 additions & 6 deletions src/backend/opencl/kernel/harris.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -126,21 +126,20 @@ void harris(unsigned *corners_out, Param &x_out, Param &y_out, Param &resp_out,

// Second order-derivatives kernel sizes
const unsigned blk_x_so =
divup(in.info.dims[3] * in.info.strides[3], HARRIS_THREADS_PER_GROUP);
divup(in.info.dims[0] * in.info.dims[1], HARRIS_THREADS_PER_GROUP);
const NDRange local_so(HARRIS_THREADS_PER_GROUP, 1);
const NDRange global_so(blk_x_so * HARRIS_THREADS_PER_GROUP, 1);

// Compute second-order derivatives
soOp(EnqueueArgs(getQueue(), global_so, local_so), *ixx.get(), *ixy.get(),
*iyy.get(), in.info.dims[3] * in.info.strides[3], *ix.get(),
*iy.get());
*iyy.get(), in.info.dims[0] * in.info.dims[1], *ix.get(), *iy.get());
CL_DEBUG_FINISH(getQueue());

// Convolve second order derivatives with proper window filter
conv_helper<T, convAccT>(ixx, ixy, iyy, filter);

cl::Buffer *d_responses =
bufferAlloc(in.info.dims[3] * in.info.strides[3] * sizeof(T));
bufferAlloc(in.info.dims[0] * in.info.dims[1] * sizeof(T));

// Harris responses kernel sizes
unsigned blk_x_hr =
Expand All @@ -159,7 +158,7 @@ void harris(unsigned *corners_out, Param &x_out, Param &y_out, Param &resp_out,

// Number of corners is not known a priori, limit maximum number of corners
// according to image dimensions
unsigned corner_lim = in.info.dims[3] * in.info.strides[3] * 0.2f;
unsigned corner_lim = in.info.dims[0] * in.info.dims[1] * 0.2f;

unsigned corners_found = 0;
cl::Buffer *d_corners_found = bufferAlloc(sizeof(unsigned));
Expand Down Expand Up @@ -221,7 +220,8 @@ void harris(unsigned *corners_out, Param &x_out, Param &y_out, Param &resp_out,
harris_idx.info.dims[k - 1] * harris_idx.info.strides[k - 1];
}

int sort_elem = harris_resp.info.strides[3] * harris_resp.info.dims[3];
int sort_elem = harris_resp.info.dims[0] * harris_resp.info.dims[1];

harris_resp.data = d_resp_corners;
// Create indices using range
harris_idx.data = bufferAlloc(sort_elem * sizeof(unsigned));
Expand Down
40 changes: 40 additions & 0 deletions test/harris.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -220,3 +220,43 @@ TEST(FloatHarris, CPP) {
<< "at: " << elIter << endl;
}
}

#define TESTS_TEMP_FORMAT(form) \
TEST(TEMP_FORMAT, form) { \
UNSUPPORTED_BACKEND(AF_BACKEND_ONEAPI); \
IMAGEIO_ENABLED_CHECK(); \
\
constexpr int MAX_CORNERS = 500; \
\
vector<dim4> inDims; \
vector<string> inFiles; \
vector<vector<float>> gold; \
\
readImageTests(string(TEST_DIR "/harris/square_0_3.test"), inDims, \
inFiles, gold); \
inFiles[0].insert(0, string(TEST_DIR "/harris/")); \
array in = loadImage(inFiles[0].c_str(), false); \
\
features out = \
harris(toTempFormat(form, in), MAX_CORNERS, 1e5f, 0.0f, 3, 0.04f); \
features gout = harris(in, MAX_CORNERS, 1e5f, 0.0f, 3, 0.04f); \
\
ASSERT_GT(MAX_CORNERS, out.getNumFeatures()); \
\
array score = out.getX() * in.dims().dims[1] + out.getY(); \
array idx, score_sorted; \
sort(score_sorted, idx, score); \
\
array gscore = gout.getX() * in.dims().dims[1] + gout.getY(); \
array gidx, gscore_sorted; \
sort(gscore_sorted, gidx, gscore); \
\
EXPECT_ARRAYS_EQ(out.getX()(idx), gout.getX()(gidx)); \
EXPECT_ARRAYS_EQ(out.getY()(idx), gout.getY()(gidx)); \
EXPECT_ARRAYS_EQ(out.getOrientation()(idx), \
gout.getOrientation()(gidx)); \
EXPECT_ARRAYS_EQ(out.getScore()(idx), gout.getScore()(gidx)); \
EXPECT_ARRAYS_EQ(out.getSize()(idx), gout.getSize()(gidx)); \
};

FOREACH_TEMP_FORMAT(TESTS_TEMP_FORMAT)