bug 描述 Describe the Bug
继 #78587 在 ROCm/HIP 后端注册了 BF16 卷积内核之后,PaddleOCR-VL-1.5(SigLIP 视觉编码器 + Ernie4.5 解码器)在 AMD GPU 上以 BF16 端到端推理仍然失败。本 issue 跟踪剩余三个相互独立的 HIP BF16 问题,#78587 与 #78664 均未涉及。
与 #78664 的关系:#78664 关注的是「BF16 类型本身在 ROCm 上的可用性」,已由 #78587(注册 BF16 conv 内核)从「卷积」一面解决。本 issue 聚焦在 PaddleOCR-VL 端到端 BF16 推理时仍然命中的 非卷积 缺口(layer_norm、softmax、PIR 融合 pass),与 #78664 不重复。
1. `layer_norm` / `layer_norm_grad` 在 HIP 上未注册 BF16 内核
`paddle/phi/kernels/gpu/layer_norm_kernel.cu` 与 `layer_norm_grad_kernel.cu` 的 `PADDLE_WITH_HIP` 注册块只覆盖 `float` + `phi::float16`。模板实现 `LayerNormKernel<T, GPUContext>` / `LayerNormGradKernel<T, GPUContext>` 本身已支持 `phi::bfloat16`,仅缺少注册。
复现:
```python
import paddle
ln = paddle.nn.LayerNorm(1024).to(dtype="bfloat16")
ln(paddle.randn([2, 64, 1024], dtype="bfloat16"))
RuntimeError: kernel layer_norm for (GPU, ALL_LAYOUT, BFLOAT16) not registered
```
2. `miopenSoftmaxForward_V2` 在 ROCm 7.x 不支持 `miopenBFloat16`
`miopenSoftmaxForward_V2(..., miopenBFloat16, ...)` 返回 `MIOPEN_STATUS_NOT_IMPLEMENTED`。`paddle/phi/kernels/gpudnn/softmax_gpudnn.h` 当 `dim` 超过 warp softmax 阈值时无条件走 MIOpen,BF16 调用直接抛 runtime error。
复现:
```python
import paddle.nn.functional as F
F.softmax(paddle.randn([4, 8192], dtype="bfloat16"), axis=-1)
RuntimeError from MIOpen: MIOPEN_STATUS_NOT_IMPLEMENTED
```
3. `conv2d_add_fuse_pass` / `conv2d_add_act_fuse_pass` 在 HIP wheel 上仍被注册
两个 PIR pass 把 `conv2d + add[+ act]` 改写为 `fused_conv2d_add_act` 算子。`fused_conv2d_add_act` 只有 cuDNN GPUDNN kernel,没有 HIP 实现。在 ROCm 上 pass 改写成功,执行时 dispatch 到不存在的 kernel 失败。注意:这与 #78587 注册的 `conv2d` 是两个独立算子,#78587 不影响该路径。
PaddleX 当前在 `paddlex/inference/models/runners/paddle_static/runner.py` 4 处用 `paddle.is_compiled_with_rocm()` 包裹 `config.delete_pass("conv2d_add_act_fuse_pass")` / `config.delete_pass("conv2d_add_fuse_pass")` 来绕过。
影响 Impact
PaddleX 目前在 `PaddleOCRVLForConditionalGeneration` 上设置 `_keep_in_fp32_modules = ["visual", "mlp_AR"]` 把视觉塔 + 多模态 projector 强制保持 FP32 来绕过(1)(2);并用 4 处 `delete_pass` 绕过(3)。结果:
- 视觉子图 FP32 → 显存近乎翻倍,吞吐下降。
- CUDA 与 ROCm 的推理路径需要分叉维护。
- 阻碍 PaddleOCR-VL 等多模态模型在 AMD GPU 上的 BF16 推理普及。
期望行为 Expected Behavior
ROCm/HIP 后端原生支持 BF16:补齐 `layer_norm` / `softmax` 注册,并把 cuDNN-only 的两个融合 pass 在 HIP 编译下从 PIR pass 注册中跳过,使 PaddleOCR-VL-1.5 在 MI300X 上可以全模型 BF16 推理。
修复 PR
#78711
环境
bug 描述 Describe the Bug
继 #78587 在 ROCm/HIP 后端注册了 BF16 卷积内核之后,PaddleOCR-VL-1.5(SigLIP 视觉编码器 + Ernie4.5 解码器)在 AMD GPU 上以 BF16 端到端推理仍然失败。本 issue 跟踪剩余三个相互独立的 HIP BF16 问题,#78587 与 #78664 均未涉及。
1. `layer_norm` / `layer_norm_grad` 在 HIP 上未注册 BF16 内核
`paddle/phi/kernels/gpu/layer_norm_kernel.cu` 与 `layer_norm_grad_kernel.cu` 的 `PADDLE_WITH_HIP` 注册块只覆盖 `float` + `phi::float16`。模板实现 `LayerNormKernel<T, GPUContext>` / `LayerNormGradKernel<T, GPUContext>` 本身已支持 `phi::bfloat16`,仅缺少注册。
复现:
```python
import paddle
ln = paddle.nn.LayerNorm(1024).to(dtype="bfloat16")
ln(paddle.randn([2, 64, 1024], dtype="bfloat16"))
RuntimeError: kernel
layer_normfor (GPU, ALL_LAYOUT, BFLOAT16) not registered```
2. `miopenSoftmaxForward_V2` 在 ROCm 7.x 不支持 `miopenBFloat16`
`miopenSoftmaxForward_V2(..., miopenBFloat16, ...)` 返回 `MIOPEN_STATUS_NOT_IMPLEMENTED`。`paddle/phi/kernels/gpudnn/softmax_gpudnn.h` 当 `dim` 超过 warp softmax 阈值时无条件走 MIOpen,BF16 调用直接抛 runtime error。
复现:
```python
import paddle.nn.functional as F
F.softmax(paddle.randn([4, 8192], dtype="bfloat16"), axis=-1)
RuntimeError from MIOpen: MIOPEN_STATUS_NOT_IMPLEMENTED
```
3. `conv2d_add_fuse_pass` / `conv2d_add_act_fuse_pass` 在 HIP wheel 上仍被注册
两个 PIR pass 把 `conv2d + add[+ act]` 改写为 `fused_conv2d_add_act` 算子。`fused_conv2d_add_act` 只有 cuDNN GPUDNN kernel,没有 HIP 实现。在 ROCm 上 pass 改写成功,执行时 dispatch 到不存在的 kernel 失败。注意:这与 #78587 注册的 `conv2d` 是两个独立算子,#78587 不影响该路径。
PaddleX 当前在 `paddlex/inference/models/runners/paddle_static/runner.py` 4 处用 `paddle.is_compiled_with_rocm()` 包裹 `config.delete_pass("conv2d_add_act_fuse_pass")` / `config.delete_pass("conv2d_add_fuse_pass")` 来绕过。
影响 Impact
PaddleX 目前在 `PaddleOCRVLForConditionalGeneration` 上设置 `_keep_in_fp32_modules = ["visual", "mlp_AR"]` 把视觉塔 + 多模态 projector 强制保持 FP32 来绕过(1)(2);并用 4 处 `delete_pass` 绕过(3)。结果:
期望行为 Expected Behavior
ROCm/HIP 后端原生支持 BF16:补齐 `layer_norm` / `softmax` 注册,并把 cuDNN-only 的两个融合 pass 在 HIP 编译下从 PIR pass 注册中跳过,使 PaddleOCR-VL-1.5 在 MI300X 上可以全模型 BF16 推理。
修复 PR
#78711
环境