Skip to content

Fixes sub array (opencl) support for confidenceCC #3668

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
christophe-murphy merged 4 commits into from
Jun 30, 2025
Merged

Fixes sub array (opencl) support for confidenceCC #3668

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
28bd6b0
Adds test helpers for temporary array formats (JIT, SUB, ...)
willyborn Jun 10, 2025
6016fd8
Fixes sub-array (opencl) support for confidenceCC
willyborn Jun 16, 2025
7d9069b
Update flood_fill.cpp
edwinsolisf Jun 27, 2025
577e8d7
Added TODO comment to change this lines after subarrays fix
edwinsolisf 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
11 changes: 9 additions & 2 deletions src/api/c/confidence_connected.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -45,8 +45,15 @@ using std::swap;
template<typename T>
Array<T> pointList(const Array<T>& in, const Array<uint>& x,
const Array<uint>& y) {
af_array xcoords = getHandle<uint>(x);
af_array ycoords = getHandle<uint>(y);

// TODO: Temporary Fix, must fix handling subarrays upstream
// Array<T> has to be a basic array, to be accepted as af_index
Array<uint> x_ = (x.getOffset() == 0 && x.isLinear()) ? x : copyArray(x);
Array<uint> y_ = (y.getOffset() == 0 && y.isLinear()) ? y : copyArray(y);

af_array xcoords = getHandle<uint>(x_);
af_array ycoords = getHandle<uint>(y_);

std::array<af_index_t, AF_MAX_DIMS> idxrs = {{{{xcoords}, false, false},
{{ycoords}, false, false},
createSpanIndex(),
Expand Down
9 changes: 5 additions & 4 deletions src/backend/opencl/kernel/flood_fill.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -23,8 +23,8 @@ kernel void init_seeds(global T *out, KParam oInfo, global const uint *seedsx,
KParam syInfo) {
uint tid = get_global_id(0);
if (tid < sxInfo.dims[0]) {
uint x = seedsx[tid];
uint y = seedsy[tid];
uint x = seedsx[tid + sxInfo.offset];
uint y = seedsy[tid + syInfo.offset];
out[(x * oInfo.strides[0] + y * oInfo.strides[1])] = VALID;
}
}
Expand Down Expand Up @@ -76,14 +76,15 @@ kernel void flood_step(global T *out, KParam oInfo, global const T *img,

T tImgVal =
img[(clamp(gx, 0, (int)(iInfo.dims[0] - 1)) * iInfo.strides[0] +
clamp(gy, 0, (int)(iInfo.dims[1] - 1)) * iInfo.strides[1])];
clamp(gy, 0, (int)(iInfo.dims[1] - 1)) * iInfo.strides[1])+
iInfo.offset];
const int isPxBtwnThresholds =
(tImgVal >= lowValue && tImgVal <= highValue);

int tid = lx + get_local_size(0) * ly;

barrier(CLK_LOCAL_MEM_FENCE);

T origOutVal = lmem[j][i];
bool isBorderPxl = (lx == 0 || ly == 0 || lx == (get_local_size(0) - 1) ||
ly == (get_local_size(1) - 1));
Expand Down
178 changes: 167 additions & 11 deletions test/arrayfire_test.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -105,17 +105,16 @@ std::string readNextNonEmptyLine(std::ifstream &file) {

std::string getBackendName(bool lower) {
af::Backend backend = af::getActiveBackend();
switch(backend) {
case AF_BACKEND_CPU:
return lower ? std::string("cpu") : std::string("CPU");
case AF_BACKEND_CUDA:
return lower ? std::string("cuda") : std::string("CUDA");
case AF_BACKEND_OPENCL:
return lower ? std::string("opencl") : std::string("OpenCL");
case AF_BACKEND_ONEAPI:
return lower ? std::string("oneapi") : std::string("oneAPI");
default:
return lower ? std::string("unknown") : std::string("Unknown");
switch (backend) {
case AF_BACKEND_CPU:
return lower ? std::string("cpu") : std::string("CPU");
case AF_BACKEND_CUDA:
return lower ? std::string("cuda") : std::string("CUDA");
case AF_BACKEND_OPENCL:
return lower ? std::string("opencl") : std::string("OpenCL");
case AF_BACKEND_ONEAPI:
return lower ? std::string("oneapi") : std::string("oneAPI");
default: return lower ? std::string("unknown") : std::string("Unknown");
}
}

Expand Down Expand Up @@ -2046,6 +2045,163 @@ INSTANTIATE(std::complex<float>);
INSTANTIATE(std::complex<double>);
#undef INSTANTIATE

af::array toTempFormat(tempFormat form, const af::array &in) {
af::array ret;
const af::dim4 &dims = in.dims();
switch (form) {
case JIT_FORMAT:
switch (in.type()) {
case b8: ret = not(in); break;
default: ret = in * 2;
}
// Make sure that the base array is <> form original
ret.eval();
switch (in.type()) {
case b8: ret = not(ret); break;
default: ret /= 2;
}
break;
case SUB_FORMAT_dim0: {
af::dim4 pdims(dims);
pdims[0] += 2;
af::array parent = af::randu(pdims, in.type());
parent(af::seq(1, dims[0]), af::span, af::span, af::span) = in;
ret = parent(af::seq(1, dims[0]), af::span, af::span, af::span);
}; break;
case SUB_FORMAT_dim1: {
af::dim4 pdims(dims);
pdims[1] += 2;
af::array parent = af::randu(pdims, in.type());
parent(af::span, af::seq(1, dims[1]), af::span, af::span) = in;
ret = parent(af::span, af::seq(1, dims[1]), af::span, af::span);
}; break;
case SUB_FORMAT_dim2: {
af::dim4 pdims(dims);
pdims[2] += 2;
af::array parent = af::randu(pdims, in.type());
parent(af::span, af::span, af::seq(1, dims[2]), af::span) = in;
ret = parent(af::span, af::span, af::seq(1, dims[2]), af::span);
}; break;
case SUB_FORMAT_dim3: {
af::dim4 pdims(dims);
pdims[3] += 2;
af::array parent = af::randu(pdims, in.type());
parent(af::span, af::span, af::span, af::seq(1, dims[3])) = in;
ret = parent(af::span, af::span, af::span, af::seq(1, dims[3]));
}; break;
case REORDERED_FORMAT: {
const dim_t idxs[4] = {0, 3, 1, 2};
// idxs[0] has to be 0, to keep the same data in mem
dim_t rev_idxs[4];
for (dim_t i = 0; i < 4; ++i) { rev_idxs[idxs[i]] = i; };
ret = af::reorder(in, idxs[0], idxs[1], idxs[2], idxs[3]);
ret = ret.copy(); // make data linear
ret = af::reorder(ret, rev_idxs[0], rev_idxs[1], rev_idxs[2],
rev_idxs[3]);
// ret has same content as in, although data is stored in
// different order
}; break;
case LINEAR_FORMAT:
default: ret = in.copy();
};
return ret;
}

void toTempFormat(tempFormat form, af_array *out, const af_array &in) {
dim_t dims[4];
af_get_dims(dims, dims + 1, dims + 2, dims + 3, in);
unsigned numdims;
af_get_numdims(&numdims, in);
af_dtype ty;
af_get_type(&ty, in);
switch (form) {
case JIT_FORMAT: {
// af_array one = nullptr, min_one = nullptr, res = nullptr;
af_array res = nullptr, two = nullptr;
ASSERT_SUCCESS(af_constant(&two, 2, numdims, dims, ty));
switch (ty) {
case b8: af_not(&res, in); break;
default:
// ret = in + af::constant(1, dims, in.type());
ASSERT_SUCCESS(af_mul(&res, in, two, false));
}
// Make sure that the base array is <> form original
ASSERT_SUCCESS(af_eval(res));
switch (ty) {
case b8: af_not(out, res); break;
default:
ASSERT_SUCCESS(af_div(out, res, two, false)); // NO EVAL!!
}
ASSERT_SUCCESS(af_release_array(two));
two = nullptr;
ASSERT_SUCCESS(af_release_array(res));
res = nullptr;
}; break;
case SUB_FORMAT_dim0: {
const dim_t pdims[4] = {dims[0] + 2, dims[1], dims[2], dims[3]};
af_array parent = nullptr;
ASSERT_SUCCESS(af_randu(&parent, std::max(1u, numdims), pdims, ty));
const af_seq idxs[4] = {af_make_seq(1, dims[0], 1), af_span,
af_span, af_span};

ASSERT_SUCCESS(af_assign_seq(out, parent, numdims, idxs, in));
ASSERT_SUCCESS(af_index(out, parent, numdims, idxs));
ASSERT_SUCCESS(af_release_array(parent));
}; break;
case SUB_FORMAT_dim1: {
const dim_t pdims[4] = {dims[0], dims[1] + 2, dims[2], dims[3]};
af_array parent = nullptr;
ASSERT_SUCCESS(af_randu(&parent, std::max(2u, numdims), pdims, ty));
const af_seq idxs[4] = {af_span, af_make_seq(1, dims[1], 1),
af_span, af_span};
ASSERT_SUCCESS(af_assign_seq(out, parent, numdims, idxs, in));
ASSERT_SUCCESS(af_index(out, parent, numdims, idxs));
ASSERT_SUCCESS(af_release_array(parent));
parent = nullptr;
}; break;
case SUB_FORMAT_dim2: {
const dim_t pdims[4] = {dims[0], dims[1], dims[2] + 2, dims[3]};
af_array parent = nullptr;
ASSERT_SUCCESS(af_randu(&parent, std::max(3u, numdims), pdims, ty));
const af_seq idxs[4] = {af_span, af_span,
af_make_seq(1, dims[2], 1), af_span};
ASSERT_SUCCESS(af_assign_seq(out, parent, numdims, idxs, in));
ASSERT_SUCCESS(af_index(out, parent, numdims, idxs));
ASSERT_SUCCESS(af_release_array(parent));
parent = nullptr;
}; break;
case SUB_FORMAT_dim3: {
const dim_t pdims[4] = {dims[0], dims[1], dims[2], dims[3] + 2};
af_array parent = nullptr;
ASSERT_SUCCESS(af_randu(&parent, std::max(4u, numdims), pdims, ty));
const af_seq idxs[4] = {af_span, af_span, af_span,
af_make_seq(1, dims[3], 1)};
ASSERT_SUCCESS(af_assign_seq(out, parent, numdims, idxs, in));
ASSERT_SUCCESS(af_index(out, parent, numdims, idxs));
ASSERT_SUCCESS(af_release_array(parent));
parent = nullptr;
}; break;
case REORDERED_FORMAT: {
const unsigned idxs[4] = {0, 3, 1, 2};
// idxs[0] has to be 0, to keep the same data in mem
dim_t rev_idxs[4];
for (dim_t i = 0; i < 4; ++i) { rev_idxs[idxs[i]] = i; };
af_array rev = nullptr;
ASSERT_SUCCESS(
af_reorder(&rev, in, idxs[0], idxs[1], idxs[2], idxs[3]));
ASSERT_SUCCESS(af_copy_array(out, rev));
ASSERT_SUCCESS(af_reorder(out, rev, rev_idxs[0], rev_idxs[1],
rev_idxs[2], rev_idxs[3]));
// ret has same content as in, although data is stored in
// different order
ASSERT_SUCCESS(af_release_array(rev));
rev = nullptr;
}; break;
case LINEAR_FORMAT:
default: af_copy_array(out, in);
};
}

int main(int argc, char **argv) {
::testing::InitGoogleTest(&argc, argv);
return RUN_ALL_TESTS();
Expand Down
59 changes: 53 additions & 6 deletions test/confidence_connected.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -92,9 +92,9 @@ void testImage(const std::string pTestFile, const size_t numSeeds,
params.iterations = iter;
params.replace = 255.0;

ASSERT_SUCCESS(af_confidence_cc(&outArray, inArray, seedxArr, seedyArr, params.radius,
params.multiplier, params.iterations,
params.replace));
ASSERT_SUCCESS(af_confidence_cc(&outArray, inArray, seedxArr, seedyArr,
params.radius, params.multiplier,
params.iterations, params.replace));
int device = 0;
ASSERT_SUCCESS(af_get_device(&device));
ASSERT_SUCCESS(af_sync(device));
Expand Down Expand Up @@ -141,9 +141,9 @@ void testData(CCCTestParams params) {
(af_dtype)af::dtype_traits<T>::af_type));

af_array outArray = 0;
ASSERT_SUCCESS(af_confidence_cc(&outArray, inArray, seedxArr, seedyArr, params.radius,
params.multiplier, params.iterations,
params.replace));
ASSERT_SUCCESS(af_confidence_cc(&outArray, inArray, seedxArr, seedyArr,
params.radius, params.multiplier,
params.iterations, params.replace));
int device = 0;
ASSERT_SUCCESS(af_get_device(&device));
ASSERT_SUCCESS(af_sync(device));
Expand Down Expand Up @@ -201,3 +201,50 @@ INSTANTIATE_TEST_SUITE_P(
<< info.param.iterations << "_replace_" << info.param.replace;
return ss.str();
});

#define TEST_FORMATS(form) \
TEST(TEMP_FORMAT, form##_2Dseed) { \
UNSUPPORTED_BACKEND(AF_BACKEND_ONEAPI); \
const string filename(string(TEST_DIR) + "/confidence_cc/donut.png"); \
const af::array image(af::loadImage(filename.c_str())); \
const af::array seed(dim4(1, 2), {10u, 8u}); \
\
const af::array out = \
af::confidenceCC(toTempFormat(form, image), \
toTempFormat(form, seed), 3, 3, 25, 255.0); \
const af::array gold = af::confidenceCC(image, seed, 3, 3, 25, 255.0); \
\
EXPECT_ARRAYS_EQ(out, gold); \
} \
\
TEST(TEMP_FORMAT, form##_2xSeed) { \
UNSUPPORTED_BACKEND(AF_BACKEND_ONEAPI); \
const string filename(string(TEST_DIR) + "/confidence_cc/donut.png"); \
const af::array image(af::loadImage(filename.c_str())); \
const af::array seedx({10u}); \
const af::array seedy({8u}); \
\
const af::array out = af::confidenceCC( \
toTempFormat(form, image), toTempFormat(form, seedx), \
toTempFormat(form, seedy), 3, 3, 25, 255.0); \
const af::array gold = \
af::confidenceCC(image, seedx, seedy, 3, 3, 25, 255.0); \
\
EXPECT_ARRAYS_EQ(out, gold); \
} \
TEST(TEMP_FORMAT, form##_vectSeed) { \
UNSUPPORTED_BACKEND(AF_BACKEND_ONEAPI); \
const string filename(string(TEST_DIR) + "/confidence_cc/donut.png"); \
const af::array image(af::loadImage(filename.c_str())); \
const unsigned seedx[1] = {10u}; \
const unsigned seedy[1] = {8u}; \
\
const af::array out = af::confidenceCC(toTempFormat(form, image), 1, \
seedx, seedy, 3, 3, 25, 255.0); \
const af::array gold = \
af::confidenceCC(image, 1, seedx, seedy, 3, 3, 25, 255.0); \
\
EXPECT_ARRAYS_EQ(out, gold); \
}

FOREACH_TEMP_FORMAT(TEST_FORMATS)
32 changes: 28 additions & 4 deletions test/testHelpers.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -244,10 +244,10 @@ bool noHalfTests(af::dtype ty);
GTEST_SKIP() << "Device doesn't support Half"

#ifdef SKIP_UNSUPPORTED_TESTS
#define UNSUPPORTED_BACKEND(backend) \
if(backend == af::getActiveBackend()) \
GTEST_SKIP() << "Skipping unsupported function on " \
+ getBackendName() + " backend"
#define UNSUPPORTED_BACKEND(backend) \
if (backend == af::getActiveBackend()) \
GTEST_SKIP() << "Skipping unsupported function on " + getBackendName() + \
" backend"
#else
#define UNSUPPORTED_BACKEND(backend)
#endif
Expand Down Expand Up @@ -653,6 +653,30 @@ ::testing::AssertionResult assertArrayEq(std::string aName, std::string bName,
const af_array a, const af_array b,
TestOutputArrayInfo *metadata);

enum tempFormat {
LINEAR_FORMAT, // Linear array (= default)
JIT_FORMAT, // Array which has JIT operations outstanding
SUB_FORMAT_dim0, // Array where only a subset is allocated for dim0
SUB_FORMAT_dim1, // Array where only a subset is allocated for dim1
SUB_FORMAT_dim2, // Array where only a subset is allocated for dim2
SUB_FORMAT_dim3, // Array where only a subset is allocated for dim3
REORDERED_FORMAT // Array where the dimensions are reordered
};
// Calls the function fn for all available formats
#define FOREACH_TEMP_FORMAT(TESTS) \
TESTS(LINEAR_FORMAT) \
TESTS(JIT_FORMAT) \
TESTS(SUB_FORMAT_dim0) \
TESTS(SUB_FORMAT_dim1) \
TESTS(SUB_FORMAT_dim2) \
TESTS(SUB_FORMAT_dim3) \
TESTS(REORDERED_FORMAT)

// formats the "in" array according to provided format. The content remains
// unchanged.
af::array toTempFormat(tempFormat form, const af::array &in);
void toTempFormat(tempFormat form, af_array *out, const af_array &in);

#ifdef __GNUC__
#pragma GCC diagnostic pop
#endif