Improve DeviceScan env test coverage and clean up stale code#8280
Open
gonidelis wants to merge 1 commit intoNVIDIA:mainfrom
Open
Improve DeviceScan env test coverage and clean up stale code#8280gonidelis wants to merge 1 commit intoNVIDIA:mainfrom
gonidelis wants to merge 1 commit intoNVIDIA:mainfrom
Conversation
* Use non-identity init values so the init actually affects scan results * Use multi-element inputs so the scan op is exercised, not just init * Remove unused cudaGetDevice/PtxVersion calls, use ptx_arch_id instead
gonidelis
commented
Apr 2, 2026
Contributor
🥳 CI Workflow Results🟩 Finished in 42m 52s: Pass: 100%/213 | Total: 1d 20h | Max: 30m 46s | Hits: 98%/124738See results here. |
gonidelis
commented
Apr 3, 2026
| for (int i = 0; i < num_items; i++) | ||
| { | ||
| REQUIRE(d_out[i] == i); | ||
| REQUIRE(d_out[i] == i + init); |
Member
Author
There was a problem hiding this comment.
use a thrust::equals(d_out, counting_iterator(init))
Member
Author
There was a problem hiding this comment.
don't teach llms how to write bad cuda code
gonidelis
commented
Apr 3, 2026
| num_items_t num_items = 1; | ||
| num_items_t num_items = 2; | ||
| c2h::device_vector<int> d_block_size(1); | ||
| block_size_check_t block_size_check{thrust::raw_pointer_cast(d_block_size.data())}; |
Member
Author
There was a problem hiding this comment.
modify the functor with if threadId.x == 0 so there is no data race
gonidelis
commented
Apr 3, 2026
| { | ||
| REQUIRE(d_out[i] == i); | ||
| } | ||
|
|
Member
Author
There was a problem hiding this comment.
we could use the stream_registry_factory_t for that but that might end up getting too involved
let's just create a custom IputTypeT for d_in that stores blockDim.x info internally and we can poll it from there
gonidelis
commented
Apr 3, 2026
| for (int i = 0; i < num_items; i++) | ||
| { | ||
| REQUIRE(d_out[i] == (i + 1 + 10.0f)); | ||
| REQUIRE(d_out[i] == (i + 1 + init)); |
Member
Author
There was a problem hiding this comment.
use thrust::equal as above
gonidelis
commented
Apr 3, 2026
| for (int i = 0; i < num_items; i++) | ||
| { | ||
| REQUIRE(d_out[i] == i); | ||
| REQUIRE(d_out[i] == i + init); |
Member
Author
There was a problem hiding this comment.
same, use thrust::equal
This file contains hidden or 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
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
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.
Taking the opportunity from a recent review from one of my env PRs I took the liberty to clean up the scan tests mainly for two reasons:
initvalues to make sure they are factored in properlyexclusive*it's just theinitvalue) to make sure without loss of generality that the results is correct for the full input range and not just the edge case