Repository navigation
[rocPRIM] Mergeback 6.4 changes - #694
Closed
NguyenNhuDi wants to merge 14 commits into
Closed
NguyenNhuDi wants to merge 14 commits into
NguyenNhuDi wants to merge 14 commits into
Conversation
The rocprim::histogram_even algorithm uses rocprim::atomic_add. There are atomic_add overloads for all of the non-half-type primitive data types we support, with the exception of double. Because of this, if you called rocprim::histogram_even and used double as the counter type, you'd get a compilation error. This changes just adds an atomic_add overload that accepts double, and tacks on a few unit test cases that exercise histogram_even with float and double counters.
* Restart on HIP language support * linking changes * rearrange --------- Co-authored-by: Lauren Wrubleski <Lauren.Wrubleski@amd.com>
* Fix for using iterators in batch_copy * Add indirect_iterator test to batch_memcpy * Add specific iterator test based on reproducer, with iterators in iterators.
Co-authored-by: Robin Voetter <robin@streamhpc.com>
This change implements a workaround to fix block_histogram while we wait for a compiler fix. The block_histogram_sort algorithm contains a loop that initializes values in a shared memory struct. Then it runs the flag_heads() algorithm using that struct. It appear that in some cases, the optimizer is reordering things such that shared memory writes in flag_heads() run before the loop initialization code completes. Changing the loop condition and removing the inner "if" statement seems to prevent the optimizer from doing this. From an algorithmic standpoint, this should be functionally equivalent to the way the code was before the change.
…3.18 (#667) Recently, we added CMake HIP Language support. That change introduced generator expressions that use the LINK_LANGUAGE keyword. This keyword is only supported in CMake 3.18+. The minimum CMake version we require is 3.16. As a result, if you built with CMake 3.16 - 3.18, you'd get a build failure. This change adds checks around all the spots where the LINK_LANGUAGE keyword is used to avoid this problem. It also adds a description of the CMake HIP Language support to the changelog (previously overlooked).
A recent cmake change added a condition check that avoids evaluating a LINK_LANGUAGE generator expression if HIP_USECXX is OFF (since LINK_LANGUAGE is unsupported in CMake < 3.18). However, it didn't account for the fact that the LINK_LANGUAGE expression is embedded in a conditional generator expression. This change adds an "else" block that ensures the "false" side of the conditional generator expression is still evaluated when the LINK_LANGUAGE code is skipped. This fixes some build issues on Windows. Also updates the readme to include recent changes to the way projects must link to rocPRIM.
…674) * Resolve "Add hipGraph compatibility for device run length encode" * Resolve "Investigate RLE (non trivial runs) failure" * Resolve "Config tuning and dynamic dispatch for device run length encode" * Resolve "Implement type traits interface for rocPRIM's compiler-extended types and third-party's types" * remove ordered block id from run length encode A missing call to __syncthreads() caused this algorithm to occasionally produce an invalid index and cause a crash. Instead of just adding that call, we can remove the use of ordered block id by the regular block id, which is already ordered per HSA guarantees. There still remains an issue with a hang, but this seems a runtime issue and will be investigated separately. * MI200 & MI300 tuning * Optimizations for MI300 unique performance * improve type traits interface docs * Move changlog items to 6.4.0 --------- Co-authored-by: Beatriz Navidad Vilches <beatriz@streamhpc.com> Co-authored-by: Yung-sheng Tu <yung-sheng@streamhpc.com> Co-authored-by: Cenxuan Tian <cenxuan@streamhpc.com> Co-authored-by: Nick Breed <nick@streamhpc.com>
Add tuned configs for gfx1200 and gfx1201 for some algorithms. A few algorithms still require some additional work to get tuning working. These will be added in a follow-up commit.
…683) Co-authored-by: Anton Gorenko <anton@streamhpc.com>
This change adds the remaining tuning configs for gfx1201.
* Support getting device ID from stream hipStreamLegacy (#687) rocprim::get_device_from_stream must retreive the device ID in different ways depending on whether a user-created stream or a built-in stream type is being used. This change modifies the check for built-in stream types so that it includes hipStreamLegacy. This prevents the function from calling hipGetStreamDeviceId(hipStreamLegacy), which currently causes a segmentation fault. * Add hipStreamLegacy HIP version check The hipStreamLegacy stream handle was added in HIP 6.1.0. Add a check to ensure that we don't reference it if rocPRIM is being built using an older version of HIP. * Add get_device_id_from_stream test for hipStreamLegacy Add a test case to ensure that passing hipStreamLegacy to rocprim::get_device_id_from_stream works correctly. Updated the changelog with a note to reflect the fact that hipStreamLegacy can now be used without issue. * Update CHANGELOG.md Co-authored-by: Jeffrey Novotny <jnovotny@amd.com> --------- Co-authored-by: Jeffrey Novotny <jnovotny@amd.com>
NguyenNhuDi
requested review from
a team,
RobsonRLemos,
lawruble13,
stanleytsang-amd and
umfranzw
as code owners
February 25, 2025 23:02
spolifroni-amd
approved these changes
Feb 26, 2025
…ses (#690) Radix sort and merge sort test cases that use the largest data sizes hang or cause an SEH exception on Windows when using the latest version of the HIP SDK. Unfortunately, the HIP SDK issue interferes with our current strategies for skipping tests that require too much memory for a device (querying available device memory or examining the return value from hipMalloc), making them unreliable. While we wait for a fix, this change introduces a workaround by reducing the maximum size these test cases can use on Windows.
Collaborator
|
@NguyenNhuDi Can you rebase and/or fix the merge conflicts here? We need to get some of these hotfixes back into develop for 6.5 |
Collaborator
Author
|
This PR has lots of merge conflicts. PR (#701 ) is the resolved version. |
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
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.
No description provided.