Author: Joseph Huber Date: 2026-08-05T10:52:09-05:00 New Revision: bf160ecef41a05cd4745cb90266f83b6c8ac97fe
URL: https://github.com/llvm/llvm-project/commit/bf160ecef41a05cd4745cb90266f83b6c8ac97fe DIFF: https://github.com/llvm/llvm-project/commit/bf160ecef41a05cd4745cb90266f83b6c8ac97fe.diff LOG: [CUDA] Add `--cuda-emit-nvcc-abi` flag to use old ABI with the new driver (#214085) Summary: We want to delete the old driver path, but we need to allow people use the legacy ABI for emitting binaries that will be compiled with NVCC, or in cases where you want to use the old clang handling in CMake or similar. Added: clang/test/CodeGenCUDA/cuda-emit-nvcc-abi.cu Modified: clang/docs/ReleaseNotes.md clang/include/clang/Basic/LangOptions.def clang/include/clang/Options/Options.td clang/lib/CodeGen/CGCUDANV.cpp clang/lib/Driver/Driver.cpp clang/lib/Driver/ToolChains/Clang.cpp clang/test/Driver/cuda-openmp-driver.cu Removed: ################################################################################ diff --git a/clang/docs/ReleaseNotes.md b/clang/docs/ReleaseNotes.md index a02ecc41b2292..7a6762df45680 100644 --- a/clang/docs/ReleaseNotes.md +++ b/clang/docs/ReleaseNotes.md @@ -498,6 +498,9 @@ features cannot lower the translation-unit ABI level; #### CUDA Support +- Added `--cuda-emit-nvcc-abi` to emit the NVCC-compatible host registration ABI + (`__cudaRegisterLinkedBinary`). + #### AIX Support #### NetBSD Support diff --git a/clang/include/clang/Basic/LangOptions.def b/clang/include/clang/Basic/LangOptions.def index 0519514327355..e8ac435b451b0 100644 --- a/clang/include/clang/Basic/LangOptions.def +++ b/clang/include/clang/Basic/LangOptions.def @@ -264,6 +264,7 @@ LANGOPT(GPUDeferDiag, 1, 0, NotCompatible, "defer host/device related diagnostic LANGOPT(GPUExcludeWrongSideOverloads, 1, 0, NotCompatible, "always exclude wrong side overloads in overloading resolution for CUDA/HIP") LANGOPT(OffloadingNewDriver, 1, 0, NotCompatible, "use the new driver for generating offloading code.") LANGOPT(OffloadViaLLVM, 1, 0, NotCompatible, "target LLVM/Offload as portable offloading runtime.") +LANGOPT(CUDANVCCABI, 1, 0, NotCompatible, "emit the NVCC-compatible host registration ABI for CUDA") LANGOPT(SYCLIsDevice , 1, 0, NotCompatible, "Generate code for SYCL device") LANGOPT(SYCLIsHost , 1, 0, NotCompatible, "SYCL host compilation") diff --git a/clang/include/clang/Options/Options.td b/clang/include/clang/Options/Options.td index e859be134096a..cf66ee3e52f2d 100644 --- a/clang/include/clang/Options/Options.td +++ b/clang/include/clang/Options/Options.td @@ -1277,6 +1277,10 @@ def offload_new_driver : Flag<["--"], "offload-new-driver">, def no_offload_new_driver : Flag<["--"], "no-offload-new-driver">, Visibility<[ClangOption, CC1Option]>, Group<f_Group>, HelpText<"Don't Use the new driver for offloading compilation.">; +def cuda_emit_nvcc_abi : Flag<["--"], "cuda-emit-nvcc-abi">, + Visibility<[ClangOption, CC1Option]>, Group<f_Group>, + MarshallingInfoFlag<LangOpts<"CUDANVCCABI">>, + HelpText<"Emit the NVCC-compatible host registration ABI for CUDA.">; def offload_device_only : Flag<["--"], "offload-device-only">, Visibility<[ClangOption, FlangOption]>, diff --git a/clang/lib/CodeGen/CGCUDANV.cpp b/clang/lib/CodeGen/CGCUDANV.cpp index 416ed935c1b30..1e688d29d15a5 100644 --- a/clang/lib/CodeGen/CGCUDANV.cpp +++ b/clang/lib/CodeGen/CGCUDANV.cpp @@ -1479,8 +1479,9 @@ llvm::Function *CGNVCUDARuntime::finalizeModule() { } return nullptr; } - if (CGM.getLangOpts().OffloadViaLLVM || - (CGM.getLangOpts().OffloadingNewDriver && RelocatableDeviceCode)) + if (!CGM.getLangOpts().CUDANVCCABI && + (CGM.getLangOpts().OffloadViaLLVM || + (CGM.getLangOpts().OffloadingNewDriver && RelocatableDeviceCode))) createOffloadingEntries(); else return makeModuleCtorFunction(); diff --git a/clang/lib/Driver/Driver.cpp b/clang/lib/Driver/Driver.cpp index 0bdc4d7e13082..3b5d0c0aad981 100644 --- a/clang/lib/Driver/Driver.cpp +++ b/clang/lib/Driver/Driver.cpp @@ -5213,9 +5213,10 @@ Driver::BuildOffloadingActions(Compilation &C, llvm::opt::DerivedArgList &Args, OffloadAction::DeviceDependences DDep; if (C.isOffloadingHostKind(Action::OFK_Cuda) && - !Args.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc, false)) { - // If we are not in RDC-mode we just emit the final CUDA fatbinary for - // each translation unit without requiring any linking. + (!Args.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc, false) || + Args.hasArg(options::OPT_cuda_emit_nvcc_abi))) { + // If we are not in RDC-mode or are targeting the NVCC ABI we just emit the + // final CUDA fatbinary for each translation unit without any linking. Action *FatbinAction = C.MakeAction<LinkJobAction>(OffloadActions, types::TY_CUDA_FATBIN); DDep.add(*FatbinAction, *C.getSingleOffloadToolChain<Action::OFK_Cuda>(), diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 41ff5c0cfe577..85fe99dbf8b69 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -8290,7 +8290,8 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, CmdArgs.push_back("-fcuda-include-gpubinary"); CmdArgs.push_back(CudaDeviceInput->getFilename()); } else if (!HostOffloadingInputs.empty()) { - if ((IsCuda || IsHIP) && !IsRDCMode) { + if ((IsCuda || IsHIP) && + (!IsRDCMode || Args.hasArg(options::OPT_cuda_emit_nvcc_abi))) { assert(HostOffloadingInputs.size() == 1 && "Only one input expected"); CmdArgs.push_back("-fcuda-include-gpubinary"); CmdArgs.push_back(HostOffloadingInputs.front().getFilename()); @@ -8305,6 +8306,8 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, if (Args.hasFlag(options::OPT_fcuda_short_ptr, options::OPT_fno_cuda_short_ptr, false)) CmdArgs.push_back("-fcuda-short-ptr"); + if (Args.hasArg(options::OPT_cuda_emit_nvcc_abi)) + CmdArgs.push_back("--cuda-emit-nvcc-abi"); } if (IsCuda || IsHIP) { diff --git a/clang/test/CodeGenCUDA/cuda-emit-nvcc-abi.cu b/clang/test/CodeGenCUDA/cuda-emit-nvcc-abi.cu new file mode 100644 index 0000000000000..cde047c8b67e7 --- /dev/null +++ b/clang/test/CodeGenCUDA/cuda-emit-nvcc-abi.cu @@ -0,0 +1,19 @@ +// RUN: echo "" > %t.fatbin +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -target-sdk-version=12.0 \ +// RUN: -fgpu-rdc --cuda-emit-nvcc-abi -emit-llvm -fcuda-include-gpubinary \ +// RUN: %t.fatbin -o - -x cuda %s \ +// RUN: | FileCheck --check-prefix=NVCC %s + +#include "Inputs/cuda.h" + +__global__ void kernel() {} +__device__ int var = 0; + +// NVCC: @[[FATBIN:[0-9]+]] = private constant {{.*}}, section "__nv_relfatbin" +// NVCC: @__cuda_fatbin_wrapper = internal constant { i32, i32, ptr, ptr } { i32 1180844977, i32 1, ptr @[[FATBIN]], ptr null }, section ".nvFatBinSegment" +// NVCC: @__fatbinwrap__nv_[[ID:[0-9a-f]+]] = alias { i32, i32, ptr, ptr }, ptr @__cuda_fatbin_wrapper +// NVCC: define internal void @__cuda_register_globals(ptr %{{.*}}) +// NVCC: define internal void @__cuda_module_ctor() +// NVCC: call void @__cudaRegisterLinkedBinary__nv_[[ID]](ptr @__cuda_register_globals, ptr @__cuda_fatbin_wrapper, ptr @{{.*}}, ptr @dummy) +// NVCC-NOT: @.offloading.entry +// NVCC-NOT: __tgt_offload_entry diff --git a/clang/test/Driver/cuda-openmp-driver.cu b/clang/test/Driver/cuda-openmp-driver.cu index a1ca987eb9d22..47e83ec426181 100644 --- a/clang/test/Driver/cuda-openmp-driver.cu +++ b/clang/test/Driver/cuda-openmp-driver.cu @@ -38,3 +38,12 @@ // GPU-BINARY: fatbinary{{.*}}"--create" "{{.*}}.fatbin" // GPU-BINARY: -cc1{{.*}}-fcuda-include-gpubinary" "{{.*}}.fatbin" + +// RUN: %clang -### -target x86_64-linux-gnu -nocudalib -fgpu-rdc --cuda-emit-nvcc-abi \ +// RUN: --offload-arch=sm_35 --offload-arch=sm_70 --cuda-path=%S/Inputs/CUDA_111/usr/local/cuda %s 2>&1 \ +// RUN: | FileCheck -check-prefix NVCC-ABI \ +// RUN: --implicit-check-not=clang-offload-packager \ +// RUN: --implicit-check-not=fembed-offload-object %s + +// NVCC-ABI: fatbinary{{.*}}"--create" "{{.*}}.fatbin" +// NVCC-ABI: -cc1{{.*}}"-fcuda-include-gpubinary" "{{.*}}.fatbin"{{.*}}"--cuda-emit-nvcc-abi" _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
