[FIX] fix illegal memory accesses for uneven rows where NITEMS > 1 #3469
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
levsnv
changed the title
levsnv
added
bug
|
rerun tests (spdlog compilation errors) |
canonizer
suggested changes
Feb 8, 2021
| for (int j = 0; j < NITEMS; ++j) curr[j] = 0; | ||
| do { | ||
| #pragma unroll | ||
| for (int j = 0; j < NITEMS; ++j) { | ||
| auto n = tree[curr[j]]; | ||
| mask &= ~(n.is_leaf() << j); | ||
| if (!n.is_leaf()) { | ||
| if (mask & (1 << j)) { |
There was a problem hiding this comment.
Suggested change
| if (mask & (1 << j)) { | |
| if ((mask & (1 << j)) != 0) { |
| @@ -172,7 +172,8 @@ struct tree_aggregator_t { | |||
| : tmp_storage(finalize_workspace) {} | |||
|
|
|||
| __device__ __forceinline__ void accumulate( | |||
| vec<NITEMS, float> single_tree_prediction, int tree) { | |||
| vec<NITEMS, float> single_tree_prediction, int tree, int num_rows) { | |||
| for (int i = num_rows; i < NITEMS; ++i) single_tree_prediction[i] = 0.0f; | |||
There was a problem hiding this comment.
I don't think the compiler will be able to unroll this loop.
Besides, I think it would be easier to change the loop in infer_one_tree(), this will also avoid a lot of changes:
vec<NITEMS, output_type> out; // initialized to 0
#pragma unroll
for (int j = 0; j < NITEMS; ++j) {
if (j < n_rows) out[j] = tree[curr[j]].template output<output_type>();
}
| raft::myAtomicAdd(votes + single_tree_prediction[item] * NITEMS + item, | ||
| 1); | ||
| for (int item = 0; item < NITEMS; ++item) { | ||
| if (item < num_rows) { |
There was a problem hiding this comment.
There's always enough shared memory allocated in votes for all NITEMS, independent of how many rows are actually processed. So no need for a conditional here.
JohnZed
added
5 - DO NOT MERGE
|
Should be superseded by #3480 |
|
Closed by #3480 |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
the bug results in memory corruption, which sometimes overwrites useful memory and triggers a test. only compute-sanitizer catches the illegal access right now, so no point adjusting array size.