From 2d8384f8bfa176b7984cc23859e98f7f11ff17cf Mon Sep 17 00:00:00 2001 From: austin1997 <18709560+austin1997@users.noreply.github.com> Date: Sat, 18 Apr 2026 03:07:19 +0000 Subject: [PATCH 1/3] [ROCm] Register BF16 kernel for layer_norm on HIP Add phi::bfloat16 to the layer_norm / layer_norm_grad kernel registrations under PADDLE_WITH_HIP so the existing templated implementation is exposed for BF16 inputs on ROCm. Matches the FLOAT16 treatment of the mean/variance output dtype (promoted to FLOAT32 for numerical stability). Unblocks BF16 inference of the PaddleOCR-VL-1.5 SigLIP-style vision encoder on MI300X (gfx942), which previously required PaddleX to keep the whole visual + mlp_AR subgraph in FP32 via _keep_in_fp32_modules. --- paddle/phi/kernels/gpu/layer_norm_grad_kernel.cu | 6 ++++-- paddle/phi/kernels/gpu/layer_norm_kernel.cu | 9 +++++++-- 2 files changed, 11 insertions(+), 4 deletions(-) diff --git a/paddle/phi/kernels/gpu/layer_norm_grad_kernel.cu b/paddle/phi/kernels/gpu/layer_norm_grad_kernel.cu index f3d4fefa011d0..da975951e6093 100644 --- a/paddle/phi/kernels/gpu/layer_norm_grad_kernel.cu +++ b/paddle/phi/kernels/gpu/layer_norm_grad_kernel.cu @@ -257,8 +257,10 @@ PD_REGISTER_KERNEL(layer_norm_grad, ALL_LAYOUT, phi::LayerNormGradKernel, float, - phi::float16) { - if (kernel_key.dtype() == phi::DataType::FLOAT16) { + phi::float16, + phi::bfloat16) { + if (kernel_key.dtype() == phi::DataType::FLOAT16 || + kernel_key.dtype() == phi::DataType::BFLOAT16) { kernel->OutputAt(1).SetDataType(phi::DataType::FLOAT32); kernel->OutputAt(2).SetDataType(phi::DataType::FLOAT32); } diff --git a/paddle/phi/kernels/gpu/layer_norm_kernel.cu b/paddle/phi/kernels/gpu/layer_norm_kernel.cu index f1285732c95ec..25463895f75da 100644 --- a/paddle/phi/kernels/gpu/layer_norm_kernel.cu +++ b/paddle/phi/kernels/gpu/layer_norm_kernel.cu @@ -786,8 +786,13 @@ template PADDLE_API void LayerNormKernel( #ifdef PADDLE_WITH_HIP // MIOPEN do not support double -PD_REGISTER_KERNEL( - layer_norm, GPU, ALL_LAYOUT, phi::LayerNormKernel, float, phi::float16) { +PD_REGISTER_KERNEL(layer_norm, + GPU, + ALL_LAYOUT, + phi::LayerNormKernel, + float, + phi::float16, + phi::bfloat16) { kernel->OutputAt(1).SetDataType(phi::DataType::UNDEFINED); kernel->OutputAt(2).SetDataType(phi::DataType::UNDEFINED); } From 9a6f793821ecb59fb81ef71cdcca21732d987fc0 Mon Sep 17 00:00:00 2001 From: austin1997 <18709560+austin1997@users.noreply.github.com> Date: Sat, 18 Apr 2026 03:07:33 +0000 Subject: [PATCH 2/3] [ROCm] Route BF16 softmax through matrix kernel (MIOpen NOT_IMPLEMENTED) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit MIOpen (as of ROCm 7.x) returns MIOPEN_STATUS_NOT_IMPLEMENTED for miopenSoftmaxForward_V2 with miopenBFloat16, so the gpudnn softmax path cannot be used for BF16 on HIP. When the input dim exceeds the warp softmax cap, route BF16 through the existing matrix softmax kernel instead of letting the call fall into the MIOpen branch. Also gate the CUDNN_VERSION < 8100 BF16 fallback specialization on !defined(PADDLE_WITH_HIP) — that branch dispatched into MIOpen too and would trip the same NOT_IMPLEMENTED failure on ROCm. --- paddle/phi/kernels/gpudnn/softmax_gpudnn.h | 13 +++++++++++-- 1 file changed, 11 insertions(+), 2 deletions(-) diff --git a/paddle/phi/kernels/gpudnn/softmax_gpudnn.h b/paddle/phi/kernels/gpudnn/softmax_gpudnn.h index 4752513c483f5..c0c26949d4764 100644 --- a/paddle/phi/kernels/gpudnn/softmax_gpudnn.h +++ b/paddle/phi/kernels/gpudnn/softmax_gpudnn.h @@ -1325,7 +1325,7 @@ void LaunchKeMatrixSoftmaxForwardKernel(const GPUContext& dev_ctx, <<>>(out, input, dim_size); } -#if CUDNN_VERSION < 8100 +#if !defined(PADDLE_WITH_HIP) && CUDNN_VERSION < 8100 template <> inline void LaunchSoftmaxForwardCudnnKernel( const GPUContext& dev_ctx, @@ -2811,7 +2811,16 @@ void SoftmaxForwardCUDAKernelDriverImpl(const GPUContext& dev_ctx, dim_log2); } } else { - if (dim >= MATRIX_SOFTMAX_THRESHOLD) { + bool use_matrix_kernel = dim >= MATRIX_SOFTMAX_THRESHOLD; +#ifdef PADDLE_WITH_HIP + // MIOpen (as of ROCm 7.x) returns MIOPEN_STATUS_NOT_IMPLEMENTED for + // miopenSoftmaxForward_V2 with miopenBFloat16. Route BF16 through the + // matrix softmax kernel for any dim that exceeds the warp-softmax cap. + if (std::is_same::value) { + use_matrix_kernel = true; + } +#endif + if (use_matrix_kernel) { LaunchKeMatrixSoftmaxForwardKernel( dev_ctx, out_data, x.data(), N, dim); } else { From ba8637e1aa191c772ed0e62169b978fcdd97707f Mon Sep 17 00:00:00 2001 From: austin1997 <18709560+austin1997@users.noreply.github.com> Date: Sat, 18 Apr 2026 03:07:49 +0000 Subject: [PATCH 3/3] [ROCm] Skip cuDNN-only conv2d fusion passes on HIP conv2d_add_fuse_pass and conv2d_add_act_fuse_pass rewrite conv2d+add[+act] into the fused_conv2d_add_act op, which has only a cuDNN GPUDNN kernel. On ROCm the rewrite succeeds but kernel dispatch later fails because no HIP kernel is registered, so PaddleX currently works around this by calling config.delete_pass("conv2d_add_act_fuse_pass") and config.delete_pass("conv2d_add_fuse_pass") under paddle.is_compiled_with_rocm() in paddlex/inference/models/runners/paddle_static/runner.py. Gate both the pass registration (REGISTER_IR_PASS / USE_PIR_PASS) and the pass-builder inclusion on PADDLE_WITH_CUDA so the rewrite never runs on HIP builds, making the PaddleX delete_pass calls unnecessary. --- paddle/fluid/inference/api/paddle_pass_builder.cc | 4 ++++ paddle/fluid/pir/transforms/gpu/conv2d_add_act_fuse_pass.cc | 6 ++++++ paddle/fluid/pir/transforms/gpu/conv2d_add_fuse_pass.cc | 6 ++++++ paddle/fluid/pir/transforms/passes.h | 2 ++ 4 files changed, 18 insertions(+) diff --git a/paddle/fluid/inference/api/paddle_pass_builder.cc b/paddle/fluid/inference/api/paddle_pass_builder.cc index 2cfab1eaf1cf3..ad7f59b11f98c 100644 --- a/paddle/fluid/inference/api/paddle_pass_builder.cc +++ b/paddle/fluid/inference/api/paddle_pass_builder.cc @@ -635,8 +635,12 @@ const std::vector kPirGpuPasses{ // Operator fusion pass "silu_fuse_pass", "conv2d_bn_fuse_pass", +#ifdef PADDLE_WITH_CUDA + // conv2d_add(_act)_fuse_pass lower to the fused_conv2d_add_act op, which + // only has a cuDNN GPUDNN kernel. Skip on ROCm/HIP. "conv2d_add_act_fuse_pass", "conv2d_add_fuse_pass", +#endif "embedding_eltwise_layernorm_fuse_pass", "fused_rotary_position_embedding_pass", "fused_flash_attn_pass", diff --git a/paddle/fluid/pir/transforms/gpu/conv2d_add_act_fuse_pass.cc b/paddle/fluid/pir/transforms/gpu/conv2d_add_act_fuse_pass.cc index d81ef58c2eecd..9adfd1282ea6e 100644 --- a/paddle/fluid/pir/transforms/gpu/conv2d_add_act_fuse_pass.cc +++ b/paddle/fluid/pir/transforms/gpu/conv2d_add_act_fuse_pass.cc @@ -349,4 +349,10 @@ std::unique_ptr CreateConv2dAddActFusePass() { } // namespace pir +// The fused_conv2d_add_act op this pass produces only has a cuDNN +// (PADDLE_WITH_CUDA) GPUDNN kernel, so the pass is a no-op on other backends. +// Skip registration on ROCm/HIP to avoid applying the rewrite and later +// failing at kernel-dispatch time. +#ifdef PADDLE_WITH_CUDA REGISTER_IR_PASS(conv2d_add_act_fuse_pass, Conv2dAddActFusePass); +#endif diff --git a/paddle/fluid/pir/transforms/gpu/conv2d_add_fuse_pass.cc b/paddle/fluid/pir/transforms/gpu/conv2d_add_fuse_pass.cc index 475eb426e1de9..1c5ff1e52bfd9 100644 --- a/paddle/fluid/pir/transforms/gpu/conv2d_add_fuse_pass.cc +++ b/paddle/fluid/pir/transforms/gpu/conv2d_add_fuse_pass.cc @@ -221,4 +221,10 @@ std::unique_ptr CreateConv2dAddFusePass() { } } // namespace pir +// The fused_conv2d_add_act op this pass produces only has a cuDNN +// (PADDLE_WITH_CUDA) GPUDNN kernel, so the pass is a no-op on other backends. +// Skip registration on ROCm/HIP to avoid applying the rewrite and later +// failing at kernel-dispatch time. +#ifdef PADDLE_WITH_CUDA REGISTER_IR_PASS(conv2d_add_fuse_pass, Conv2dAddFusePass); +#endif diff --git a/paddle/fluid/pir/transforms/passes.h b/paddle/fluid/pir/transforms/passes.h index 5b34175ac4f54..31f0e514cfa54 100644 --- a/paddle/fluid/pir/transforms/passes.h +++ b/paddle/fluid/pir/transforms/passes.h @@ -34,8 +34,10 @@ USE_PIR_PASS(matmul_add_act_fuse_pass); USE_PIR_PASS(silu_fuse_pass); USE_PIR_PASS(fc_elementwise_layernorm_fuse_pass); USE_PIR_PASS(conv2d_bn_fuse_pass); +#ifdef PADDLE_WITH_CUDA USE_PIR_PASS(conv2d_add_fuse_pass); USE_PIR_PASS(conv2d_add_act_fuse_pass); +#endif USE_PIR_PASS(embedding_eltwise_layernorm_fuse_pass); USE_PIR_PASS(add_norm_fuse_pass); USE_PIR_PASS(group_norm_silu_fuse_pass);