Skip to content

Commit 253b8c1

Browse files
lamb-jassistant-librarian[bot]
authored andcommitted
[rocm-libraries] ROCm/rocm-libraries#7665 (commit 4767702)
[rocPRIM] build benchmark_device_adjacent_difference at -O1 (#7665) Build `benchmark_device_adjacent_difference` at `-O1` to work around ROCm/llvm-project#2616: clang/llc hangs indefinitely in AMDGPU SelectionDAG ISel on this TU at `-O2`/`-O3` for `gfx950`. The hang is a DAG combine freeze cycle in shared SDAG code. Confirmed: with `-O1` the TU builds in ~25 s instead of hanging. Other opt levels and other targets are unaffected. Revert this commit once the upstream codegen fix lands and is in TheRock. Addresses ROCM-24721
1 parent 3361c4b commit 253b8c1

2 files changed

Lines changed: 9 additions & 0 deletions

File tree

CHANGELOG.md

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -108,6 +108,10 @@ Full documentation for rocPRIM is available at [https://rocm.docs.amd.com/projec
108108
* Added `rocprim::apply` which applies a function to a `rocprim::tuple`.
109109

110110

111+
### Known issues
112+
113+
* benchmark_device_adjacent_difference build hangs due to https://github.com/ROCm/llvm-project/issues/2616. Workaround is to build with -O1.
114+
111115
### Changed
112116

113117
* Changed tests to support `ptr-to-const` output in `/test/rocprim/test_device_batch_memcpy.cpp`.

benchmark/CMakeLists.txt

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -172,6 +172,11 @@ add_rocprim_benchmark(benchmark_block_scan.cpp)
172172
add_rocprim_benchmark(benchmark_block_sort.cpp)
173173
add_rocprim_benchmark(benchmark_config_dispatch.cpp)
174174
add_rocprim_benchmark(benchmark_device_adjacent_difference.cpp)
175+
# Workaround for ROCm/llvm-project#2616: DAG combine cycle hangs codegen at
176+
# -O2/-O3 on gfx950 for this TU. Build at -O1 until the codegen fix lands.
177+
if(TARGET benchmark_device_adjacent_difference)
178+
target_compile_options(benchmark_device_adjacent_difference PRIVATE -O1)
179+
endif()
175180
add_rocprim_benchmark(benchmark_device_adjacent_find.cpp)
176181
add_rocprim_benchmark(benchmark_device_batch_memcpy.cpp)
177182
add_rocprim_benchmark(benchmark_device_binary_search.cpp)

0 commit comments

Comments
 (0)