Skip to content

Changes that should not cause crash, but do. #525

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
danpovey wants to merge 1 commit into
base: master
Choose a base branch
from
Open

Changes that should not cause crash, but do. #525

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
2 changes: 1 addition & 1 deletion k2/csrc/intersect_pruned.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -846,7 +846,7 @@ class MultiGraphDenseIntersectPruned {
// Set the forward log-like of the dest state to the largest of any
// of those of the incoming arcs. Note: we initialized this in
// lambda_init_loglike above.
AtomicMax(&(kept_states_data[state_idx01].forward_loglike),
atomicMax(&(kept_states_data[state_idx01].forward_loglike),
end_loglike_int);
});
}
Expand Down
6 changes: 1 addition & 5 deletions k2/csrc/utils.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -601,14 +601,10 @@ __host__ __device__ __forceinline__ float OrderedIntToFloat(int32_t i) {
host version of Cuda's atomicMax function, marked __host__ (the default) for
clarity. So we can use this in lambdas that run on both host and device.
*/
__host__ __device__ __forceinline__ int32_t AtomicMax(int32_t *address, int32_t val) {
#if defined(__CUDA_ARCH__)
return atomicMax(address, val);
#else
__host__ __forceinline__ int32_t atomicMax(int32_t *address, int32_t val) {
int32_t old = *address;
if (old < val) *address = val;
return old;
#endif
}

// have to figure out if there's a better place to put this
Expand Down