Commit 82065bff authored by one's avatar one
Browse files

[mre] Move MREs to a standalone dir, add dtk-hip-target-mismatch

parent 3a1e2ad8
# dcc 对 `void` 函数误报 `-Wreturn-type`
## 问题描述
DTK26.04 的 `dcc` 编译无返回的 `void` 函数会报错:
```text
void-return.cpp:1:11: error: void function is missing a return statement [-Werror,-Wreturn-type]
void f() {}
^
```
## 复现方式
镜像:
......@@ -14,16 +24,6 @@ harbor.sourcefind.cn:5443/dcu/admin/base/pytorch:2.7.1-ubuntu22.04-dtk26.04-py3.
dcc -std=c++17 -Werror=return-type -c void-return.cpp -o void-return.o
```
## 实际结果
DTK 26.04 的 `dcc` 会报错:
```text
void-return.cpp:1:11: error: void function is missing a return statement [-Werror,-Wreturn-type]
void f() {}
^
```
## 预期结果
这段源码应该能够正常编译。`void f() {}` 是合法的 C++ 代码,`-Wreturn-type` 不应该仅仅因为 `void` 函数没有显式写出 `return;` 就给出诊断。
......
cmake_minimum_required(VERSION 3.23)
project(hip_target_mismatch_minimal CXX)
find_package(hip CONFIG REQUIRED)
add_library(provider_gfx936 OBJECT provider.cpp)
target_compile_options(provider_gfx936 PRIVATE
-x hip
-fgpu-rdc
--offload-arch=gfx936
)
add_executable(consumer
consumer.cpp
$<TARGET_OBJECTS:provider_gfx936>
)
target_link_libraries(consumer PRIVATE hip::device)
target_compile_options(consumer PRIVATE -fgpu-rdc)
target_link_options(consumer PRIVATE -fgpu-rdc)
# HIP target mismatch 最小复现
## 问题描述
使用 Spack 在如下环境编译 Nalu-Wind 依赖的 Trilinos+Kokkos 时遇到编译错误:
- Target: gfx936
- Image: harbor.sourcefind.cn:5443/dcu/admin/base/pytorch:2.7.1-ubuntu22.04-dtk26.04-py3.11
报错信息:
```text
lld: error: undefined protected symbol: kokkos_impl_hip_constant_memory_buffer
>>> referenced 21505 more times
dcc: error: amdgcn-link command failed with exit code 1
```
## 复现方式
拷贝目录中的文件到容器中,然后运行脚本:
```sh
./run.sh
```
容器中使用 DTK26 原版 `hip-config.cmake` 时,预期会有编译错误:
```text
lld: error: undefined protected symbol: kokkos_impl_hip_constant_memory_buffer
>>> referenced by lto.tmp:(read_constant(unsigned long*))
clang++: error: amdgcn-link command failed with exit code 1
```
## 原因分析
这个复现说明问题的关键不是 Kokkos 或 Trilinos 本身,而是两个 HIP 编译单元的 GPU target 集合不一致:
- `provider.cpp` 定义 `kokkos_impl_hip_constant_memory_buffer`,并固定只按 `gfx936` 编译。
- `consumer.cpp` 声明并引用同名 constant memory symbol,编译 target 来自 `hip::device`
- 两个编译单元都使用 `-fgpu-rdc`,最终链接也使用 `-fgpu-rdc`,从而触发 RDC device link 阶段的符号解析。
使用 DTK26 原版 `hip-config.cmake` 时,`hip::device` 会向 consumer 编译和最终链接注入 `gfx906;gfx926;gfx928;gfx936;gfx938` 五个 target。此时 provider 只提供 `gfx936` 的 device symbol 定义,而 consumer 在其它 target 上也引用了这个符号,因此 device link 报错。
使用修复后的 `hip-config.cmake` 时,`hip::device` 在容器中自动检测为 `gfx936`,provider、consumer 和最终 device link 的 target 一致,因此可以正常链接。
#include <hip/hip_runtime.h>
__device__ __constant__ extern unsigned long kokkos_impl_hip_constant_memory_buffer[];
__global__ void read_constant(unsigned long *out) {
out[0] = kokkos_impl_hip_constant_memory_buffer[0];
return;
}
int main() {
return 0;
}
__device__ __constant__ unsigned long kokkos_impl_hip_constant_memory_buffer[1];
#!/usr/bin/env bash
set -eu
src_dir="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
build_dir="${src_dir}/build"
rm -rf "${build_dir}"
cmake \
-S "${src_dir}" \
-B "${build_dir}" \
-DCMAKE_BUILD_TYPE=Release \
-DCMAKE_VERBOSE_MAKEFILE=ON \
-DCMAKE_PREFIX_PATH=/opt/dtk/hip \
-DCMAKE_CXX_COMPILER=/opt/dtk/llvm/bin/clang++
cmake --build "${build_dir}" --target consumer --verbose
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment