Skip to content

Minor Bug Fixes #1605

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
shehzan10 merged 4 commits into from
Sep 29, 2016
Merged

Minor Bug Fixes #1605

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
1b18226
Fix double free corruption when release arrays in af_assign_seq
shehzan10 Sep 23, 2016
0ed6ccc
Fix syncthreads in cuda nearest neighbour
shehzan10 Sep 21, 2016
ab61f45
Remove rogue printf from bilateral kernel
shehzan10 Sep 24, 2016
38cba47
Add key.eval to CPU scan by key
shehzan10 Sep 26, 2016
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
10 changes: 8 additions & 2 deletions src/api/c/assign.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -127,7 +127,10 @@ af_err af_assign_seq(af_array *out,
AF_CHECK(af_assign_seq(&tmp_out, tmp_in, ndims, index, rhs));
AF_CHECK(af_moddims(out, tmp_out, lInfo.ndims(), lInfo.dims().get()));
AF_CHECK(af_release_array(tmp_in));
AF_CHECK(af_release_array(tmp_out));
// This can run into a double free issue if tmp_in == tmp_out
// The condition ensures release only if both are different
// Issue found on Tegra X1
if(tmp_in != tmp_out) AF_CHECK(af_release_array(tmp_out));
return AF_SUCCESS;
}

Expand Down Expand Up @@ -244,7 +247,10 @@ af_err af_assign_gen(af_array *out,
AF_CHECK(af_assign_gen(&tmp_out, tmp_in, ndims, indexs, rhs_));
AF_CHECK(af_moddims(out, tmp_out, lInfo.ndims(), lInfo.dims().get()));
AF_CHECK(af_release_array(tmp_in));
AF_CHECK(af_release_array(tmp_out));
// This can run into a double free issue if tmp_in == tmp_out
// The condition ensures release only if both are different
// Issue found on Tegra X1
if(tmp_in != tmp_out) AF_CHECK(af_release_array(tmp_out));
return AF_SUCCESS;
}

Expand Down
2 changes: 2 additions & 0 deletions src/backend/cpu/scan_by_key.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -31,6 +31,8 @@ namespace cpu
kernel::scan_dim_by_key<op, Ti, Tk, To, 4> func4(inclusive_scan);

in.eval();
key.eval();

switch (in.ndims()) {
case 1:
getQueue().enqueue(func1, out, 0, key, 0, in, 0, dim);
Expand Down
2 changes: 1 addition & 1 deletion src/backend/cuda/kernel/nearest_neighbour.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -411,8 +411,8 @@ __global__ void select_matches(
s_dist[sid] = dist;
s_idx[sid] = s_idx[sid + i];
}
__syncthreads();
}
Copy link
Copy Markdown
Member

@pavanky pavanky Sep 23, 2016

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

__syncthreads() should go here. This is a block reduction step.

Copy link
Copy Markdown
Member Author

@shehzan10 shehzan10 Sep 23, 2016

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Ah! my bad. i thought i moved it there.

Copy link
Copy Markdown
Member Author

@shehzan10 shehzan10 Sep 23, 2016

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Corrected.

__syncthreads();
}

// Store best matches and indexes to training dataset
Expand Down
1 change: 0 additions & 1 deletion src/backend/opencl/kernel/bilateral.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -50,7 +50,6 @@ void bilateral(Param out, const Param in, float s_sigma, float c_sigma)

std::call_once( compileFlags[device], [device] () {
bool use_native_exp = getActivePlatform() != AFCL_PLATFORM_POCL;
printf("NATIVE_EXP: %d\n", use_native_exp);
std::ostringstream options;
options << " -D inType=" << dtype_traits<inType>::getName()
<< " -D outType=" << dtype_traits<outType>::getName();
Expand Down