8000 Fixes sub array (opencl) support for confidenceCC (#3668) · arrayfire/arrayfire@f01e6fe · GitHub
[go: up one dir, main page]
Skip to content

Commit f01e6fe

Browse files
willybornedwinsolisf
willyborn
and
edwinsolisf
authored
Fixes sub array (opencl) support for confidenceCC (#3668)
* Adds test helpers for temporary array formats (JIT, SUB, ...) * Fixes sub-array (opencl) support for confidenceCC * Update flood_fill.cpp Removed unnecessary #include * Added TODO comment to change this lines after subarrays fix --------- Co-authored-by: Edwin Lester Solís Fuentes <68087165+edwinsolisf@users.noreply.github.com>
1 parent eaa49ca commit f01e6fe

File tree

3 files changed

+67
-12
lines changed

3 files changed

+67
-12
lines changed

‎src/api/c/confidence_connected.cpp

Lines changed: 9 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -45,8 +45,15 @@ using std::swap;
4545
template<typename T>
4646
Array<T> pointList(const Array<T>& in, const Array<uint>& x,
4747
const Array<uint>& y) {
48-
af_array xcoords = getHandle<uint>(x);
49-
af_array ycoords = getHandle<uint>(y);
48+
49+
// TODO: Temporary Fix, must fix handling subarrays upstream
50+
// Array<T> has to be a basic array, to be accepted as af_index
51+
Array<uint> x_ = (x.getOffset() == 0 && x.isLinear()) ? x : copyArray(x);
52+
Array<uint> y_ = (y.getOffset() == 0 && y.isLinear()) ? y : copyArray(y);
53+
54+
af_array xcoords = getHandle<uint>(x_);
55+
af_array ycoords = getHandle<uint>(y_);
56+
5057
std::array<af_index_t, AF_MAX_DIMS> idxrs = {{{{xcoords}, false, false},
5158
{{ycoords}, false, false},
5259
createSpanIndex(),

‎src/backend/opencl/kernel/flood_fill.cl

Lines changed: 5 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -23,8 +23,8 @@ kernel void init_seeds(global T *out, KParam oInfo, global const uint *seedsx,
2323
KParam syInfo) {
2424
uint tid = get_global_id(0);
2525
if (tid < sxInfo.dims[0]) {
26-
uint x = seedsx[tid];
27-
uint y = seedsy[tid];
26+
uint x = seedsx[tid + sxInfo.offset];
27+
uint y = seedsy[tid + syInfo.offset];
2828
out[(x * oInfo.strides[0] + y * oInfo.strides[1])] = VALID;
2929
}
3030
}
@@ -76,14 +76,15 @@ kernel void flood_step(global T *out, KParam oInfo, global const T *img,
7676

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

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

8586
barrier(CLK_LOCAL_MEM_FENCE);
86-
87+
8788
T origOutVal = lmem[j][i];
8889
bool isBorderPxl = (lx == 0 || ly == 0 || lx == (get_local_size(0) - 1) ||
8990
ly == (get_local_size(1) - 1));

‎test/confidence_connected.cpp

Lines changed: 53 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -92,9 +92,9 @@ void testImage(const std::string pTestFile, const size_t numSeeds,
9292
params.iterations = iter;
9393
params.replace = 255.0;
9494

95-
ASSERT_SUCCESS(af_confidence_cc(&outArray, inArray, seedxArr, seedyArr, params.radius,
96-
params.multiplier, params.iterations,
97-
params.replace));
95+
ASSERT_SUCCESS(af_confidence_cc(&outArray, inArray, seedxArr, seedyArr,
96+
params.radius, params.multiplier,
97+
params.iterations, params.replace));
9898
int device = 0;
9999
ASSERT_SUCCESS(af_get_device(&device));
100100
ASSERT_SUCCESS(af_sync(device));
@@ -141,9 +141,9 @@ void testData(CCCTestParams params) {
141141
(af_dtype)af::dtype_traits<T>::af_type));
142142

143143
af_array outArray = 0;
144-
ASSERT_SUCCESS(af_confidence_cc(&outArray, inArray, seedxArr, seedyArr, params.radius,
145-
params.multiplier, params.iterations,
146-
params.replace));
144+
ASSERT_SUCCESS(af_confidence_cc(&outArray, inArray, seedxArr, seedyArr,
145+
params.radius, params.multiplier,
146+
params.iterations, params.replace));
147147
int device = 0;
148148
ASSERT_SUCCESS(af_get_device(&device));
149149
ASSERT_SUCCESS(af_sync(device));
@@ -201,3 +201,50 @@ INSTANTIATE_TEST_SUITE_P(
201201
<< info.param.iterations << "_replace_" << info.param.replace;
202202
return ss.str();
203203
});
204+
205+
#define TEST_FORMATS(form) \
206+
TEST(TEMP_FORMAT, form##_2Dseed) { \
207+
UNSUPPORTED_BACKEND(AF_BACKEND_ONEAPI); \
208+
const string filename(string(TEST_DIR) + "/confidence_cc/donut.png"); \
209+
const af::array image(af::loadImage(filename.c_str())); \
210+
const af::array seed(dim4(1, 2), {10u, 8u}); \
211+
\
212+
const af::array out = \
213+
af::confidenceCC(toTempFormat(form, image), \
214+
toTempFormat(form, seed), 3, 3, 25, 255.0); \
215+
const af::array gold = af::confidenceCC(image, seed, 3, 3, 25, 255.0); \
216+
\
217+
EXPECT_ARRAYS_EQ(out, gold); \
218+
} \
219+
\
220+
TEST(TEMP_FORMAT, form##_2xSeed) { \
221+
UNSUPPORTED_BACKEND(AF_BACKEND_ONEAPI); \
222+
const string filename(string(TEST_DIR) + "/confidence_cc/donut.png"); \
223+
const af::array image(af::loadImage(filename.c_str())); \
224+
const af::array seedx({10u}); \
225+
const af::array seedy({8u}); \
226+
\
227+
const af::array out = af::confidenceCC( \
228+
A93C toTempFormat(form, image), toTempFormat(form, seedx), \
229+
toTempFormat(form, seedy), 3, 3, 25, 255.0); \
230+
const af::array gold = \
231+
af::confidenceCC(image, seedx, seedy, 3, 3, 25, 255.0); \
232+
\
233+
EXPECT_ARRAYS_EQ(out, gold); \
234+
} \
235+
TEST(TEMP_FORMAT, form##_vectSeed) { \
236+
UNSUPPORTED_BACKEND(AF_BACKEND_ONEAPI); \
237+
const string filename(string(TEST_DIR) + "/confidence_cc/donut.png"); \
238+
const af::array image(af::loadImage(filename.c_str())); \
239+
const unsigned seedx[1] = {10u}; \
240+
const unsigned seedy[1] = {8u}; \
241+
\
242+
const af::array out = af::confidenceCC(toTempFormat(form, image), 1, \
243+
seedx, seedy, 3, 3, 25, 255.0); \
244+
const af::array gold = \
245+
af::confidenceCC(image, 1, seedx, seedy, 3, 3, 25, 255.0); \
246+
\
247+
EXPECT_ARRAYS_EQ(out, gold); \
248+
}
249+
250+
FOREACH_TEMP_FORMAT(TEST_FORMATS)

0 commit comments

Comments
 (0)
0