diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 4fe9cc4990eb7..244d28084117a 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -9171,6 +9171,32 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, CmdArgs.push_back( Args.MakeArgString("--host-triple=" + getToolChain().getTripleString())); + const Driver &D = getToolChain().getDriver(); + const llvm::Triple TheTriple = getToolChain().getTriple(); + std::string HostCPU = tools::getCPUName(D, Args, TheTriple); + if (!HostCPU.empty()) + CmdArgs.push_back(Args.MakeArgString("--host-cpu=" + HostCPU)); + + ArgStringList HostFeatureArgs; + getTargetFeatures(D, TheTriple, Args, HostFeatureArgs, /*ForAS=*/false, + /*IsAux=*/false); + + // getTargetFeatures() returns arguments in the form ["-target-feature", "+f", + // "-target-feature", "+d"] but clang-linker-wrapper expects + // --host-features=+f,+d format, so we need to extract only the feature + // strings ("+f", "+d") and skip the "-target-feature" flags. + std::vector HostFeatures; + // Features always come in pairs: ["-target-feature", "value"], so take every + // odd element + for (size_t i = 1; i < HostFeatureArgs.size(); i += 2) { + HostFeatures.push_back(HostFeatureArgs[i]); + } + + if (!HostFeatures.empty()) { + std::string HostFeaturesStr = llvm::join(HostFeatures, ","); + CmdArgs.push_back(Args.MakeArgString("--host-features=" + HostFeaturesStr)); + } + // CMake hack, suppress passing verbose arguments for the special-case HIP // non-RDC mode compilation. This confuses default CMake implicit linker // argument parsing when the language is set to HIP and the system linker is diff --git a/clang/test/Driver/linker-wrapper-host-features.c b/clang/test/Driver/linker-wrapper-host-features.c new file mode 100644 index 0000000000000..f6a004e2e8d84 --- /dev/null +++ b/clang/test/Driver/linker-wrapper-host-features.c @@ -0,0 +1,77 @@ +// Test that RISC-V host target features are correctly passed to the linker wrapper compilation +// and applied to wrapper object generation. +// UNSUPPORTED: system-windows +// REQUIRES: riscv-registered-target +// REQUIRES: nvptx-registered-target + +// Simple program that requires OpenMP offloading +int main() { +#pragma omp target + { + // Device code + } + return 0; +} + +// Verify that the driver invokes clang-linker-wrapper with correct host features. + +// Test lp64d (double-float) ABI +// RUN: %clang -target riscv64-linux-gnu -mabi=lp64d \ +// RUN: -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda --offload-arch=sm_80 \ +// RUN: --libomptarget-nvptx-bc-path=%S/Inputs/libomptarget/libomptarget-nvptx-test.bc \ +// RUN: %s -o %t.lp64d -### 2>&1 | FileCheck %s --check-prefix=LP64D-DRIVER + +// Check that the driver calls clang-linker-wrapper with correct host features for lp64d +// LP64D-DRIVER: clang-linker-wrapper{{.*}}--host-triple=riscv64{{.*}}--host-features={{.*}}+f{{.*}}+d{{.*}} + +// Test lp64f (single-float) ABI +// RUN: %clang -target riscv64-linux-gnu -mabi=lp64f \ +// RUN: -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda --offload-arch=sm_80 \ +// RUN: --libomptarget-nvptx-bc-path=%S/Inputs/libomptarget/libomptarget-nvptx-test.bc \ +// RUN: %s -o %t.lp64f -### 2>&1 | FileCheck %s --check-prefix=LP64F-DRIVER + +// LP64F-DRIVER: clang-linker-wrapper{{.*}}--host-triple=riscv64{{.*}}--host-features={{.*}}+f{{.*}}-d{{.*}} + +// Test lp64 (soft-float) ABI +// RUN: %clang -target riscv64-linux-gnu -mabi=lp64 \ +// RUN: -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda --offload-arch=sm_80 \ +// RUN: --libomptarget-nvptx-bc-path=%S/Inputs/libomptarget/libomptarget-nvptx-test.bc \ +// RUN: %s -o %t.lp64 -### 2>&1 | FileCheck %s --check-prefix=LP64-DRIVER + +// LP64-DRIVER: clang-linker-wrapper{{.*}}--host-triple=riscv64{{.*}}--host-features={{.*}}-f{{.*}}-d{{.*}} + +// Verify that clang-linker-wrapper applies RISC-V host features correctly when creating wrapper objects. +// We do this by checking the ELF ABI flags in the generated wrapper object files. + +// Create test objects for linker-wrapper testing +// RUN: %clang -cc1 %s -triple nvptx64-nvidia-cuda -emit-llvm-bc -o %t.device.bc +// RUN: llvm-offload-binary -o %t.openmp.out \ +// RUN: --image=file=%t.device.bc,kind=openmp,triple=nvptx64-nvidia-cuda,arch=sm_70 +// RUN: %clang -cc1 %s -triple riscv64-unknown-linux-gnu -emit-obj -o %t.riscv.host.o -fembed-offload-object=%t.openmp.out + +// Test lp64 (soft-float) ABI - should generate ELF flags 0x0 +// RUN: rm -rf %t.tmpdir1 && mkdir %t.tmpdir1 +// RUN: cd %t.tmpdir1 && clang-linker-wrapper --host-triple=riscv64-unknown-linux-gnu \ +// RUN: --host-features=-f,-d --linker-path=/usr/bin/ld %t.riscv.host.o -o %t.lp64.out --dry-run --save-temps 2>&1 +// RUN: cd %t.tmpdir1 && find . -name "*.wrapper*.o" -exec llvm-readobj --file-headers "{}" ";" \ +// RUN: | FileCheck %s --check-prefix=SOFT-FLOAT-OBJ + +// SOFT-FLOAT-OBJ: Flags [ (0x0) + +// Test lp64f (single-float) ABI - should generate ELF flags 0x2 +// RUN: rm -rf %t.tmpdir2 && mkdir %t.tmpdir2 +// RUN: cd %t.tmpdir2 && clang-linker-wrapper --host-triple=riscv64-unknown-linux-gnu \ +// RUN: --host-features=+f,-d --linker-path=/usr/bin/ld %t.riscv.host.o -o %t.lp64f.out --dry-run --save-temps 2>&1 +// RUN: cd %t.tmpdir2 && find . -name "*.wrapper*.o" -exec llvm-readobj --file-headers "{}" ";" \ +// RUN: | FileCheck %s --check-prefix=SINGLE-FLOAT-OBJ + +// SINGLE-FLOAT-OBJ: Flags [ (0x2) + +// Test lp64d (double-float) ABI - should generate ELF flags 0x4 +// RUN: rm -rf %t.tmpdir3 && mkdir %t.tmpdir3 +// RUN: cd %t.tmpdir3 && clang-linker-wrapper --host-triple=riscv64-unknown-linux-gnu \ +// RUN: --host-features=+f,+d --linker-path=/usr/bin/ld %t.riscv.host.o -o %t.lp64d.out --dry-run --save-temps 2>&1 +// RUN: cd %t.tmpdir3 && find . -name "*.wrapper*.o" -exec llvm-readobj --file-headers "{}" ";" \ +// RUN: | FileCheck %s --check-prefix=DOUBLE-FLOAT-OBJ + +// DOUBLE-FLOAT-OBJ: Flags [ (0x4) diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index bfeca17d2147e..9317c9b666ab1 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -641,7 +641,8 @@ Expected writeOffloadFile(const OffloadFile &File) { // Compile the module to an object file using the appropriate target machine for // the host triple. -Expected compileModule(Module &M, OffloadKind Kind) { +Expected compileModule(Module &M, const ArgList &Args, + OffloadKind Kind) { llvm::TimeTraceScope TimeScope("Compile module"); std::string Msg; const Target *T = TargetRegistry::lookupTarget(M.getTargetTriple(), Msg); @@ -650,8 +651,8 @@ Expected compileModule(Module &M, OffloadKind Kind) { auto Options = codegen::InitTargetOptionsFromCodeGenFlags(M.getTargetTriple()); - StringRef CPU = ""; - StringRef Features = ""; + StringRef CPU = Args.getLastArgValue(OPT_host_cpu_EQ, ""); + StringRef Features = Args.getLastArgValue(OPT_host_features_EQ, ""); std::unique_ptr TM( T->createTargetMachine(M.getTargetTriple(), CPU, Features, Options, Reloc::PIC_, M.getCodeModel())); @@ -746,7 +747,7 @@ wrapDeviceImages(ArrayRef> Buffers, WriteBitcodeToFile(M, OS); } - auto FileOrErr = compileModule(M, Kind); + auto FileOrErr = compileModule(M, Args, Kind); if (!FileOrErr) return FileOrErr.takeError(); return *FileOrErr; diff --git a/clang/tools/clang-linker-wrapper/LinkerWrapperOpts.td b/clang/tools/clang-linker-wrapper/LinkerWrapperOpts.td index 87f911c749bf6..23fd7f8ed9e30 100644 --- a/clang/tools/clang-linker-wrapper/LinkerWrapperOpts.td +++ b/clang/tools/clang-linker-wrapper/LinkerWrapperOpts.td @@ -20,6 +20,15 @@ def host_triple_EQ : Joined<["--"], "host-triple=">, Flags<[WrapperOnlyOption]>, MetaVarName<"">, HelpText<"Triple to use for the host compilation">; +def host_cpu_EQ : Joined<["--"], "host-cpu=">, + Flags<[WrapperOnlyOption]>, + MetaVarName<"">, + HelpText<"CPU name to use for the host compilation">; +def host_features_EQ + : Joined<["--"], "host-features=">, + Flags<[WrapperOnlyOption]>, + MetaVarName<"">, + HelpText<"Target features to use for the host compilation">; def device_linker_args_EQ : Joined<["--"], "device-linker=">, Flags<[WrapperOnlyOption]>, MetaVarName<" or =">, HelpText<"Arguments to pass to the device linker invocation">;