-
Notifications
You must be signed in to change notification settings - Fork 19
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.
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
feat: detray ranges #303
feat: detray ranges #303
Conversation
Still needs a lot of polishing, but at least it seems to work without killing performance. Also needs some CUDA testing |
8fee52a
to
1e64d4d
Compare
Why do you pick a random surface from the detector? |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Nice. I don't know well what is going on exactly in the core directory but the unit tests look solid.
|
||
void test_single(const dindex value, dindex& check) { | ||
dindex* result{nullptr}; | ||
cudaMallocManaged(&result, sizeof(dindex)); |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Maybe you can transfer the single value with vecmem pointer but it's trivial
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Ok, I did not know that existed :D
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Could you maybe 'point' me there?
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Hmm I tried to write a test kernel to check this out, but vecmem unique_ptr header is not compiled in device code in my PC :/... But I have seen codes that call unique_ptr in the device (acts-project/traccc#230) so there might be something wrong in my environment setup. Anyway it's not important so we can go with this
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I am not sure a unique pointer makes much sense for a single index? Or is the unique pointer not meant to do memory management in vecmem?
This was originally intended for the local navigation, because that will return surface indices that are not contiguous in the detector surface container. So I either take the enclosing subrange or pick them one by one ( |
72d78ba
to
74c55d7
Compare
After reimplementing some of the stl stuff, the CUDA tests seem fine now. There is some stupid code in the |
dee83f3
to
7daa0a4
Compare
7daa0a4
to
ae46e65
Compare
Ok, let me merge this as discussed. This should not disturb too much of the rest of the code (I hope) and I start to need some of the type traits... |
@niermann999 Are you sure that this PR passed GitLab CI? It seems it was refused |
This was possibly a bug in thrust with older version of CUDA. |
No, it refuses the pipeline for me, so I check on one of our machines myself. There (CUDA 11.8), everything was fine... |
Thank you :) |
Refactor iterator infrastructure and add iterators that allow chained iteration over different grid bins, as well as to pick random surfaces from the detector store instead of just ranges.
Does not do lazy evaluation though and no pipe operator