[llvm-branch-commits] [clang] [Clang][LTO][AMDGPU] start using new pass manager for backend codegen (PR #210250)
Vikram Hegde via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Tue Aug 18 00:58:08 PDT 2026
https://github.com/vikramRH updated https://github.com/llvm/llvm-project/pull/210250
>From 0182ee47e5f83baee047c5737f839063df53128e Mon Sep 17 00:00:00 2001
From: vikhegde <vikram.hegde at amd.com>
Date: Thu, 16 Jul 2026 16:27:28 +0530
Subject: [PATCH] [Clang][LTO][AMDGPU] Force disable backend NPM when GIsel is
requested
---
clang/lib/Driver/ToolChains/Clang.cpp | 27 +++++++++++++++++--
.../CodeGenCUDA/atomics-remarks-gfx90a.cu | 5 ++++
clang/test/CodeGenHIP/ballot.cpp | 1 +
clang/test/Driver/hip-options.hip | 22 +++++++++++++++
4 files changed, 53 insertions(+), 2 deletions(-)
diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp
index 54583fe3abbd8..39a6d39504689 100644
--- a/clang/lib/Driver/ToolChains/Clang.cpp
+++ b/clang/lib/Driver/ToolChains/Clang.cpp
@@ -5337,6 +5337,17 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
}
}
+ // TODO: remove this once GlobalISel is supported in NewPM.
+ if (getToolChain().getTriple().isAMDGCN()) {
+ if (const Arg *A = Args.getLastArg(options::OPT_fglobal_isel,
+ options::OPT_fno_global_isel)) {
+ if (A->getOption().matches(options::OPT_fglobal_isel)) {
+ CmdArgs.push_back(
+ Args.MakeArgString("-fenable-new-pm-codegen=force-disable"));
+ }
+ }
+ }
+
if (IsCuda && !IsCudaDevice && !UsesLLVMOffloading) {
// We need to figure out which CUDA version we're compiling for, as that
// determines how we load and launch GPU kernels.
@@ -9985,12 +9996,24 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA,
CmdArgs.push_back(Args.MakeArgString(
"--device-linker=" + TC->getTripleString() + "=" + Arg));
+ // TODO: remove this once GlobalISel is supported in NewPM.
+ if (TC->getTriple().isAMDGCN()) {
+ if (const Arg *A = Args.getLastArg(options::OPT_fglobal_isel,
+ options::OPT_fno_global_isel)) {
+ if (A->getOption().matches(options::OPT_fglobal_isel)) {
+ CmdArgs.push_back(Args.MakeArgString(
+ "--device-linker=" + TC->getTripleString() +
+ "=-plugin-opt=-enable-npm-for-backend=force-disable"));
+ }
+ }
+ }
+
// Forward the LTO mode for this toolchain.
auto DeviceLTOMode = TC->getLTOMode(ToolChainArgs, Kind);
- if (DeviceLTOMode == LTOK_Full)
+ if (DeviceLTOMode == LTOK_Full) {
CmdArgs.push_back(Args.MakeArgString(
"--device-compiler=" + TC->getTripleString() + "=-flto=full"));
- else if (DeviceLTOMode == LTOK_Thin) {
+ } else if (DeviceLTOMode == LTOK_Thin) {
CmdArgs.push_back(Args.MakeArgString(
"--device-compiler=" + TC->getTripleString() + "=-flto=thin"));
if (TC->getTriple().isAMDGPU()) {
diff --git a/clang/test/CodeGenCUDA/atomics-remarks-gfx90a.cu b/clang/test/CodeGenCUDA/atomics-remarks-gfx90a.cu
index 92869a96d6e8c..3819dfb964dd1 100644
--- a/clang/test/CodeGenCUDA/atomics-remarks-gfx90a.cu
+++ b/clang/test/CodeGenCUDA/atomics-remarks-gfx90a.cu
@@ -2,6 +2,10 @@
// RUN: -Rpass=atomic-expand -S -o - 2>&1 | \
// RUN: FileCheck %s --check-prefix=GFX90A-CAS
+// RUN: %clang_cc1 %s -triple=amdgpu9.0a-amd-amdhsa -fenable-new-pm-codegen=force-on -fcuda-is-device \
+// RUN: -Rpass=atomic-expand -S -o - 2>&1 | \
+// RUN: FileCheck %s --check-prefix=GFX90A-CAS
+
// REQUIRES: amdgpu-registered-target
#include "Inputs/cuda.h"
@@ -11,6 +15,7 @@
// GFX90A-CAS-LABEL: _Z14atomic_add_casPf
// GFX90A-CAS: flat_atomic_cmpswap
// GFX90A-CAS: s_cbranch_execnz
+
__device__ float atomic_add_cas(float *p) {
return __atomic_fetch_add(p, 1.0f, memory_order_relaxed);
}
diff --git a/clang/test/CodeGenHIP/ballot.cpp b/clang/test/CodeGenHIP/ballot.cpp
index 69480fcf31b2d..9d4fa838a6897 100644
--- a/clang/test/CodeGenHIP/ballot.cpp
+++ b/clang/test/CodeGenHIP/ballot.cpp
@@ -1,6 +1,7 @@
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgpu9.00-amd-amdhsa -aux-triple x86_64-pc-windows-msvc -x hip -emit-llvm -fcuda-is-device -o - %s | FileCheck %s
// RUN: %clang_cc1 -triple amdgpu9.00-amd-amdhsa -aux-triple x86_64-pc-windows-msvc -x hip -S -fcuda-is-device -o - %s | FileCheck %s --check-prefix=GFX9
+// RUN: %clang_cc1 -triple amdgpu9.00-amd-amdhsa -fenable-new-pm-codegen=force-on -aux-triple x86_64-pc-windows-msvc -x hip -S -fcuda-is-device -o - %s | FileCheck %s --check-prefix=GFX9
// Unlike OpenCL, HIP depends on the C++ interpration of "unsigned long", which
// is 64 bits long on Linux and 32 bits long on Windows. The return type of the
diff --git a/clang/test/Driver/hip-options.hip b/clang/test/Driver/hip-options.hip
index 25d5d7f7673d1..2effb646ee394 100644
--- a/clang/test/Driver/hip-options.hip
+++ b/clang/test/Driver/hip-options.hip
@@ -32,6 +32,28 @@
// BUFF: "-cc1" "-triple" "x86_64-unknown-linux-gnu"{{.*}} "-E" {{.*}}
// BUFF: "-cc1" "-triple" "x86_64-unknown-linux-gnu"{{.*}}"-x" "ir"
+// GlobalISel is not yet supported by the New PM CodeGen pipeline, so the driver
+// force-disables it at cc1 when GlobalISel is requested.
+// RUN: %clang -### -nogpulib -nogpuinc --target=x86_64-unknown-linux-gnu -fglobal-isel %s 2>&1 | FileCheck -check-prefix=NPM_GISEL %s
+// NPM_GISEL: "-cc1"{{.*}} "-triple" "amdgcn-amd-amdhsa"{{.*}} "-fenable-new-pm-codegen=force-disable"
+
+// Check that the force-disable reaches the device cc1 when clang runs the backend
+// directly instead of the default lld/LTO path.
+// RUN: %clang -### -nogpulib -nogpuinc --target=x86_64-unknown-linux-gnu --offload-arch=gfx900 -fglobal-isel --cuda-device-only -c %s 2>&1 | FileCheck -check-prefix=NPM_GISEL_BACKEND %s
+// NPM_GISEL_BACKEND: "-cc1"{{.*}} "-triple" "amdgcn-amd-amdhsa"{{.*}} "-fenable-new-pm-codegen=force-disable"{{.*}} "-emit-obj"
+
+// Without GlobalISel the driver leaves the New PM CodeGen option untouched at cc1.
+// RUN: %clang -### -nogpulib -nogpuinc --target=x86_64-unknown-linux-gnu %s 2>&1 | FileCheck -check-prefix=NPM_DEFAULT %s
+// NPM_DEFAULT-NOT: -fenable-new-pm-codegen
+
+// GlobalISel also force-disables the New PM for the device LTO backend.
+// RUN: %clang -### -nogpulib -nogpuinc --target=x86_64-unknown-linux-gnu -foffload-lto=full -fglobal-isel %s 2>&1 | FileCheck -check-prefix=NPM_LTO_GISEL %s
+// NPM_LTO_GISEL: "--device-linker=amdgcn-amd-amdhsa=-plugin-opt=-enable-npm-for-backend=force-disable"
+
+// Without GlobalISel the device LTO backend option is left untouched.
+// RUN: %clang -### -nogpulib -nogpuinc --target=x86_64-unknown-linux-gnu -foffload-lto=full %s 2>&1 | FileCheck -check-prefix=NPM_LTO_DEFAULT %s
+// NPM_LTO_DEFAULT-NOT: -enable-npm-for-backend
+
// RUN: %clang -### -x hip -nogpulib -nogpuinc --target=x86_64-pc-windows-msvc -fms-extensions \
// RUN: -mllvm -amdgpu-early-inline-all=true %s 2>&1 | \
// RUN: FileCheck -check-prefix=MLLVM %s
More information about the llvm-branch-commits
mailing list