Develop stream 2024-06-26

[Snektron]

@Naraenda is taking this over since @Snektron is on holiday.

This PR adds the following merge commits. Emphasis on important changes.

• 46faf89 test: make tests pass with hardened libstc++ (resolves Failing tests with libstdc++ assertions due to unsigned overflow #570)
• 39c9d9b ci(gitlab-ci.yml): re-enable most of the debug builds
• c6ea6aa test: improved test of device histogram_even for half/bfloat16
• a40cbdb fix: fix warnings when building rocprim from rocthrust when using hipcc
• 10c6175 feat: add partial_sort and partial_sort_copy algorithms
• 4a26848 feat: add deterministic versions of rocprim::inclusive_scan, rocprim::exclusive_scan, rocprim::inclusive_scan_by_key, rocprim::exclusive_scan_by_key, and rocprim::reduce_by_key
• 9c7c297 style: format attributes on a separate line
• fc2de52 perf: remove oracle in nth element
• a4ae29c ci(gitlab-ci.yml): compress autotuning artifacts before uploading them
• fab7fd6 refactor: improve code-reuse of rocprim::detail::merge_path
• f586d3e test: add extra tests to nth-element
• d09bae6 feat: add nth-element algorithm
• 43c7ce7 perf: add auto tuning to reduce by key
• d63cd11 feat: add large index support to device segmented reduce
• 110f2d0 test: fix sizes generated by benchmark utility for random segments
• c523e13 ci: build everything using cmake-minimum and cmake-latest workflows
• d4aec9d test: fix benchmark minmax underoverflow
• 941984e test: disable naive batch memcpy in benchmarks
• 94e17b4 docs: document thread level methods
• f90e4e7 refactor: only autotune sleep-based lookback scan for gfx908
• 7b6b04b docs: add documentation for autotuning
• e2aec7e test: improve test logging
• 1617f92 refactor: deprecate raw_storage and replace it by uninitialized_array due to possible undefined behavior with raw_storage
• f7209de refactor: deprecate use of rocprim::thread_load and rocprim::thread_store
• 917b151 fix: fix block histogram test failures
• 441a934 docs: document thread types
• 2455a40 ci(gitlab-ci.yml): disable debug builds due to long compilation times
• 9fc5d18 fix: remove memory fence workaround on gfx10 and gfx11
• c07e212 fix: fix unroll warnings in debug min size releases
• ad746b4 perf: add tuning to device partition
• 9148c77 test: add benchmark flag to specify seed
• 9e849db docs: add contribution guidelines
• 71ba9ce fix: prevent rocprim::adjacent_difference benchmarks from using more than 8 gb vram
• d3b3441 test: fix windows test assertions
• c5a9d07 fix: disable force inline on windows
• 58c96d9 refactor: remove device device warpsize usage in host function in radix sort benchmark generator
• fc2ac73 feat: add an overload of rocprim::match_any which accepts the maximum number of bits as a run-time parameter
• 61ca0d1 test: change default size in adjacent difference benchmark
• 246fe37 fix: various fixes for batch memcpy
• de58f35 perf: add auto tuning to transform

[Naraenda]

Rebased and added a fix that resolves #570

[Naraenda]

Fixed doxygen errors in 8d5b5e8

[Naraenda]

4651b07 Noticed CI failures after targeting a later commit in develop. rocm-cmake was recently-ish changed and now requires building. This commit resolves that when the rocm-cmake is freshly downloaded through FetchContent.

6edfe13 Should fix one of the test failures on MI200 with very large inputs.

Also rebased it, again.

CC @stanleytsang-amd

[stanleytsang-amd]

4651b07 Noticed CI failures after targeting a later commit in develop. rocm-cmake was recently-ish changed and now requires building. This commit resolves that when the rocm-cmake is freshly downloaded through FetchContent.

6edfe13 Should fix one of the test failures on MI200 with very large inputs.

Also rebased it, again.

CC @stanleytsang-amd

Thanks Nara, I will rerun CI right now.

[Naraenda]

I managed to test this on a MI200. It turns out, for reasons I'm not fully aware of, that running hipStreamSynchronize after the cooperative kernel launch is required. Moving that statement out of the debug synchronous branch seems to resolve the hanging on MI200 (4314c76). I'm checking the implementation in ROCm/CLR, but nothing out of the ordinary seems to jump out.

CC @stanleytsang-amd

[stanleytsang-amd]

I managed to test this on a MI200. It turns out, for reasons I'm not fully aware of, that running hipStreamSynchronize after the cooperative kernel launch is required. Moving that statement out of the debug synchronous branch seems to resolve the hanging on MI200 (4314c76). I'm checking the implementation in ROCm/CLR, but nothing out of the ordinary seems to jump out.

CC @stanleytsang-amd

Re-running CI now. It does seem strange that an explicit hipStreamSynchronize is required here since this function only deals with one stream and all kernels are being executed on that stream and I don't see any data race conditions here. Maybe we can make a small reproducer and file a bug report with the HIP team?

[stanleytsang-amd]

I managed to test this on a MI200. It turns out, for reasons I'm not fully aware of, that running hipStreamSynchronize after the cooperative kernel launch is required. Moving that statement out of the debug synchronous branch seems to resolve the hanging on MI200 (4314c76). I'm checking the implementation in ROCm/CLR, but nothing out of the ordinary seems to jump out.

CC @stanleytsang-amd

It seems as though the hang still persists on MI200 on CI, even with the hipStreamSychronize().

[Naraenda]

Rebased & resolved conflicts

[stanleytsang-amd]

Can you update the packaging version at https://github.com/ROCm/rocPRIM/pull/575/files#diff-1e7de1ae2d059d21e1dd75d5812d5a34b0222cef273b7c3a2af62eb747f9d20aL133 to 3.3.0 as well? Thanks.

[Naraenda]

Can you update the packaging version at https://github.com/ROCm/rocPRIM/pull/575/files#diff-1e7de1ae2d059d21e1dd75d5812d5a34b0222cef273b7c3a2af62eb747f9d20aL133 to 3.3.0 as well? Thanks.

Done in 8d06855.

[stanleytsang-amd]

Can you update the packaging version at https://github.com/ROCm/rocPRIM/pull/575/files#diff-1e7de1ae2d059d21e1dd75d5812d5a34b0222cef273b7c3a2af62eb747f9d20aL133 to 3.3.0 as well? Thanks.

Done in 8d06855.

Thanks.