https://github.com/vikramRH updated https://github.com/llvm/llvm-project/pull/210250
>From 0182ee47e5f83baee047c5737f839063df53128e Mon Sep 17 00:00:00 2001 From: vikhegde <[email protected]> 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 _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
