[hipCUB][Code Coverage] Increase Code Coverage for hipCUB#172
Conversation
NguyenNhuDi
requested review from
RobsonRLemos,
lawruble13,
stanleytsang-amd and
umfranzw
as code owners
umfranzw
left a comment
umfranzw
left a comment
There was a problem hiding this comment.
Nice work - it's great to have more test coverage. I just have a few questions and comments on a few things.
| { | ||
| const size_t offset = bi * items_per_block; | ||
| const size_t i0 = offset + ti * items_per_thread + ii; | ||
| const size_t i1 = offset + host_ranks[i0] % block_size * items_per_thread |
There was a problem hiding this comment.
I have a feeling I'm not understanding this correctly, but I just wanted to double-check. On line 907, host_ranks[i0] will be -1 in some cases (set by the loop on line 893). I believe that in this case we don't want to set host_expected (since it was already set to 0 in the loop above) - is that right?. If so, don't we need to check (on line 909) if host_ranks[i0] is -1, and if so, avoid setting host_expected?
| 0, | ||
| 0, | ||
| device_input); | ||
| HIP_CHECK(hipPeekAtLastError()); |
There was a problem hiding this comment.
Just a note that as of HIP 7.0, the error returned by this function might not be from the kernel launch above - it could be from any HIP API call executed before this (eg. a kernel launch in a previous test). If you need to be certain that the error is from the kernel call above, you can add a call to hipGetLastError() before the call to hipLaunchKernelGGL. This will clear the internally recorded HIP error. It might also be good to use hipGetLastError on line 1155 (in place of hipPeekAtLastError), so that (because it clears the internally recorded error) any error that occurs won't affect following tests.
There are a number of other hipPeekAtLastError calls in tests in this PR that might also benefit from this kind of change.
…into zenguyen-hipcub/increase-code-coverage-rs-7.0
| HIP_CHECK( | ||
| hipMemcpy(device_input, input.data(), input.size() * sizeof(type), hipMemcpyHostToDevice)); | ||
|
|
||
| HIP_CHECK(hipPeekAtLastError()); |
There was a problem hiding this comment.
Could we change lines 170 and 182 from hipPeekAtLastError to hipGetLastError? That way the internally recorded error will be cleared before and after the kernel call.
|
Clsoing this PR to make a new one to rebase to develop. New PR here: #411 |
Re making PR #172 to merge to develop instead of release-staging
Re making PR ROCm#172 to merge to develop instead of release-staging
Re making PR #172 to merge to develop instead of release-staging
Re making PR #172 to merge to develop instead of release-staging
Co-authored-by: BrianHarrisonAMD <169072757+BrianHarrisonAMD@users.noreply.github.com>
Co-authored-by: BrianHarrisonAMD <169072757+BrianHarrisonAMD@users.noreply.github.com> [ROCm/hipDNN commit: 5c67ec1]
4 participants
This PR is opened for the purpose of merging into the release staging branch.
This PR should not be merged until PR #171 is merged. This is because the additional unit tests found that a function in hipCUB was making a wrong backend call and this will cause the unit tests to fail during compilation.