Rename -mlink-cuda-bitcode to -mlink-builtin-bitcode

[arsenm]

Summary

The same semantics work for OpenCL, and probably any offload
language. Keep the old name around as an alias.

[arsenm]

Forgot to commit part

[arsenm]

r340193

include/clang/Driver/CC1Options.td

 def mregparm : Separate<["-"], "mregparm">,
   HelpText<"Limit the number of registers available for integer arguments">;
 def munwind_tables : Flag<["-"], "munwind-tables">,
   HelpText<"Generate unwinding tables for all functions">;
 def mconstructor_aliases : Flag<["-"], "mconstructor-aliases">,
   HelpText<"Emit complete constructors and destructors as aliases when possible">;
 def mlink_bitcode_file : Separate<["-"], "mlink-bitcode-file">,
   HelpText<"Link the given bitcode file before performing optimizations.">;
-def mlink_cuda_bitcode : Separate<["-"], "mlink-cuda-bitcode">,
+
+def mlink_builtin_bitcode : Separate<["-"], "mlink-builtin-bitcode">,
   HelpText<"Link and internalize needed symbols from the given bitcode file "
            "before performing optimizations.">;
+def mlink_cuda_bitcode : Separate<["-"], "mlink-cuda-bitcode">,
+  Alias<mlink_builtin_bitcode>;
+
 def vectorize_loops : Flag<["-"], "vectorize-loops">,
   HelpText<"Run the Loop vectorization passes">;
 def vectorize_slp : Flag<["-"], "vectorize-slp">,
   HelpText<"Run the SLP vectorization passes">;
 def dependent_lib : Joined<["--"], "dependent-lib=">,
   HelpText<"Add dependent library">;
 def linker_option : Joined<["--"], "linker-option=">,
   HelpText<"Add linker option">;

lib/Driver/ToolChains/Cuda.cpp

if (LibDeviceFile.empty()) {
     if (DeviceOffloadingKind == Action::OFK_OpenMP &&
         DriverArgs.hasArg(options::OPT_S))
       return;
 
     getDriver().Diag(diag::err_drv_no_cuda_libdevice) << GpuArch;
     return;
   }
 
-  CC1Args.push_back("-mlink-cuda-bitcode");
+  CC1Args.push_back("-mlink-builtin-bitcode");
   CC1Args.push_back(DriverArgs.MakeArgString(LibDeviceFile));
 
   // Libdevice in CUDA-7.0 requires PTX version that's more recent than LLVM
   // defaults to. Use PTX4.2 by default, which is the PTX version that came with
   // CUDA-7.0.
   const char *PtxFeature = "+ptx42";
   if (CudaInstallation.version() >= CudaVersion::CUDA_91) {
     // CUDA-9.1 uses new instructions that are only available in PTX6.1+
if (DeviceOffloadingKind == Action::OFK_OpenMP) {
 
     std::string LibOmpTargetName =
       "libomptarget-nvptx-" + GpuArch.str() + ".bc";
     bool FoundBCLibrary = false;
     for (StringRef LibraryPath : LibraryPaths) {
       SmallString<128> LibOmpTargetFile(LibraryPath);
       llvm::sys::path::append(LibOmpTargetFile, LibOmpTargetName);
       if (llvm::sys::fs::exists(LibOmpTargetFile)) {
-        CC1Args.push_back("-mlink-cuda-bitcode");
+        CC1Args.push_back("-mlink-builtin-bitcode");
         CC1Args.push_back(DriverArgs.MakeArgString(LibOmpTargetFile));
         FoundBCLibrary = true;
         break;
       }
     }
     if (!FoundBCLibrary)
       getDriver().Diag(diag::warn_drv_omp_offload_target_missingbcruntime)
           << LibOmpTargetName;

lib/Frontend/CompilerInvocation.cpp

if (A->getOption().getID() == OPT_compress_debug_sections) {
                      .Default(llvm::DebugCompressionType::None);
       Opts.setCompressDebugSections(DCT);
     }
   }
 
   Opts.RelaxELFRelocations = Args.hasArg(OPT_mrelax_relocations);
   Opts.DebugCompilationDir = Args.getLastArgValue(OPT_fdebug_compilation_dir);
   for (auto *A :
-       Args.filtered(OPT_mlink_bitcode_file, OPT_mlink_cuda_bitcode)) {
+       Args.filtered(OPT_mlink_bitcode_file, OPT_mlink_builtin_bitcode)) {
     CodeGenOptions::BitcodeFileToLink F;
     F.Filename = A->getValue();
-    if (A->getOption().matches(OPT_mlink_cuda_bitcode)) {
+    if (A->getOption().matches(OPT_mlink_builtin_bitcode)) {
       F.LinkFlags = llvm::Linker::Flags::LinkOnlyNeeded;
       // When linking CUDA bitcode, propagate function attributes so that
       // e.g. libdevice gets fast-math attrs if we're building with fast-math.
       F.PropagateAttrs = true;
       F.Internalize = true;
     }
     Opts.LinkBitcodeFiles.push_back(F);
   }

test/CodeGenCUDA/link-device-bitcode.cu

 // Test for linking with CUDA's libdevice as outlined in
 // http://llvm.org/docs/NVPTXUsage.html#linking-with-libdevice
 //
 // REQUIRES: nvptx-registered-target
 //
 // Prepare bitcode file to link with
 // RUN: %clang_cc1 -triple nvptx-unknown-cuda -emit-llvm-bc \
 // RUN:    -disable-llvm-passes -o %t.bc %S/Inputs/device-code.ll
 // RUN: %clang_cc1 -triple nvptx-unknown-cuda -emit-llvm-bc \
 // RUN:    -disable-llvm-passes -o %t-2.bc %S/Inputs/device-code-2.ll
 //
 // Make sure function in device-code gets linked in and internalized.
 // RUN: %clang_cc1 -triple nvptx-unknown-cuda -fcuda-is-device \
+// RUN:    -mlink-builtin-bitcode %t.bc  -emit-llvm \
+// RUN:    -disable-llvm-passes -o - %s \
+// RUN:    | FileCheck %s -check-prefix CHECK-IR
+
+// Make sure legacy flag name works
+// RUN: %clang_cc1 -triple nvptx-unknown-cuda -fcuda-is-device \
 // RUN:    -mlink-cuda-bitcode %t.bc  -emit-llvm \
 // RUN:    -disable-llvm-passes -o - %s \
 // RUN:    | FileCheck %s -check-prefix CHECK-IR
 //
 // Make sure we can link two bitcode files.
 // RUN: %clang_cc1 -triple nvptx-unknown-cuda -fcuda-is-device \
-// RUN:    -mlink-cuda-bitcode %t.bc -mlink-cuda-bitcode %t-2.bc \
+// RUN:    -mlink-builtin-bitcode %t.bc -mlink-builtin-bitcode %t-2.bc \
 // RUN:    -emit-llvm -disable-llvm-passes -o - %s \
 // RUN:    | FileCheck %s -check-prefix CHECK-IR -check-prefix CHECK-IR-2
 //
 // Make sure function in device-code gets linked but is not internalized
 // without -fcuda-uses-libdevice
 // RUN: %clang_cc1 -triple nvptx-unknown-cuda -fcuda-is-device \
 // RUN:    -mlink-bitcode-file %t.bc -emit-llvm \
 // RUN:    -disable-llvm-passes -o - %s \
 // RUN:    | FileCheck %s -check-prefix CHECK-IR-NLD
 //
 // Make sure NVVMReflect pass is enabled in NVPTX back-end.
 // RUN: %clang_cc1 -triple nvptx-unknown-cuda -fcuda-is-device \
-// RUN:    -mlink-cuda-bitcode %t.bc -S -o /dev/null %s \
+// RUN:    -mlink-builtin-bitcode %t.bc -S -o /dev/null %s \
 // RUN:    -mllvm -debug-pass=Structure 2>&1 \
 // RUN:    | FileCheck %s -check-prefix CHECK-REFLECT
 
 #include "Inputs/cuda.h"
 
 __device__ float device_mul_or_add(float a, float b);
 extern "C" __device__ double __nv_sin(double x);
 extern "C" __device__ double __nv_exp(double x);

test/CodeGenCUDA/propagate-metadata.cu

 // Check that when we link a bitcode module into a file using
-// -mlink-cuda-bitcode, we apply the same attributes to the functions in that
+// -mlink-builtin-bitcode, we apply the same attributes to the functions in that
 // bitcode module as we apply to functions we generate.
 //
 // In particular, we check that ftz and unsafe-math are propagated into the
 // bitcode library as appropriate.
 //
 // In addition, we set -ftrapping-math on the bitcode library, but then set
 // -fno-trapping-math on the main compilations, and ensure that the latter flag
 // overrides the flag on the bitcode library.
 
 // Build the bitcode library.  This is not built in CUDA mode, otherwise it
 // might have incompatible attributes.  This mirrors how libdevice is built.
 // RUN: %clang_cc1 -x c++ -emit-llvm-bc -ftrapping-math -DLIB \
 // RUN:   %s -o %t.bc -triple nvptx-unknown-unknown
 
-// RUN: %clang_cc1 -x cuda %s -emit-llvm -mlink-cuda-bitcode %t.bc -o - \
+// RUN: %clang_cc1 -x cuda %s -emit-llvm -mlink-builtin-bitcode %t.bc -o - \
 // RUN:   -fno-trapping-math -fcuda-is-device -triple nvptx-unknown-unknown \
 // RUN: | FileCheck %s --check-prefix=CHECK --check-prefix=NOFTZ --check-prefix=NOFAST
 
-// RUN: %clang_cc1 -x cuda %s -emit-llvm -mlink-cuda-bitcode %t.bc \
+// RUN: %clang_cc1 -x cuda %s -emit-llvm -mlink-builtin-bitcode %t.bc \
 // RUN:   -fno-trapping-math -fcuda-flush-denormals-to-zero -o - \
 // RUN:   -fcuda-is-device -triple nvptx-unknown-unknown \
 // RUN: | FileCheck %s --check-prefix=CHECK --check-prefix=FTZ \
 // RUN:   --check-prefix=NOFAST
 
-// RUN: %clang_cc1 -x cuda %s -emit-llvm -mlink-cuda-bitcode %t.bc \
+// RUN: %clang_cc1 -x cuda %s -emit-llvm -mlink-builtin-bitcode %t.bc \
 // RUN:   -fno-trapping-math -fcuda-flush-denormals-to-zero -o - \
 // RUN:   -fcuda-is-device -menable-unsafe-fp-math -triple nvptx-unknown-unknown \
 // RUN: | FileCheck %s --check-prefix=CHECK --check-prefix=FAST
 
 // Wrap everything in extern "C" so we don't have to worry about name mangling
 // in the IR.
 extern "C" {
 #ifdef LIB

test/Driver/cuda-detect.cu

 // CHECK: Found CUDA installation: {{.*}}/Inputs/CUDA/usr/local/cuda
 // NO-LIBDEVICE: Found CUDA installation: {{.*}}/Inputs/CUDA-nolibdevice/usr/local/cuda
 // NOCUDA-NOT: Found CUDA installation:
 
 // MISSINGLIBDEVICE: error: cannot find libdevice for sm_20.
 
 // COMMON: "-triple" "nvptx-nvidia-cuda"
 // COMMON-SAME: "-fcuda-is-device"
-// LIBDEVICE-SAME: "-mlink-cuda-bitcode"
-// NOLIBDEVICE-NOT: "-mlink-cuda-bitcode"
+// LIBDEVICE-SAME: "-mlink-builtin-bitcode"
+// NOLIBDEVICE-NOT: "-mlink-builtin-bitcode"
 // LIBDEVICE20-SAME: libdevice.compute_20.10.bc
 // LIBDEVICE30-SAME: libdevice.compute_30.10.bc
 // LIBDEVICE35-SAME: libdevice.compute_35.10.bc
 // LIBDEVICE50-SAME: libdevice.compute_50.10.bc
 // NOLIBDEVICE-NOT: libdevice.compute_{{.*}}.bc
 // LIBDEVICE-SAME: "-target-feature" "+ptx42"
 // NOLIBDEVICE-NOT: "-target-feature" "+ptx42"
 // CUDAINC-SAME: "-internal-isystem" "{{.*}}/Inputs/CUDA{{[_0-9]+}}/usr/local/cuda/include"

test/Driver/openmp-offload-gpu.c

 
 /// Check that the runtime bitcode library is part of the compile line. Create a bogus
 /// bitcode library and add it to the LIBRARY_PATH.
 // RUN:   env LIBRARY_PATH=%S/Inputs/libomptarget %clang -### -fopenmp=libomp -fopenmp-targets=nvptx64-nvidia-cuda \
 // RUN:   -Xopenmp-target -march=sm_20 --cuda-path=%S/Inputs/CUDA_80/usr/local/cuda \
 // RUN:   -fopenmp-relocatable-target -save-temps -no-canonical-prefixes %s 2>&1 \
 // RUN:   | FileCheck -check-prefix=CHK-BCLIB %s
 
-// CHK-BCLIB: clang{{.*}}-triple{{.*}}nvptx64-nvidia-cuda{{.*}}-mlink-cuda-bitcode{{.*}}libomptarget-nvptx-sm_20.bc
+// CHK-BCLIB: clang{{.*}}-triple{{.*}}nvptx64-nvidia-cuda{{.*}}-mlink-builtin-bitcode{{.*}}libomptarget-nvptx-sm_20.bc
 // CHK-BCLIB-NOT: {{error:|warning:}}
 
 /// ###########################################################################
 
 /// Check that the warning is thrown when the libomptarget bitcode library is not found.
 /// Libomptarget requires sm_35 or newer so an sm_20 bitcode library should never exist.
 // RUN:   %clang -### -fopenmp=libomp -fopenmp-targets=nvptx64-nvidia-cuda \
 // RUN:   -Xopenmp-target -march=sm_20 --cuda-path=%S/Inputs/CUDA_80/usr/local/cuda \