Avoid evaluating CMake LINK_LANGUAGE generator expressions in CMake <3.18 by umfranzw · Pull Request #667 · ROCm/rocPRIM

umfranzw · 2024-12-30T20:48:16Z

Avoid evaluating CMake LINK_LANGUAGE generator expressions in CMake <3.18

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).

spolifroni-amd

Leaving a rewrite for the changelog. The item that I changed should also be moved to the "Changed" section.

spolifroni-amd · 2025-01-06T18:19:23Z

 * Added the parallel `search` and `find_end` device functions similar to `std::search` and `std::find_end`, these functions search for the first and last occurrence of the sequence respectively.
 * Added a parallel device-level function, `rocprim::search_n`, similar to the C++ Standard Library `std::search_n` algorithm.
 * Added new constructors and a `base` function, and added `constexpr` specifier to all functions in `rocprim::reverse_iterator` to improve parity with the C++17 `std::reverse_iterator`.
+* If you are using CMake 3.18+, you can now build using CMake HIP language support. This allows CMake to treat HIP as any other language (like C or CXX), removing the need to set the CXX environment variable. Instead, CMake attempts to find your HIP-aware compiler on its own. To build in this way, please set the USE\_HIPCXX option to ON (-DUSE\_HIPCXX=ON) when invoking CMake.


This should be moved to the "Changed" section.

Suggested change

* If you are using CMake 3.18+, you can now build using CMake HIP language support. This allows CMake to treat HIP as any other language (like C or CXX), removing the need to set the CXX environment variable. Instead, CMake attempts to find your HIP-aware compiler on its own. To build in this way, please set the USE\_HIPCXX option to ON (-DUSE\_HIPCXX=ON) when invoking CMake.

* You can use CMake HIP language support with CMake 3.18 and later. To use HIP language support, run `cmake` with `-DUSE_HIPCXX=ON` instead of setting the `CXX` variable to `HIP`.

Thanks, I've moved the bullet point to the "Changed" section and made the suggested change, which one small caveat: the CXX variable was previously set to the path to a HIP-aware compiler rather than HIP.

…3.18 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).

Snektron · 2025-01-17T09:59:03Z

Can this also be cherry picked on to develop?

…3.18 (ROCm#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).

* Avoid evaluating CMake LINK_LANGUAGE generator expressions in CMake <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). * Fix Windows CMake link issue (#675) 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. * 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 * Fix static asserts and remove template specialization * Fix tests triggering static asserts * Fix benchmarks triggering static asserts * Set explicit return types * Resolve "Investigate RLE (non trivial runs) failure" * Revert "Fix benchmarks triggering static asserts" This reverts commit a1634a9. * Revert "Fix tests triggering static asserts" This reverts commit 24a36cb. * Revert "Fix static asserts and remove template specialization" This reverts commit b2089ee. * Fix formatting and copyright date * Temp fix: histogram dynamic memory failure on windows * add ROCPRIM_ENABLE_ASSERTS to cmake to enable asserts in release build set ROCPRIM_ENABLE_ASSERTS for CI tests * Add missing docs for `block_run_length_decode` * Resolve "Add documentation and test for out-of-range behaviour on merge path search" add test for out of bounds diagonal test for rocprim::merge_path_search document behaviour for out of bounds diagonal test for rocprim::merge_path_search * Resolve "Replace __launch_bounds__ with ROCPRIM_LAUNCH_BOUNDS in rocPRIM API" * Update rmake.py copyright * fix double buffers in device radix sort * Resolve "Replace hipMalloc, hipMallocHelper, and hipFree in tests with device_ptr" * fix device adjacent difference tuning benchmark * Fix warp dpp reduction * Fix static assertion for int128 in ONNX Runtime with GNU extensions * fix GenerateResourceSpec for gfx90a * Resolve "ROCPRIM_TARGET_ARCH deprecation" * Fix warnings when generating sphinx docs * Performance improvements with shuffled warp exchange * log more information during configure * Update documentation generation instructions * Resolve "Replace hipMalloc, hipMallocHelper, and hipFree in benchmarks and tests with device_ptr" Part 4 * Resolve "search_n optimization" * CI changes to speedup development * Remove select_arch and select_arch_case * fix missing equality operator and function constness * Print arguments and timing info in `batch_memcpy_func` when `debug_synchronous` is true * Resolve "Replace hipMalloc, hipMallocHelper, and hipFree in benchmarks and tests with device_ptr" Part 3 * Resolve "Replace hipMalloc, hipMallocHelper, and hipFree in benchmarks and tests with device_ptr" Part 1 * Resolve "Replace hipMalloc, hipMallocHelper, and hipFree in benchmarks and tests with device_ptr" Part 2 * Fix rules overwrite for other then MR pipelines * Resolve "Add stream, memset, array, and unique_ptr support to device_ptr" * Resolve "Reduce amount of test runs based on environment variable." * Resolve "Get clarity on skipping of adjacent difference tests" * Remove windows CI check * Resolve "Add UnrolledThreadLoad and UnrolledThreadCopy functions" * Resolve "General changes CCCL 2.6" * Reduce CI artifact retention time * Resolve "Add thread_load_volatile_pointer to rocprim" * Fix test:install on rocm/dev-ubuntu * Resolve "Nth element - Investigate smaller index sizes where possible" * Enable int128 tests on Windows * Resolve "Benchmark and tuning for 128-bit integers" * Resolve "Verify sorting on device" * Resolve "Expose lookback_scan_state::create and lookback_scan_state::get_storage_size as ROCPRIM_HOST_DEVICE." * add naive in-place api for device merge * To remove error in CI change it to AMDGPU_TARGETS * fix device_ptr assign move and void types * Resolve "Add block- and warp-level inclusive_scan with initial value support" * Fix C++ includes showing up in the docs * Resolve "Performance regressions during tuning: Adjacent find" * CI skip label fix * fix doxygen warning in merge.rst * Resolve "Autotuning for segmented reduce" * Resolve "Update thread_store for parity" * Resolve "Fix compile errors in "rocprim-release-tune" mode" * Documentation to CI upload * Resolve "Remove ASM from thread_load and thread_store" * Add documentation steps to developer guidelines * Resolve "Refactor lookback_scan_state API" * Adopt symbol visibility rules in rocPRIM * remove nm debug messages from cmake * Remove extra dashes in the documentation its \param and \tparam * update changelog * fix(test_device_merge_inplace.cpp): prevent merge inplace tests to tests to grow too big * Fix test_linking: install the test's libraries too * chore(rmake.py): bump copyright * fix(example_type_traits_interface.cpp): explicit return from main * Fix unintentionally exported shared libs of test_linking rocm_install adds all libraries to ${PROJECT_NAME_LOWER}-targets (i.e. rocprim-targets.cmake) unless a custom export file is specified. --------- Co-authored-by: Wayne Franz <wayfranz@amd.com> 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: Jaap Blok <jaap@streamhpc.com> Co-authored-by: Nick Breed <nick@streamhpc.com> Co-authored-by: Bence Parajdi <bence@streamhpc.com> Co-authored-by: Sander Bos <sander@streamhpc.com> Co-authored-by: Nara Prasetya <nara@streamhpc.com> Co-authored-by: Anton Gorenko <anton@streamhpc.com> Co-authored-by: Milo Lurati <milo@streamhpc.com> Co-authored-by: Saiyang Zhang <saiyang@streamhpc.com> Co-authored-by: Enrico Degregori <enrico@streamhpc.com> Co-authored-by: Borys Petrov <borys@streamhpc.com>

* Add atomic_add overload for double (#652) 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. * Make temp_storage functions __host__ and __device__ (#658) * Enable CMake HIP language (#655) * Restart on HIP language support * linking changes * rearrange --------- Co-authored-by: Lauren Wrubleski <Lauren.Wrubleski@amd.com> * Fix for rocprim::batch_copy when using iterators (#662) * 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. * Fix warp dpp reduction (#660) Co-authored-by: Robin Voetter <robin@streamhpc.com> * Workaround for shared memory write ordering compiler issue (#672) 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. * Avoid evaluating CMake LINK_LANGUAGE generator expressions in CMake <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). * Fix Windows CMake link issue (#675) 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. * MI300 radix sort, run length encode, and reduce by key optimizations (#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 (#681) 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. * Fix static assertion for int128 in ONNX Runtime with GNU extensions (#683) Co-authored-by: Anton Gorenko <anton@streamhpc.com> * Add remaining gfx1201 tuning configs (#684) This change adds the remaining tuning configs for gfx1201. * Accept hip stream legacy (#689) * 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> * Workaround memory allocation issues on Windows for large size test cases (#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. * used 678 version of test_device_run_length_encode * got rid of duplicated code * removed ; on if statement --------- Co-authored-by: Wayne Franz <wayfranz@amd.com> Co-authored-by: Nick Breed <78807921+NB4444@users.noreply.github.com> Co-authored-by: Lauren Wrubleski <Lauren.Wrubleski@amd.com> Co-authored-by: Beatriz Navidad Vilches <61422851+Beanavil@users.noreply.github.com> Co-authored-by: Robin Voetter <robin@streamhpc.com> 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> Co-authored-by: Anton Gorenko <anton@streamhpc.com> Co-authored-by: Jeffrey Novotny <jnovotny@amd.com>

* Avoid evaluating CMake LINK_LANGUAGE generator expressions in CMake <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). * Fix Windows CMake link issue (#675) 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. * 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 * Fix static asserts and remove template specialization * Fix tests triggering static asserts * Fix benchmarks triggering static asserts * Set explicit return types * Resolve "Investigate RLE (non trivial runs) failure" * Revert "Fix benchmarks triggering static asserts" This reverts commit a1634a9. * Revert "Fix tests triggering static asserts" This reverts commit 24a36cb. * Revert "Fix static asserts and remove template specialization" This reverts commit b2089ee. * Fix formatting and copyright date * Temp fix: histogram dynamic memory failure on windows * add ROCPRIM_ENABLE_ASSERTS to cmake to enable asserts in release build set ROCPRIM_ENABLE_ASSERTS for CI tests * Add missing docs for `block_run_length_decode` * Resolve "Add documentation and test for out-of-range behaviour on merge path search" add test for out of bounds diagonal test for rocprim::merge_path_search document behaviour for out of bounds diagonal test for rocprim::merge_path_search * Resolve "Replace __launch_bounds__ with ROCPRIM_LAUNCH_BOUNDS in rocPRIM API" * Update rmake.py copyright * fix double buffers in device radix sort * Resolve "Replace hipMalloc, hipMallocHelper, and hipFree in tests with device_ptr" * fix device adjacent difference tuning benchmark * Fix warp dpp reduction * Fix static assertion for int128 in ONNX Runtime with GNU extensions * fix GenerateResourceSpec for gfx90a * Resolve "ROCPRIM_TARGET_ARCH deprecation" * Fix warnings when generating sphinx docs * Performance improvements with shuffled warp exchange * log more information during configure * Update documentation generation instructions * Resolve "Replace hipMalloc, hipMallocHelper, and hipFree in benchmarks and tests with device_ptr" Part 4 * Resolve "search_n optimization" * CI changes to speedup development * Remove select_arch and select_arch_case * fix missing equality operator and function constness * Print arguments and timing info in `batch_memcpy_func` when `debug_synchronous` is true * Resolve "Replace hipMalloc, hipMallocHelper, and hipFree in benchmarks and tests with device_ptr" Part 3 * Resolve "Replace hipMalloc, hipMallocHelper, and hipFree in benchmarks and tests with device_ptr" Part 1 * Resolve "Replace hipMalloc, hipMallocHelper, and hipFree in benchmarks and tests with device_ptr" Part 2 * Fix rules overwrite for other then MR pipelines * Resolve "Add stream, memset, array, and unique_ptr support to device_ptr" * Resolve "Reduce amount of test runs based on environment variable." * Resolve "Get clarity on skipping of adjacent difference tests" * Remove windows CI check * Resolve "Add UnrolledThreadLoad and UnrolledThreadCopy functions" * Resolve "General changes CCCL 2.6" * Reduce CI artifact retention time * Resolve "Add thread_load_volatile_pointer to rocprim" * Fix test:install on rocm/dev-ubuntu * Resolve "Nth element - Investigate smaller index sizes where possible" * Enable int128 tests on Windows * Resolve "Benchmark and tuning for 128-bit integers" * Resolve "Verify sorting on device" * Resolve "Expose lookback_scan_state::create and lookback_scan_state::get_storage_size as ROCPRIM_HOST_DEVICE." * add naive in-place api for device merge * To remove error in CI change it to AMDGPU_TARGETS * fix device_ptr assign move and void types * Resolve "Add block- and warp-level inclusive_scan with initial value support" * Fix C++ includes showing up in the docs * Resolve "Performance regressions during tuning: Adjacent find" * CI skip label fix * fix doxygen warning in merge.rst * Resolve "Autotuning for segmented reduce" * Resolve "Update thread_store for parity" * Resolve "Fix compile errors in "rocprim-release-tune" mode" * Documentation to CI upload * Resolve "Remove ASM from thread_load and thread_store" * Add documentation steps to developer guidelines * Resolve "Refactor lookback_scan_state API" * Adopt symbol visibility rules in rocPRIM * remove nm debug messages from cmake * Remove extra dashes in the documentation its \param and \tparam * update changelog * fix(test_device_merge_inplace.cpp): prevent merge inplace tests to tests to grow too big * Fix test_linking: install the test's libraries too * chore(rmake.py): bump copyright * fix(example_type_traits_interface.cpp): explicit return from main * Fix unintentionally exported shared libs of test_linking rocm_install adds all libraries to ${PROJECT_NAME_LOWER}-targets (i.e. rocprim-targets.cmake) unless a custom export file is specified. --------- Co-authored-by: Wayne Franz <wayfranz@amd.com> 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: Jaap Blok <jaap@streamhpc.com> Co-authored-by: Nick Breed <nick@streamhpc.com> Co-authored-by: Bence Parajdi <bence@streamhpc.com> Co-authored-by: Sander Bos <sander@streamhpc.com> Co-authored-by: Nara Prasetya <nara@streamhpc.com> Co-authored-by: Anton Gorenko <anton@streamhpc.com> Co-authored-by: Milo Lurati <milo@streamhpc.com> Co-authored-by: Saiyang Zhang <saiyang@streamhpc.com> Co-authored-by: Enrico Degregori <enrico@streamhpc.com> Co-authored-by: Borys Petrov <borys@streamhpc.com> [ROCm/rocPRIM commit: 678701d]

* Add atomic_add overload for double (#652) 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. * Make temp_storage functions __host__ and __device__ (#658) * Enable CMake HIP language (#655) * Restart on HIP language support * linking changes * rearrange --------- Co-authored-by: Lauren Wrubleski <Lauren.Wrubleski@amd.com> * Fix for rocprim::batch_copy when using iterators (#662) * 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. * Fix warp dpp reduction (#660) Co-authored-by: Robin Voetter <robin@streamhpc.com> * Workaround for shared memory write ordering compiler issue (#672) 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. * Avoid evaluating CMake LINK_LANGUAGE generator expressions in CMake <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). * Fix Windows CMake link issue (#675) 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. * MI300 radix sort, run length encode, and reduce by key optimizations (#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 (#681) 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. * Fix static assertion for int128 in ONNX Runtime with GNU extensions (#683) Co-authored-by: Anton Gorenko <anton@streamhpc.com> * Add remaining gfx1201 tuning configs (#684) This change adds the remaining tuning configs for gfx1201. * Accept hip stream legacy (#689) * 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> * Workaround memory allocation issues on Windows for large size test cases (#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. * used 678 version of test_device_run_length_encode * got rid of duplicated code * removed ; on if statement --------- Co-authored-by: Wayne Franz <wayfranz@amd.com> Co-authored-by: Nick Breed <78807921+NB4444@users.noreply.github.com> Co-authored-by: Lauren Wrubleski <Lauren.Wrubleski@amd.com> Co-authored-by: Beatriz Navidad Vilches <61422851+Beanavil@users.noreply.github.com> Co-authored-by: Robin Voetter <robin@streamhpc.com> 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> Co-authored-by: Anton Gorenko <anton@streamhpc.com> Co-authored-by: Jeffrey Novotny <jnovotny@amd.com> [ROCm/rocPRIM commit: 07fda10]

umfranzw requested a review from NguyenNhuDi December 30, 2024 20:48

umfranzw requested review from a team, RobsonRLemos, lawruble13 and stanleytsang-amd as code owners December 30, 2024 20:48

spolifroni-amd suggested changes Jan 6, 2025

View reviewed changes

umfranzw force-pushed the fix_cmake_link_language_rebased branch from 10d475c to 6903b30 Compare January 6, 2025 18:56

umfranzw requested a review from spolifroni-amd January 16, 2025 19:08

stanleytsang-amd approved these changes Jan 16, 2025

View reviewed changes

umfranzw merged commit 0788f75 into ROCm:release-staging/rocm-rel-6.4 Jan 16, 2025

Snektron mentioned this pull request Jan 17, 2025

MI300 radix sort, run length encode, and reduce by key optimizations #674

Merged

Navigation Menu

Repository navigation

Avoid evaluating CMake LINK_LANGUAGE generator expressions in CMake <3.18 - #667

Avoid evaluating CMake LINK_LANGUAGE generator expressions in CMake <3.18#667
umfranzw merged 1 commit into
ROCm:release-staging/rocm-rel-6.4from
umfranzw:fix_cmake_link_language_rebased

umfranzw commented Dec 30, 2024

Uh oh!

spolifroni-amd left a comment

Uh oh!

spolifroni-amd Jan 6, 2025

Uh oh!

umfranzw Jan 6, 2025 •

edited

Loading

Uh oh!

Snektron commented Jan 17, 2025

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

	* If you are using CMake 3.18+, you can now build using CMake HIP language support. This allows CMake to treat HIP as any other language (like C or CXX), removing the need to set the CXX environment variable. Instead, CMake attempts to find your HIP-aware compiler on its own. To build in this way, please set the USE\_HIPCXX option to ON (-DUSE\_HIPCXX=ON) when invoking CMake.
	* You can use CMake HIP language support with CMake 3.18 and later. To use HIP language support, run `cmake` with `-DUSE_HIPCXX=ON` instead of setting the `CXX` variable to `HIP`.

Conversation

umfranzw commented Dec 30, 2024

Uh oh!

spolifroni-amd left a comment

Choose a reason for hiding this comment

Uh oh!

spolifroni-amd Jan 6, 2025

Choose a reason for hiding this comment

Uh oh!

umfranzw Jan 6, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Snektron commented Jan 17, 2025

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

umfranzw Jan 6, 2025 •

edited

Loading