[clang] Pass target-features to clang-linker-wrapper

Author: jzc

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<StringRef> 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

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)

clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp

@@ -641,7 +641,8 @@ Expected<StringRef> writeOffloadFile(const OffloadFile &File) {
 
 // Compile the module to an object file using the appropriate target machine for
 // the host triple.
-Expected<StringRef> compileModule(Module &M, OffloadKind Kind) {
+Expected<StringRef> 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<StringRef> 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<TargetMachine> TM(
       T->createTargetMachine(M.getTargetTriple(), CPU, Features, Options,
                              Reloc::PIC_, M.getCodeModel()));
@@ -746,7 +747,7 @@ wrapDeviceImages(ArrayRef<std::unique_ptr<MemoryBuffer>> Buffers,
     WriteBitcodeToFile(M, OS);
   }
 
-  auto FileOrErr = compileModule(M, Kind);
+  auto FileOrErr = compileModule(M, Args, Kind);
   if (!FileOrErr)
     return FileOrErr.takeError();
   return *FileOrErr;

clang/tools/clang-linker-wrapper/LinkerWrapperOpts.td

@@ -20,6 +20,15 @@ def host_triple_EQ : Joined<["--"], "host-triple=">,
                      Flags<[WrapperOnlyOption]>,
                      MetaVarName<"<triple>">,
                      HelpText<"Triple to use for the host compilation">;
+def host_cpu_EQ : Joined<["--"], "host-cpu=">,
+                  Flags<[WrapperOnlyOption]>,
+                  MetaVarName<"<cpu>">,
+                  HelpText<"CPU name to use for the host compilation">;
+def host_features_EQ
+    : Joined<["--"], "host-features=">,
+      Flags<[WrapperOnlyOption]>,
+      MetaVarName<"<features>">,
+      HelpText<"Target features to use for the host compilation">;
 def device_linker_args_EQ : Joined<["--"], "device-linker=">,
   Flags<[WrapperOnlyOption]>, MetaVarName<"<value> or <triple>=<value>">,
   HelpText<"Arguments to pass to the device linker invocation">;