Merge remote-tracking branch 'origin/master' into runtime_compile_v5

.circleci/config.yml

@@ -4,6 +4,7 @@ jobs:
     docker:
       - image: gcc
     environment:
+      - LLVM_VERSION: 5.0.0
       - HOST_LDC_VERSION: 1.3.0
     steps:
       - checkout
@@ -24,15 +25,22 @@ jobs:
             ninja --version
             gdb --version
             python -c "import lit; lit.main();" --version | head -n 1
-      - run:
-          name: Install LLVM nightly
-          command: |
-            add-apt-repository -y 'deb http://apt.llvm.org/unstable/ llvm-toolchain-5.0 main'
-            wget -O - http://apt.llvm.org/llvm-snapshot.gpg.key | apt-key add -
-            apt update
-            apt install -y llvm-5.0 llvm-5.0-dev libclang-common-5.0-dev libfuzzer-5.0-dev
       - restore_cache:
-          key: host-ldc-{{ .Environment.HOST_LDC_VERSION }}
+          keys:
+            - llvm-5.0.0
+            - host-ldc-{{ .Environment.HOST_LDC_VERSION }}
+      - run:
+          name: Install LLVM 5.0.0
+          command: |
+            if [[ ! -d llvm-$LLVM_VERSION ]]; then
+              wget -O llvm-$LLVM_VERSION.tar.xz http://releases.llvm.org/$LLVM_VERSION/clang+llvm-$LLVM_VERSION-linux-x86_64-ubuntu16.04.tar.xz
+              mkdir llvm-$LLVM_VERSION
+              tar -xpf llvm-$LLVM_VERSION.tar.xz --strip 1 -C llvm-$LLVM_VERSION
+            fi
+      - save_cache:
+          key: llvm-5.0.0
+          paths:
+            - llvm-5.0.0
       - run:
           name: Install LDC host compiler
           command: |
@@ -50,7 +58,7 @@ jobs:
             export HOST_LDMD=$PWD/ldc2-$HOST_LDC_VERSION-linux-x86_64/bin/ldmd2
             mkdir bootstrap
             cd bootstrap
-            cmake -G Ninja -DBUILD_SHARED_LIBS=OFF -DD_COMPILER=$HOST_LDMD ..
+            cmake -G Ninja -DLLVM_ROOT_DIR=$PWD/../llvm-$LLVM_VERSION -DBUILD_SHARED_LIBS=OFF -DD_COMPILER=$HOST_LDMD ..
             ninja -j3
             bin/ldc2 -version
             cd ..
@@ -60,7 +68,7 @@ jobs:
             export HOST_LDMD=$PWD/bootstrap/bin/ldmd2
             mkdir build
             cd build
-            cmake -G Ninja -DLDC_INSTALL_LTOPLUGIN=ON -DLDC_INSTALL_LLVM_RUNTIME_LIBS=ON -DD_COMPILER=$HOST_LDMD ..
+            cmake -G Ninja -DLLVM_ROOT_DIR=$PWD/../llvm-$LLVM_VERSION -DLDC_INSTALL_LTOPLUGIN=ON -DLDC_INSTALL_LLVM_RUNTIME_LIBS=ON -DD_COMPILER=$HOST_LDMD ..
             ninja -j3 all all-test-runners
             bin/ldc2 -version
             cd ..

.travis.yml

@@ -4,6 +4,9 @@ sudo: false
 
 matrix:
   include:
+    - os: linux
+      d: ldc
+      env: LLVM_VERSION=5.0.0 OPTS="-DLIB_SUFFIX=64"
     - os: linux
       d: ldc
       env: LLVM_VERSION=4.0.1 OPTS="-DLIB_SUFFIX=64"
@@ -18,7 +21,7 @@ matrix:
       env: LLVM_VERSION=3.7.1 OPTS="-DTEST_COVERAGE=ON"
     - os: osx
       d: ldc-beta
-      env: LLVM_VERSION=6.0.0-2 OPTS="-DBUILD_SHARED_LIBS=OFF" LLVM_SPIRV_AVAILABLE=ON
+      env: LLVM_VERSION=6.0.0-4 OPTS="-DBUILD_SHARED_LIBS=OFF" LLVM_SPIRV_AVAILABLE=ON
     - os: osx
       d: ldc
       env: LLVM_VERSION=4.0.0 OPTS="-DBUILD_SHARED_LIBS=ON"
@@ -27,7 +30,8 @@ matrix:
 
 cache:
   directories:
-    - llvm-spirv-6.0.0-2
+    - llvm-spirv-6.0.0-4
+    - llvm-5.0.0
     - llvm-4.0.1
     - llvm-4.0.0
     - llvm-3.9.1
@@ -44,7 +48,9 @@ addons:
 before_install:
   -
     if [ "${TRAVIS_OS_NAME}" = "linux" ]; then
-      if [ "${LLVM_VERSION}" = "4.0.1" ]; then
+      if [ "${LLVM_VERSION}" = "5.0.0" ]; then
+        export LLVM_ARCH="linux-x86_64-ubuntu14.04";
+      elif [ "${LLVM_VERSION}" = "4.0.1" ]; then
         export LLVM_ARCH="x86_64-linux-gnu-debian8";
       else
         export LLVM_ARCH="x86_64-linux-gnu-ubuntu-14.04";
@@ -54,7 +60,7 @@ before_install:
     fi;
     if [ -z "$(ls -A llvm-$LLVM_VERSION)" ]; then
       if [ -n "${LLVM_SPIRV_AVAILABLE}" ]; then
-        wget -O llvm-spirv-$LLVM_VERSION.tar.bz2 https://github.com/thewilsonator/llvm/releases/download/pre-intrinsics/LLVM-6.0.0svn-Darwin.tar.bz2;
+        wget -O llvm-spirv-$LLVM_VERSION.tar.bz2 https://github.com/thewilsonator/llvm/releases/download/pre-intrinsics/LLVM-6.0.0svn-Darwin-2.tar.bz2;
         mkdir llvm-spirv-$LLVM_VERSION;
         tar -xjf llvm-spirv-$LLVM_VERSION.tar.bz2 --strip 1 -C llvm-spirv-$LLVM_VERSION;
       else

CMakeLists.txt

@@ -52,7 +52,7 @@ endfunction()
 #
 
 # Version information
-set(LDC_VERSION "1.4.0") # May be overridden by git hash tag
+set(LDC_VERSION "1.5.0") # May be overridden by git hash tag
 set(DMDFE_MAJOR_VERSION   2)
 set(DMDFE_MINOR_VERSION   0)
 set(DMDFE_PATCH_VERSION   74)
@@ -339,6 +339,7 @@ set(DRV_SRC
     driver/cache.cpp
     driver/cl_options.cpp
     driver/cl_options_sanitizers.cpp
+    driver/cl_options-llvm.cpp
     driver/codegenerator.cpp
     driver/configfile.cpp
     driver/dcomputecodegenerator.cpp
@@ -358,6 +359,7 @@ set(DRV_HDR
     driver/cache_pruning.h
     driver/cl_options.h
     driver/cl_options_sanitizers.h
+    driver/cl_options-llvm.h
     driver/codegenerator.h
     driver/configfile.h
     driver/dcomputecodegenerator.h
@@ -464,27 +466,29 @@ endif()
 #
 # LLD integration (requires LLVM >= 3.9 with LLD headers & libs)
 #
-set(LDC_WITH_LLD OFF)
+if(NOT DEFINED LDC_WITH_LLD)
-if(LDC_LLVM_VER GREATER 308)
+    if(LDC_LLVM_VER GREATER 308)
-    # check for LLD header
+        # check for LLD header
-    unset(LDC_WITH_LLD)
+        if(NOT MSVC)
-    if(NOT MSVC)
+            set(CMAKE_REQUIRED_FLAGS -std=c++11)
-        set(CMAKE_REQUIRED_FLAGS -std=c++11)
+        endif()
-    endif()
+        set(CMAKE_REQUIRED_INCLUDES ${LLVM_INCLUDE_DIRS})
-    set(CMAKE_REQUIRED_INCLUDES ${LLVM_INCLUDE_DIRS})
+        CHECK_INCLUDE_FILE_CXX(lld/Driver/Driver.h LDC_WITH_LLD)
-    CHECK_INCLUDE_FILE_CXX(lld/Driver/Driver.h LDC_WITH_LLD)
+        unset(CMAKE_REQUIRED_FLAGS)
-    unset(CMAKE_REQUIRED_FLAGS)
+        unset(CMAKE_REQUIRED_INCLUDES)
-    unset(CMAKE_REQUIRED_INCLUDES)
+    else()
-    if(LDC_WITH_LLD)
+        set(LDC_WITH_LLD OFF)
-        message(STATUS "Building LDC with LLD support")
-        append("-DLDC_WITH_LLD" LDC_CXXFLAGS)
     endif()
 endif()
+if(LDC_WITH_LLD)
+    message(STATUS "Building LDC with LLD support")
+    append("-DLDC_WITH_LLD" LDC_CXXFLAGS)
+endif()
 
 #
 # Enable building with riscv-llvm, for full RISC-V support.
 #
-option(RISCV_LLVM_DEV, "full RISC-V support with riscv-llvm")
+option(RISCV_LLVM_DEV "full RISC-V support with riscv-llvm")
 mark_as_advanced(RISCV_LLVM_DEV)
 if(RISCV_LLVM_DEV)
     append("-DRISCV_LLVM_DEV" LDC_CXXFLAGS)

README.md

@@ -1,6 +1,7 @@
 LDC – the LLVM-based D Compiler
 ===============================
 
+[![Build Status](https://circleci.com/gh/ldc-developers/ldc/tree/master.svg?style=svg)][6]
 [![Build Status](https://semaphoreci.com/api/v1/ldc-developers/ldc/branches/master/shields_badge.svg)][4]
 [![Build Status](https://travis-ci.org/ldc-developers/ldc.png?branch=master)][1]
 [![Build Status](https://ci.appveyor.com/api/projects/status/2cfhvg79782n4nth/branch/master?svg=true)][5]
@@ -95,3 +96,4 @@ Feedback of any kind is very much appreciated!
 [3]: https://www.bountysource.com/trackers/283332-ldc?utm_source=283332&utm_medium=shield&utm_campaign=TRACKER_BADGE "Bountysource"
 [4]: https://semaphoreci.com/ldc-developers/ldc "Semaphore CI Build Status"
 [5]: https://ci.appveyor.com/project/kinke/ldc/history "AppVeyor CI Build Status"
+[6]: https://circleci.com/gh/ldc-developers/ldc/tree/master "Circle CI Build Status"

ddmd/mtype.d

@@ -2532,7 +2532,10 @@ extern (C++) abstract class Type : RootObject
         }
         else if (ident == Id.__xalignof)
         {
-            e = new IntegerExp(loc, alignsize(), Type.tsize_t);
+            const explicitAlignment = alignment();
+            const naturalAlignment = alignsize();
+            const actualAlignment = (explicitAlignment == STRUCTALIGN_DEFAULT ? naturalAlignment : explicitAlignment);
+            e = new IntegerExp(loc, actualAlignment, Type.tsize_t);
         }
         else if (ident == Id._init)
         {

driver/archiver.cpp

@@ -50,7 +50,7 @@ bool Symtab = true;
 bool Deterministic = true;
 bool Thin = false;
 
-void fail(Twine Error) { outs() << "llvm-ar: " << Error << ".\n"; }
+void fail(Twine Error) { errs() << "llvm-ar: " << Error << ".\n"; }
 
 void fail(std::error_code EC, std::string Context = {}) {
   if (Context.empty())
@@ -79,6 +79,12 @@ int addMember(std::vector<NewArchiveMember> &Members, StringRef FileName,
     fail(std::move(Error), FileName);
     return 1;
   }
+
+#if LDC_LLVM_VER >= 500
+  // Use the basename of the object path for the member name.
+  NMOrErr->MemberName = sys::path::filename(NMOrErr->MemberName);
+#endif
+
   if (Pos == -1)
     Members.push_back(std::move(*NMOrErr));
   else
@@ -197,13 +203,15 @@ int performWriteOperation(object::Archive *OldArchive,
   else
     Kind = getKindFromMember(NewMembers.front());
 
-  const auto Result =
+  auto Result =
       writeArchive(ArchiveName, NewMembers, Symtab, Kind, Deterministic, Thin,
                    std::move(OldArchiveBuf));
 
 #if LDC_LLVM_VER >= 600
   if (Result) {
-    fail("error writing '" + ArchiveName + "': " + Result.message());
+    handleAllErrors(std::move(Result), [](ErrorInfoBase &EIB) {
+      fail("error writing '" + ArchiveName + "': " + EIB.message());
+    });
     return 1;
   }
 #else
@@ -221,7 +229,7 @@ int performWriteOperation() {
   auto Buf = MemoryBuffer::getFile(ArchiveName, -1, false);
   std::error_code EC = Buf.getError();
   if (EC && EC != errc::no_such_file_or_directory) {
-    fail("error opening '" + ArchiveName + "': " + EC.message() + "!");
+    fail("error opening '" + ArchiveName + "': " + EC.message());
     return 1;
   }
 
@@ -230,9 +238,7 @@ int performWriteOperation() {
     object::Archive Archive(Buf.get()->getMemBufferRef(), Err);
     EC = errorToErrorCode(std::move(Err));
     if (EC) {
-      fail(
+      fail(EC, ("error loading '" + ArchiveName + "': " + EC.message()).str());
-          EC,
-          ("error loading '" + ArchiveName + "': " + EC.message() + "!").str());
       return 1;
     }
     return performWriteOperation(&Archive, std::move(Buf.get()));

driver/cache.cpp

@@ -302,14 +302,18 @@ void outputIR2ObjRelevantCmdlineArgs(llvm::raw_ostream &hash_os) {
   // sharing the cache).
   outputOptimizationSettings(hash_os);
   opts::outputSanitizerSettings(hash_os);
-  hash_os << opts::mCPU;
+  hash_os << opts::getCPUStr();
-  for (auto &attr : opts::mAttrs) {
+  hash_os << opts::getFeaturesStr();
-    hash_os << attr;
+  hash_os << opts::floatABI;
-  }
+#if LDC_LLVM_VER >= 309
-  hash_os << opts::mFloatABI;
+  const auto relocModel = opts::getRelocModel();
-  hash_os << opts::mRelocModel;
+  if (relocModel.hasValue())
-  hash_os << opts::mCodeModel;
+    hash_os << relocModel.getValue();
-  hash_os << opts::disableFpElim;
+#else
+  hash_os << opts::getRelocModel();
+#endif
+  hash_os << opts::getCodeModel();
+  hash_os << opts::disableFPElim();
 }
 
 // Output to `hash_os` all environment flags that influence object code output

driver/cl_options-llvm.cpp

@@ -0,0 +1,64 @@
+//===-- cl_options-llvm.cpp -----------------------------------------------===//
+//
+//                         LDC – the LLVM D compiler
+//
+// This file is distributed under the BSD-style LDC license. See the LICENSE
+// file for details.
+//
+//===----------------------------------------------------------------------===//
+
+#include "driver/cl_options-llvm.h"
+
+// Pull in command-line options and helper functions from special LLVM header
+// shared by multiple LLVM tools.
+#include "llvm/CodeGen/CommandFlags.h"
+
+static cl::opt<bool>
+    DisableRedZone("disable-red-zone", cl::ZeroOrMore,
+                   cl::desc("Do not emit code that uses the red zone."));
+
+// Now expose the helper functions (with static linkage) via external wrappers
+// in the opts namespace, including some additional helper functions.
+namespace opts {
+
+std::string getArchStr() { return ::MArch; }
+
+#if LDC_LLVM_VER >= 309
+Optional<Reloc::Model> getRelocModel() { return ::getRelocModel(); }
+#else
+Reloc::Model getRelocModel() { return ::RelocModel; }
+#endif
+
+CodeModel::Model getCodeModel() { return ::CMModel; }
+
+bool disableFPElim() { return ::DisableFPElim; }
+
+bool disableRedZone() { return ::DisableRedZone; }
+
+bool printTargetFeaturesHelp() {
+  if (MCPU == "help")
+    return true;
+  return std::any_of(MAttrs.begin(), MAttrs.end(),
+                     [](const std::string &a) { return a == "help"; });
+}
+
+TargetOptions InitTargetOptionsFromCodeGenFlags() {
+  return ::InitTargetOptionsFromCodeGenFlags();
+}
+
+std::string getCPUStr() { return ::getCPUStr(); }
+std::string getFeaturesStr() { return ::getFeaturesStr(); }
+} // namespace opts
+
+#if LDC_WITH_LLD && LDC_LLVM_VER >= 500
+// LLD 5.0 uses the shared header too (for LTO) and exposes some wrappers in
+// the lld namespace. Define them here to prevent the LLD object from being
+// linked in with its conflicting command-line options.
+namespace lld {
+TargetOptions InitTargetOptionsFromCodeGenFlags() {
+  return ::InitTargetOptionsFromCodeGenFlags();
+}
+
+CodeModel::Model GetCodeModelFromCMModel() { return CMModel; }
+}
+#endif // LDC_WITH_LLD && LDC_LLVM_VER >= 500

driver/cl_options-llvm.h

@@ -0,0 +1,36 @@
+//===-- driver/cl_options-llvm.h - LLVM command line options ----*- C++ -*-===//
+//
+//                         LDC – the LLVM D compiler
+//
+// This file is distributed under the BSD-style LDC license. See the LICENSE
+// file for details.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LDC_DRIVER_CL_OPTIONS_LLVM_H
+#define LDC_DRIVER_CL_OPTIONS_LLVM_H
+
+#include "llvm/ADT/Optional.h"
+#include "llvm/Support/CommandLine.h"
+#include "llvm/Support/CodeGen.h"
+#include "llvm/Target/TargetOptions.h"
+
+namespace opts {
+
+std::string getArchStr();
+#if LDC_LLVM_VER >= 309
+llvm::Optional<llvm::Reloc::Model> getRelocModel();
+#else
+llvm::Reloc::Model getRelocModel();
+#endif
+llvm::CodeModel::Model getCodeModel();
+bool disableFPElim();
+bool disableRedZone();
+bool printTargetFeaturesHelp();
+
+llvm::TargetOptions InitTargetOptionsFromCodeGenFlags();
+std::string getCPUStr();
+std::string getFeaturesStr();
+}
+
+#endif

driver/cl_options.cpp

@@ -134,10 +134,6 @@ static cl::opt<ubyte, true> debugInfo(
         clEnumValN(3, "gline-tables-only", "Add line tables only")),
     cl::location(global.params.symdebug), cl::init(0));
 
-static cl::opt<unsigned, true>
-    dwarfVersion("dwarf-version", cl::desc("Dwarf version"), cl::ZeroOrMore,
-                 cl::location(global.params.dwarfVersion), cl::Hidden);
-
 cl::opt<bool> noAsm("noasm", cl::desc("Disallow use of inline assembler"),
                     cl::ZeroOrMore);
 
@@ -175,12 +171,6 @@ static cl::opt<bool, true>
                        cl::desc("Remove generated object files on success"),
                        cl::location(global.params.cleanupObjectFiles));
 
-// Disabling Red Zone
-cl::opt<bool, true>
-    disableRedZone("disable-red-zone", cl::ZeroOrMore,
-                   cl::desc("Do not emit code that uses the red zone."),
-                   cl::location(global.params.disableRedZone));
-
 // DDoc options
 static cl::opt<bool, true> doDdoc("D", cl::desc("Generate documentation"),
                                   cl::location(global.params.doDocComments),
@@ -296,21 +286,10 @@ cl::opt<std::string>
                         "'-deps' alone prints module dependencies "
                         "(imports/file/version/debug/lib)"));
 
-cl::opt<std::string> mArch("march", cl::ZeroOrMore,
-                           cl::desc("Architecture to generate code for:"));
-
 cl::opt<bool> m32bits("m32", cl::desc("32 bit target"), cl::ZeroOrMore);
 
 cl::opt<bool> m64bits("m64", cl::desc("64 bit target"), cl::ZeroOrMore);
 
-cl::opt<std::string>
-    mCPU("mcpu", cl::ZeroOrMore, cl::value_desc("cpu-name"), cl::init(""),
-         cl::desc("Target a specific cpu type (-mcpu=help for details)"));
-
-cl::list<std::string>
-    mAttrs("mattr", cl::CommaSeparated, cl::value_desc("a1,+a2,-a3,..."),
-           cl::desc("Target specific attributes (-mattr=help for details)"));
-
 cl::opt<std::string> mTargetTriple("mtriple", cl::ZeroOrMore,
                                    cl::desc("Override target triple"));
 
@@ -325,54 +304,7 @@ static cl::list<std::string, StringsAdapter> modFileAliasStrings(
     cl::value_desc("<package.module>=<filespec>"),
     cl::location(modFileAliasStringsStore));
 
-cl::opt<llvm::Reloc::Model> mRelocModel(
+FloatABI::Type floatABI; // Storage for the dynamically created float-abi option.
-    "relocation-model", cl::desc("Relocation model"), cl::ZeroOrMore,
-#if LDC_LLVM_VER < 309
-    cl::init(llvm::Reloc::Default),
-#endif
-    clEnumValues(
-#if LDC_LLVM_VER < 309
-        clEnumValN(llvm::Reloc::Default, "default",
-                   "Target default relocation model"),
-#endif
-        clEnumValN(llvm::Reloc::Static, "static", "Non-relocatable code"),
-        clEnumValN(llvm::Reloc::PIC_, "pic",
-                   "Fully relocatable, position independent code"),
-        clEnumValN(llvm::Reloc::DynamicNoPIC, "dynamic-no-pic",
-                   "Relocatable external references, non-relocatable code")));
-
-cl::opt<llvm::CodeModel::Model> mCodeModel(
-    "code-model", cl::desc("Code model"), cl::ZeroOrMore,
-#if LDC_LLVM_VER < 600
-    cl::init(llvm::CodeModel::Default),
-    clEnumValues(
-        clEnumValN(llvm::CodeModel::Default, "default",
-                   "Target default code model"),
-#else
-    cl::init(llvm::CodeModel::Small),
-    clEnumValues(
-#endif
-        clEnumValN(llvm::CodeModel::Small, "small", "Small code model"),
-        clEnumValN(llvm::CodeModel::Kernel, "kernel", "Kernel code model"),
-        clEnumValN(llvm::CodeModel::Medium, "medium", "Medium code model"),
-        clEnumValN(llvm::CodeModel::Large, "large", "Large code model")));
-
-cl::opt<FloatABI::Type> mFloatABI(
-    "float-abi", cl::desc("ABI/operations to use for floating-point types:"),
-    cl::ZeroOrMore, cl::init(FloatABI::Default),
-    clEnumValues(
-        clEnumValN(FloatABI::Default, "default",
-                   "Target default floating-point ABI"),
-        clEnumValN(FloatABI::Soft, "soft",
-                   "Software floating-point ABI and operations"),
-        clEnumValN(FloatABI::SoftFP, "softfp",
-                   "Soft-float ABI, but hardware floating-point instructions"),
-        clEnumValN(FloatABI::Hard, "hard",
-                   "Hardware floating-point ABI and instructions")));
-
-cl::opt<bool>
-    disableFpElim("disable-fp-elim", cl::ZeroOrMore,
-                  cl::desc("Disable frame pointer elimination optimization"));
 
 static cl::opt<bool, true, FlagParser<bool>>
     asserts("asserts", cl::ZeroOrMore, cl::desc("(*) Enable assertions"),
@@ -437,12 +369,10 @@ cl::opt<bool> disableLinkerStripDead(
 // Math options
 bool fFastMath; // Storage for the dynamically created ffast-math option.
 llvm::FastMathFlags defaultFMF;
-void setDefaultMathOptions(llvm::TargetMachine &target) {
+void setDefaultMathOptions(llvm::TargetOptions &targetOptions) {
   if (fFastMath) {
     defaultFMF.setUnsafeAlgebra();
-
+    targetOptions.UnsafeFPMath = true;
-    llvm::TargetOptions &TO = target.Options;
-    TO.UnsafeFPMath = true;
   }
 }
 
@@ -519,6 +449,11 @@ cl::list<std::string>
                              " list. Use 'ocl-xy0' for OpenCL x.y, and "
                              "'cuda-xy0' for CUDA CC x.y"),
                      cl::value_desc("targets"));
+cl::opt<std::string>
+    dcomputeFilePrefix("mdcompute-file-prefix",
+                       cl::desc("Prefix to prepend to the generated kernel files."),
+                       cl::init("kernels"),
+                       cl::value_desc("prefix"));
 #endif
 
 #if defined(LDC_RUNTIME_COMPILE)
@@ -562,6 +497,7 @@ void createClashingOptions() {
   // is a clash in the command line options.
   renameAndHide("color", "llvm-color");
   renameAndHide("ffast-math", "llvm-ffast-math");
+  renameAndHide("float-abi", "llvm-float-abi");
 
   // Step 2. Add the LDC options.
   new cl::opt<bool, true, FlagParser<bool>>(
@@ -569,6 +505,19 @@ void createClashingOptions() {
       cl::desc("(*) Force colored console output"));
   new cl::opt<bool, true>("ffast-math", cl::ZeroOrMore, cl::location(fFastMath),
                           cl::desc("Set @fastmath for all functions."));
+  new cl::opt<FloatABI::Type, true>(
+      "float-abi", cl::desc("ABI/operations to use for floating-point types:"),
+      cl::ZeroOrMore, cl::location(floatABI), cl::init(FloatABI::Default),
+      clEnumValues(
+          clEnumValN(FloatABI::Default, "default",
+                     "Target default floating-point ABI"),
+          clEnumValN(FloatABI::Soft, "soft",
+                     "Software floating-point ABI and operations"),
+          clEnumValN(
+              FloatABI::SoftFP, "softfp",
+              "Soft-float ABI, but hardware floating-point instructions"),
+          clEnumValN(FloatABI::Hard, "hard",
+                     "Hardware floating-point ABI and instructions")));
 }
 
 /// Hides command line options exposed from within LLVM that are unlikely
@@ -605,6 +554,15 @@ void hideLLVMOptions() {
       "verify-region-info", "verify-scev", "verify-scev-maps",
       "x86-early-ifcvt", "x86-use-vzeroupper", "x86-recip-refinement-steps",
 
+      "thread-model", "exception-model", "enable-fp-mad",
+      "enable-unsafe-fp-math", "enable-no-infs-fp-math",
+      "enable-no-nans-fp-math", "enable-no-trapping-fp-math",
+      "denormal-fp-math", "recip", "nozero-initialized-in-bss", "tailcallopt",
+      "stack-symbol-ordering", "stack-alignment", "enable-pie", "use-ctors",
+      "emulated-tls", "unique-section-names", "jump-table-type", "meabi",
+      "debugger-tune", "asm-instrumentation", "mc-relax-all",
+      "incremental-linker-compatible", "asm-show-inst", "pie-copy-relocations",
+
       // We enable -fdata-sections/-ffunction-sections by default where it makes
       // sense for reducing code size, so hide them to avoid confusion.
       //
@@ -613,7 +571,20 @@ void hideLLVMOptions() {
       // on the target triple (and thus we do not know it until after the
       // command
       // line has been parsed).
-      "fdata-sections", "ffunction-sections"};
+      "fdata-sections", "ffunction-sections", "data-sections",
+      "function-sections"};
+
+  // pulled in from shared LLVM headers, but unused or not desired in LDC
+  static const char *const removedOptions[] = {"disable-tail-calls",
+                                               "fatal-warnings",
+                                               "filetype",
+                                               "no-deprecated-warn",
+                                               "no-warn",
+                                               "stackrealign",
+                                               "start-after",
+                                               "stop-after",
+                                               "trap-func",
+                                               "W"};
 
   llvm::StringMap<cl::Option *> &map = cl::getRegisteredOptions();
   for (const auto name : hiddenOptions) {
@@ -624,6 +595,13 @@ void hideLLVMOptions() {
       it->second->setHiddenFlag(cl::Hidden);
     }
   }
+
+  for (const auto name : removedOptions) {
+    auto it = map.find(name);
+    if (it != map.end()) {
+      map.erase(it);
+    }
+  }
 }
 
 } // namespace opts

driver/cl_options.h

@@ -15,6 +15,7 @@
 #ifndef LDC_DRIVER_CL_OPTIONS_H
 #define LDC_DRIVER_CL_OPTIONS_H
 
+#include "driver/cl_options-llvm.h"
 #include "driver/targetmachine.h"
 #include "gen/cl_helpers.h"
 #include "llvm/ADT/SmallVector.h"
@@ -69,24 +70,18 @@ extern cl::opt<std::string> cacheDir;
 extern cl::list<std::string> linkerSwitches;
 extern cl::list<std::string> ccSwitches;
 
-extern cl::opt<std::string> mArch;
 extern cl::opt<bool> m32bits;
 extern cl::opt<bool> m64bits;
-extern cl::opt<std::string> mCPU;
-extern cl::list<std::string> mAttrs;
 extern cl::opt<std::string> mTargetTriple;
 extern cl::opt<std::string> mABI;
-extern cl::opt<llvm::Reloc::Model> mRelocModel;
+extern FloatABI::Type floatABI;
-extern cl::opt<llvm::CodeModel::Model> mCodeModel;
-extern cl::opt<bool> disableFpElim;
-extern cl::opt<FloatABI::Type> mFloatABI;
 extern cl::opt<bool> linkonceTemplates;
 extern cl::opt<bool> disableLinkerStripDead;
 
 // Math options
 extern bool fFastMath;
 extern llvm::FastMathFlags defaultFMF;
-void setDefaultMathOptions(llvm::TargetMachine &target);
+void setDefaultMathOptions(llvm::TargetOptions &targetOptions);
 
 extern cl::opt<BOUNDSCHECK> boundsCheck;
 extern bool nonSafeBoundsChecks;
@@ -124,6 +119,7 @@ extern cl::opt<std::string> saveOptimizationRecord;
 #endif
 #if LDC_LLVM_SUPPORTED_TARGET_SPIRV || LDC_LLVM_SUPPORTED_TARGET_NVPTX
 extern cl::list<std::string> dcomputeTargets;
+extern cl::opt<std::string> dcomputeFilePrefix;
 #endif
 
 #if defined(LDC_RUNTIME_COMPILE)

driver/cl_options_sanitizers.cpp

@@ -198,8 +198,18 @@ void outputSanitizerSettings(llvm::raw_ostream &hash_os) {
 }
 
 bool functionIsInSanitizerBlacklist(FuncDeclaration *funcDecl) {
-  return sanitizerBlacklist &&
+  if (!sanitizerBlacklist)
-         sanitizerBlacklist->inSection("fun", mangleExact(funcDecl));
+    return false;
+
+#if LDC_LLVM_VER >= 600
+  // TODO: LLVM 6.0 supports sections (e.g. "[address]") in the blacklist file
+  // to only blacklist a function for a particular sanitizer. We could make use
+  // of that too.
+  return sanitizerBlacklist->inSection("" /* section name */, "fun",
+                                       mangleExact(funcDecl));
+#else
+  return sanitizerBlacklist->inSection("fun", mangleExact(funcDecl));
+#endif
 }
 
 } // namespace opts

driver/codegenerator.cpp

@@ -31,12 +31,18 @@ extern Module *g_entrypointModule;
 /// The module that contains the actual D main() (_Dmain) definition.
 extern Module *g_dMainModule;
 
+#if LDC_LLVM_VER < 600
+namespace llvm {
+  using ToolOutputFile = tool_output_file;
+}
+#endif
+
 namespace {
 
-std::unique_ptr<llvm::tool_output_file>
+std::unique_ptr<llvm::ToolOutputFile>
 createAndSetDiagnosticsOutputFile(IRState &irs, llvm::LLVMContext &ctx,
                                   llvm::StringRef filename) {
-  std::unique_ptr<llvm::tool_output_file> diagnosticsOutputFile;
+  std::unique_ptr<llvm::ToolOutputFile> diagnosticsOutputFile;
 
 #if LDC_LLVM_VER >= 400
   // Set LLVM Diagnostics outputfile if requested
@@ -50,7 +56,7 @@ createAndSetDiagnosticsOutputFile(IRState &irs, llvm::LLVMContext &ctx,
     }
 
     std::error_code EC;
-    diagnosticsOutputFile = llvm::make_unique<llvm::tool_output_file>(
+    diagnosticsOutputFile = llvm::make_unique<llvm::ToolOutputFile>(
         diagnosticsFilename, EC, llvm::sys::fs::F_None);
     if (EC) {
       irs.dmodule->error("Could not create file %s: %s",
@@ -260,7 +266,7 @@ void CodeGenerator::writeAndFreeLLModule(const char *filename) {
   llvm::Metadata *IdentNode[] = {llvm::MDString::get(ir_->context(), Version)};
   IdentMetadata->addOperand(llvm::MDNode::get(ir_->context(), IdentNode));
 
-  std::unique_ptr<llvm::tool_output_file> diagnosticsOutputFile =
+  std::unique_ptr<llvm::ToolOutputFile> diagnosticsOutputFile =
       createAndSetDiagnosticsOutputFile(*ir_, context_, filename);
 
   writeModule(&ir_->module, filename);

driver/dcomputecodegenerator.cpp

@@ -17,47 +17,67 @@
 #include <string>
 #include <algorithm>
 
+#if !(LDC_LLVM_SUPPORTED_TARGET_SPIRV || LDC_LLVM_SUPPORTED_TARGET_NVPTX)
+
+DComputeCodeGenManager::DComputeCodeGenManager(llvm::LLVMContext &c) : ctx(c) {}
+void DComputeCodeGenManager::emit(Module *) {}
+void DComputeCodeGenManager::writeModules() {}
+DComputeCodeGenManager::~DComputeCodeGenManager() {}
+
+#else
+
 DComputeTarget *
 DComputeCodeGenManager::createComputeTarget(const std::string &s) {
-  int v;
+#if LDC_LLVM_SUPPORTED_TARGET_SPIRV
 #define OCL_VALID_VER_INIT 100, 110, 120, 200, 210, 220
   const std::array<int, 6> valid_ocl_versions = {{OCL_VALID_VER_INIT}};
-#define CUDA_VALID_VER_INIT 100, 110, 120, 130, 200, 210, 300, 350, 370,\
- 500, 520, 600, 610, 620
-  const std::array<int, 14> vaild_cuda_versions = {{CUDA_VALID_VER_INIT}};
 
   if (s.substr(0, 4) == "ocl-") {
-    v = atoi(s.c_str() + 4);
+    const int v = atoi(s.c_str() + 4);
     if (std::find(valid_ocl_versions.begin(), valid_ocl_versions.end(), v) !=
         valid_ocl_versions.end()) {
       return createOCLTarget(ctx, v);
     }
-  } else if (s.substr(0, 5) == "cuda-") {
+  }
-    v = atoi(s.c_str() + 5);
+#endif
 
-    if (std::find(vaild_cuda_versions.begin(), vaild_cuda_versions.end(), v) !=
+#if LDC_LLVM_SUPPORTED_TARGET_NVPTX
-        vaild_cuda_versions.end()) {
+#define CUDA_VALID_VER_INIT 100, 110, 120, 130, 200, 210, 300, 350, 370,\
+ 500, 520, 600, 610, 620
+  const std::array<int, 14> valid_cuda_versions = {{CUDA_VALID_VER_INIT}};
+
+  if (s.substr(0, 5) == "cuda-") {
+    const int v = atoi(s.c_str() + 5);
+    if (std::find(valid_cuda_versions.begin(), valid_cuda_versions.end(), v) !=
+        valid_cuda_versions.end()) {
       return createCUDATarget(ctx, v);
     }
   }
+#endif
+
 #define XSTR(x) #x
 #define STR(x) XSTR((x))
 
   error(Loc(),
         "unrecognised or invalid DCompute targets: the format is ocl-xy0 "
-        "for OpenCl x.y and cuda-xy0 for CUDA CC x.y. Valid versions "
+        "for OpenCl x.y and cuda-xy0 for CUDA CC x.y."
-        "for OpenCl are " STR(OCL_VALID_VER_INIT) ". Valid versions for CUDA "
+#if LDC_LLVM_SUPPORTED_TARGET_SPIRV
-        "are " STR(CUDA_VALID_VER_INIT));
+        " Valid versions for OpenCl are " STR(OCL_VALID_VER_INIT) "."
+#endif
+#if LDC_LLVM_SUPPORTED_TARGET_NVPTX
+        " Valid versions for CUDA are " STR(CUDA_VALID_VER_INIT)
+#endif
+  );
   fatal();
   return nullptr;
 }
 
 DComputeCodeGenManager::DComputeCodeGenManager(llvm::LLVMContext &c) : ctx(c) {
-#if LDC_LLVM_SUPPORTED_TARGET_SPIRV || LDC_LLVM_SUPPORTED_TARGET_NVPTX
   for (auto &option : opts::dcomputeTargets) {
     targets.push_back(createComputeTarget(option));
   }
-#endif
+  oldGIR = gIR;
+  oldGTargetMachine = gTargetMachine;
 }
 
 void DComputeCodeGenManager::emit(Module *m) {
@@ -72,3 +92,10 @@ void DComputeCodeGenManager::writeModules() {
     target->writeModule();
   }
 }
+
+DComputeCodeGenManager::~DComputeCodeGenManager() {
+  gIR = oldGIR;
+  gTargetMachine = oldGTargetMachine;
+}
+
+#endif // LDC_LLVM_SUPPORTED_TARGET_SPIRV || LDC_LLVM_SUPPORTED_TARGET_NVPTX

driver/dcomputecodegenerator.h

@@ -13,6 +13,10 @@
 #include "gen/dcompute/target.h"
 #include "llvm/ADT/SmallVector.h"
 
+namespace llvm {
+  class TargetMachine;
+}
+
 // gets run on modules marked @compute
 // All @compute D modules are emitted into one LLVM module once per target.
 class DComputeCodeGenManager {
@@ -20,12 +24,14 @@ class DComputeCodeGenManager {
   llvm::LLVMContext &ctx;
   llvm::SmallVector<DComputeTarget *, 2> targets;
   DComputeTarget *createComputeTarget(const std::string &s);
-
+  IRState *oldGIR = nullptr;
+  llvm::TargetMachine *oldGTargetMachine = nullptr;
 public:
   void emit(Module *m);
   void writeModules();
 
   DComputeCodeGenManager(llvm::LLVMContext &c);
+  ~DComputeCodeGenManager();
 };
 
 #endif

driver/ldmd.cpp

@@ -107,7 +107,13 @@ char *concat(const char *a, int b) {
  */
 int execute(const std::string &exePath, const char **args) {
   std::string errorMsg;
-  int rc = ls::ExecuteAndWait(exePath, args, nullptr, nullptr, 0, 0, &errorMsg);
+  int rc = ls::ExecuteAndWait(exePath, args, nullptr,
+#if LDC_LLVM_VER >= 600
+                              {},
+#else
+                              nullptr,
+#endif
+                              0, 0, &errorMsg);
   if (!errorMsg.empty()) {
     error("Error executing %s: %s", exePath.c_str(), errorMsg.c_str());
   }

driver/linker-gcc.cpp

@@ -118,8 +118,9 @@ void ArgsBuilder::addLTOGoldPluginFlags() {
   if (opts::isUsingThinLTO())
     addLdFlag("-plugin-opt=thinlto");
 
-  if (!opts::mCPU.empty())
+  const auto cpu = gTargetMachine->getTargetCPU();
-    addLdFlag(llvm::Twine("-plugin-opt=mcpu=") + opts::mCPU);
+  if (!cpu.empty())
+    addLdFlag(llvm::Twine("-plugin-opt=mcpu=") + cpu);
 
   // Use the O-level passed to LDC as the O-level for LTO, but restrict it to
   // the [0, 3] range that can be passed to the linker plugin.

driver/main.cpp

@@ -104,18 +104,6 @@ static cl::opt<bool> linkDebugLib(
     "link-debuglib", cl::ZeroOrMore,
     cl::desc("Link with libraries specified in -debuglib, not -defaultlib"));
 
-#if LDC_LLVM_VER >= 309
-static inline llvm::Optional<llvm::Reloc::Model> getRelocModel() {
-  if (mRelocModel.getNumOccurrences()) {
-    llvm::Reloc::Model R = mRelocModel;
-    return R;
-  }
-  return llvm::None;
-}
-#else
-static inline llvm::Reloc::Model getRelocModel() { return mRelocModel; }
-#endif
-
 // This function exits the program.
 void printVersion(llvm::raw_ostream &OS) {
   OS << "LDC - the LLVM D compiler (" << global.ldc_version << "):\n";
@@ -372,8 +360,7 @@ void parseCommandLine(int argc, char **argv, Strings &sourceFiles,
                               const_cast<char **>(allArguments.data()),
                               "LDC - the LLVM D compiler\n");
 
-  helpOnly = mCPU == "help" ||
+  helpOnly = opts::printTargetFeaturesHelp();
-             (std::find(mAttrs.begin(), mAttrs.end(), "help") != mAttrs.end());
   if (helpOnly) {
     auto triple = llvm::Triple(cfg_triple);
     std::string errMsg;
@@ -599,19 +586,12 @@ void parseCommandLine(int argc, char **argv, Strings &sourceFiles,
     error(Loc(), "-lib and -shared switches cannot be used together");
   }
 
-#if LDC_LLVM_VER >= 309
-  if (global.params.dll && !mRelocModel.getNumOccurrences()) {
-#else
-  if (global.params.dll && mRelocModel == llvm::Reloc::Default) {
-#endif
-    mRelocModel = llvm::Reloc::PIC_;
-  }
-
   if (soname.getNumOccurrences() > 0 && !global.params.dll) {
     error(Loc(), "-soname can be used only when building a shared library");
   }
 
   global.params.hdrStripPlainFunctions = !opts::hdrKeepAllBodies;
+  global.params.disableRedZone = opts::disableRedZone();
 }
 
 void initializePasses() {
@@ -1014,7 +994,8 @@ int cppmain(int argc, char **argv) {
   }
 
   // Set up the TargetMachine.
-  if ((m32bits || m64bits) && (!mArch.empty() || !mTargetTriple.empty())) {
+  const auto arch = getArchStr();
+  if ((m32bits || m64bits) && (!arch.empty() || !mTargetTriple.empty())) {
     error(Loc(), "-m32 and -m64 switches cannot be used together with -march "
                  "and -mtriple switches");
   }
@@ -1029,9 +1010,21 @@ int cppmain(int argc, char **argv) {
     fatal();
   }
 
+  auto relocModel = getRelocModel();
+#if LDC_LLVM_VER >= 309
+  if (global.params.dll && !relocModel.hasValue()) {
+#else
+  if (global.params.dll && relocModel == llvm::Reloc::Default) {
+#endif
+    relocModel = llvm::Reloc::PIC_;
+  }
+
   gTargetMachine = createTargetMachine(
-      mTargetTriple, mArch, mCPU, mAttrs, bitness, mFloatABI, getRelocModel(),
+      mTargetTriple, arch, opts::getCPUStr(), opts::getFeaturesStr(), bitness,
-      mCodeModel, codeGenOptLevel(), disableFpElim, disableLinkerStripDead);
+      floatABI, relocModel, opts::getCodeModel(), codeGenOptLevel(),
+      disableLinkerStripDead);
+
+  opts::setDefaultMathOptions(gTargetMachine->Options);
 
 #if LDC_LLVM_VER >= 308
   static llvm::DataLayout DL = gTargetMachine->createDataLayout();
@@ -1047,6 +1040,7 @@ int cppmain(int argc, char **argv) {
     global.params.isLP64 = gDataLayout->getPointerSizeInBits() == 64;
     global.params.is64bit = triple->isArch64Bit();
     global.params.hasObjectiveC = objc_isSupported(*triple);
+    global.params.dwarfVersion = gTargetMachine->Options.MCOptions.DwarfVersion;
     // mscoff enables slightly different handling of interface functions
     // in the front end
     global.params.mscoff = triple->isKnownWindowsMSVCEnvironment();
@@ -1054,8 +1048,6 @@ int cppmain(int argc, char **argv) {
       global.obj_ext = "obj";
   }
 
-  opts::setDefaultMathOptions(*gTargetMachine);
-
   // allocate the target abi
   gABI = TargetABI::getTarget();
 

driver/targetmachine.cpp

@@ -15,6 +15,7 @@
 
 #include "driver/cl_options.h"
 #include "driver/targetmachine.h"
+#include "llvm/ADT/StringExtras.h"
 #include "llvm/ADT/StringSwitch.h"
 #include "llvm/ADT/Triple.h"
 #include "llvm/MC/SubtargetFeature.h"
@@ -27,6 +28,7 @@
 #include "llvm/Target/TargetOptions.h"
 #include "llvm/IR/Module.h"
 #include "mars.h"
+#include "driver/cl_options.h"
 #include "gen/logger.h"
 
 static const char *getABI(const llvm::Triple &triple) {
@@ -197,28 +199,13 @@ static std::string getAArch64TargetCPU(const llvm::Triple &triple) {
   return "generic";
 }
 
-/// Returns the LLVM name of the target CPU to use given the provided
+/// Returns the LLVM name of the default CPU for the provided target triple.
-/// -mcpu argument and target triple.
+static std::string getTargetCPU(const llvm::Triple &triple) {
-static std::string getTargetCPU(const std::string &cpu,
-                                const llvm::Triple &triple) {
-  if (!cpu.empty()) {
-    if (cpu != "native") {
-      return cpu;
-    }
-
-    // FIXME: Reject attempts to use -mcpu=native unless the target matches
-    // the host.
-    std::string hostCPU = llvm::sys::getHostCPUName();
-    if (!hostCPU.empty() && hostCPU != "generic") {
-      return hostCPU;
-    }
-  }
-
   switch (triple.getArch()) {
   default:
     // We don't know about the specifics of this platform, just return the
     // empty string and let LLVM decide.
-    return cpu;
+    return "";
   case llvm::Triple::x86:
   case llvm::Triple::x86_64:
     return getX86TargetCPU(triple);
@@ -355,17 +342,18 @@ const llvm::Target *lookupTarget(const std::string &arch, llvm::Triple &triple,
 }
 
 llvm::TargetMachine *
-createTargetMachine(std::string targetTriple, std::string arch, std::string cpu,
+createTargetMachine(const std::string targetTriple, const std::string arch,
-                    std::vector<std::string> attrs,
+                    std::string cpu, const std::string featuresString,
-                    ExplicitBitness::Type bitness, FloatABI::Type floatABI,
+                    const ExplicitBitness::Type bitness,
+                    FloatABI::Type floatABI,
 #if LDC_LLVM_VER >= 309
                     llvm::Optional<llvm::Reloc::Model> relocModel,
 #else
                     llvm::Reloc::Model relocModel,
 #endif
-                    llvm::CodeModel::Model codeModel,
+                    const llvm::CodeModel::Model codeModel,
-                    llvm::CodeGenOpt::Level codeGenOptLevel,
+                    const llvm::CodeGenOpt::Level codeGenOptLevel,
-                    bool noFramePointerElim, bool noLinkerStripDead) {
+                    const bool noLinkerStripDead) {
   // Determine target triple. If the user didn't explicitly specify one, use
   // the one set at LLVM configure time.
   llvm::Triple triple;
@@ -398,47 +386,40 @@ createTargetMachine(std::string targetTriple, std::string arch, std::string cpu,
     fatal();
   }
 
-  // Package up features to be passed to target/subtarget.
-  llvm::SubtargetFeatures features;
-  features.getDefaultSubtargetFeatures(triple);
-  if (cpu == "native") {
-    llvm::StringMap<bool> hostFeatures;
-    if (llvm::sys::getHostCPUFeatures(hostFeatures)) {
-      for (const auto &hf : hostFeatures) {
-        features.AddFeature(
-            std::string(hf.second ? "+" : "-").append(hf.first()));
-      }
-    }
-  }
-  for (auto &attr : attrs) {
-    features.AddFeature(attr);
-  }
-
   // With an empty CPU string, LLVM will default to the host CPU, which is
   // usually not what we want (expected behavior from other compilers is
   // to default to "generic").
-  cpu = getTargetCPU(cpu, triple);
+  if (cpu.empty() || cpu == "generic") {
+    cpu = getTargetCPU(triple);
+    if (cpu.empty())
+      cpu = "generic";
+  }
+
+  // Package up features to be passed to target/subtarget.
+  llvm::SmallVector<llvm::StringRef, 8> features;
+
+  // NOTE: needs a persistent (non-temporary) string
+  auto splitAndAddFeatures = [&features](llvm::StringRef str) {
+    str.split(features, ",", -1, /* KeepEmpty */ false);
+  };
+
+  llvm::SubtargetFeatures defaultSubtargetFeatures;
+  defaultSubtargetFeatures.getDefaultSubtargetFeatures(triple);
+  const std::string defaultSubtargetFeaturesString =
+      defaultSubtargetFeatures.getString();
+  splitAndAddFeatures(defaultSubtargetFeaturesString);
+
+  splitAndAddFeatures(featuresString);
 
   // cmpxchg16b is not available on old 64bit CPUs. Enable code generation
   // if the user did not make an explicit choice.
   if (cpu == "x86-64") {
-    const char *cx16_plus = "+cx16";
+    const bool has_cx16 =
-    const char *cx16_minus = "-cx16";
+        std::any_of(features.begin(), features.end(),
-    bool cx16 = false;
+                    [](llvm::StringRef f) { return f.substr(1) == "cx16"; });
-    for (auto &attr : attrs) {
+    if (!has_cx16) {
-      if (attr == cx16_plus || attr == cx16_minus) {
+      features.push_back("+cx16");
-        cx16 = true;
-      }
     }
-    if (!cx16) {
-      features.AddFeature(cx16_plus);
-    }
-  }
-
-  if (Logger::enabled()) {
-    Logger::println("Targeting '%s' (CPU '%s' with features '%s')",
-                    triple.str().c_str(), cpu.c_str(),
-                    features.getString().c_str());
   }
 
 // Handle cases where LLVM picks wrong default relocModel
@@ -473,26 +454,28 @@ createTargetMachine(std::string targetTriple, std::string arch, std::string cpu,
     }
   }
 
-  if (floatABI == FloatABI::Default) {
+  llvm::TargetOptions targetOptions = opts::InitTargetOptionsFromCodeGenFlags();
+  if (targetOptions.MCOptions.ABIName.empty())
+    targetOptions.MCOptions.ABIName = getABI(triple);
+
+  auto ldcFloatABI = floatABI;
+  if (ldcFloatABI == FloatABI::Default) {
     switch (triple.getArch()) {
     default: // X86, ...
-      floatABI = FloatABI::Hard;
+      ldcFloatABI = FloatABI::Hard;
       break;
     case llvm::Triple::arm:
     case llvm::Triple::thumb:
-      floatABI = getARMFloatABI(triple, getLLVMArchSuffixForARM(cpu));
+      ldcFloatABI = getARMFloatABI(triple, getLLVMArchSuffixForARM(cpu));
       break;
     }
   }
 
-  llvm::TargetOptions targetOptions;
+  switch (ldcFloatABI) {
-  targetOptions.MCOptions.ABIName = getABI(triple);
-
-  switch (floatABI) {
   default:
     llvm_unreachable("Floating point ABI type unknown.");
   case FloatABI::Soft:
-    features.AddFeature("+soft-float");
+    features.push_back("+soft-float");
     targetOptions.FloatABIType = llvm::FloatABI::Soft;
     break;
   case FloatABI::SoftFP:
@@ -514,11 +497,20 @@ createTargetMachine(std::string targetTriple, std::string arch, std::string cpu,
     targetOptions.DataSections = true;
   }
 
-  return target->createTargetMachine(triple.str(), cpu, features.getString(),
+  const std::string finalFeaturesString =
-                                     targetOptions, relocModel, codeModel,
+      llvm::join(features.begin(), features.end(), ",");
+
+  if (Logger::enabled()) {
+    Logger::println("Targeting '%s' (CPU '%s' with features '%s')",
+                    triple.str().c_str(), cpu.c_str(),
+                    finalFeaturesString.c_str());
+  }
+
+  return target->createTargetMachine(triple.str(), cpu, finalFeaturesString,
+                                     targetOptions, relocModel, opts::getCodeModel(),
                                      codeGenOptLevel);
 }
-    
+
 ComputeBackend::Type getComputeTargetType(llvm::Module* m) {
   llvm::Triple::ArchType a = llvm::Triple(m->getTargetTriple()).getArch();
   if (a == llvm::Triple::spir || a == llvm::Triple::spir64)

driver/targetmachine.h

@@ -52,18 +52,19 @@ ComputeBackend::Type getComputeTargetType(llvm::Module*);
  * parameters and the host platform defaults.
  *
  * Does not depend on any global state.
- */
+*/
-llvm::TargetMachine *createTargetMachine(
+llvm::TargetMachine *
-    std::string targetTriple, std::string arch, std::string cpu,
+createTargetMachine(std::string targetTriple, std::string arch, std::string cpu,
-    std::vector<std::string> attrs, ExplicitBitness::Type bitness,
+                    std::string featuresString, ExplicitBitness::Type bitness,
-    FloatABI::Type floatABI,
+                    FloatABI::Type floatABI,
 #if LDC_LLVM_VER >= 309
-    llvm::Optional<llvm::Reloc::Model> relocModel,
+                    llvm::Optional<llvm::Reloc::Model> relocModel,
 #else
-    llvm::Reloc::Model relocModel,
+                    llvm::Reloc::Model relocModel,
 #endif
-    llvm::CodeModel::Model codeModel, llvm::CodeGenOpt::Level codeGenOptLevel,
+                    llvm::CodeModel::Model codeModel,
-    bool noFramePointerElim, bool noLinkerStripDead);
+                    llvm::CodeGenOpt::Level codeGenOptLevel,
+                    bool noLinkerStripDead);
 
 /**
  * Returns the Mips ABI which is used for code generation.

driver/tool.cpp

@@ -171,7 +171,12 @@ int executeToolAndWait(const std::string &tool_,
   // Execute tool.
   std::string errstr;
   if (int status = llvm::sys::ExecuteAndWait(tool, &realargs[0], nullptr,
-                                             nullptr, 0, 0, &errstr)) {
+#if LDC_LLVM_VER >= 600
+                                             {},
+#else
+                                             nullptr,
+#endif
+                                             0, 0, &errstr)) {
     error(Loc(), "%s failed with status: %d", tool.c_str(), status);
     if (!errstr.empty()) {
       error(Loc(), "message: %s", errstr.c_str());

gen/dcompute/target.cpp

@@ -7,19 +7,21 @@
 //
 //===----------------------------------------------------------------------===//
 
+#if LDC_LLVM_SUPPORTED_TARGET_SPIRV || LDC_LLVM_SUPPORTED_TARGET_NVPTX
+
 #include "ddmd/dsymbol.h"
 #include "ddmd/mars.h"
 #include "ddmd/module.h"
 #include "ddmd/scope.h"
 #include "driver/linker.h"
 #include "driver/toobj.h"
+#include "driver/cl_options.h"
 #include "gen/dcompute/target.h"
 #include "gen/llvmhelpers.h"
 #include "gen/runtime.h"
 #include <string>
 
 void DComputeTarget::doCodeGen(Module *m) {
-
   // process module members
   for (unsigned k = 0; k < m->members->dim; k++) {
     Dsymbol *dsym = (*m->members)[k];
@@ -45,7 +47,7 @@ void DComputeTarget::writeModule() {
 
   std::string filename;
   llvm::raw_string_ostream os(filename);
-  os << "kernels_" << short_name << tversion << '_'
+  os << opts::dcomputeFilePrefix << '_' << short_name << tversion << '_'
      << (global.params.is64bit ? 64 : 32) << '.' << binSuffix;
 
   const char *path = FileName::combine(global.params.objdir, os.str().c_str());
@@ -56,3 +58,5 @@ void DComputeTarget::writeModule() {
   delete _ir;
   _ir = nullptr;
 }
+
+#endif

gen/dcompute/target.h

@@ -9,10 +9,12 @@
 
 #ifndef LDC_GEN_DCOMPUTE_TARGET_H
 #define LDC_GEN_DCOMPUTE_TARGET_H
+
 #include "gen/irstate.h"
 #include "llvm/IR/LLVMContext.h"
 #include "llvm/IR/Function.h"
 #include <array>
+
 namespace llvm {
 class Module;
 class Function;
@@ -51,7 +53,12 @@ public:
   virtual void addKernelMetadata(FuncDeclaration *df, llvm::Function *llf) = 0;
 };
 
+#if LDC_LLVM_SUPPORTED_TARGET_NVPTX
 DComputeTarget *createCUDATarget(llvm::LLVMContext &c, int sm);
+#endif
+
+#if LDC_LLVM_SUPPORTED_TARGET_SPIRV
 DComputeTarget *createOCLTarget(llvm::LLVMContext &c, int oclver);
+#endif
 
 #endif

gen/dcompute/targetCUDA.cpp

@@ -7,6 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
+#if LDC_LLVM_SUPPORTED_TARGET_NVPTX
+
 #include "gen/dcompute/target.h"
 #include "gen/dcompute/druntime.h"
 #include "gen/metadata.h"
@@ -54,7 +56,7 @@ public:
         is64 ? "nvptx64-nvidia-cuda" : "nvptx-nvidia-cuda",
         is64 ? "nvptx64" : "nvptx", "sm_" + ldc::to_string(tversion / 10), {},
         is64 ? ExplicitBitness::M64 : ExplicitBitness::M32, ::FloatABI::Hard,
-        llvm::Reloc::Static, llvm::CodeModel::Medium, codeGenOptLevel(), false,
+        llvm::Reloc::Static, llvm::CodeModel::Medium, codeGenOptLevel(),
         false);
   }
 
@@ -77,3 +79,5 @@ public:
 DComputeTarget *createCUDATarget(llvm::LLVMContext &c, int sm) {
   return new TargetCUDA(c, sm);
 };
+
+#endif // LDC_LLVM_SUPPORTED_TARGET_NVPTX

gen/dcompute/targetOCL.cpp

@@ -7,6 +7,8 @@
 // See the LICENSE file for details.
 //===----------------------------------------------------------------------===//
 
+#if LDC_LLVM_SUPPORTED_TARGET_SPIRV
+
 #include "gen/dcompute/target.h"
 #include "gen/dcompute/druntime.h"
 #include "ddmd/template.h"
@@ -198,3 +200,5 @@ public:
 DComputeTarget *createOCLTarget(llvm::LLVMContext &c, int oclver) {
   return new TargetOCL(c, oclver);
 }
+
+#endif // LDC_LLVM_SUPPORTED_TARGET_SPIRV

gen/dibuilder.cpp

@@ -113,12 +113,26 @@ ldc::DIScope ldc::DIBuilder::GetCurrentScope() {
   return fn->diLexicalBlocks.top();
 }
 
-void ldc::DIBuilder::Declare(const Loc &loc, llvm::Value *var,
+// Sets the memory address for a debuginfo variable.
+void ldc::DIBuilder::Declare(const Loc &loc, llvm::Value *storage,
                              ldc::DILocalVariable divar,
                              ldc::DIExpression diexpr) {
   unsigned charnum = (loc.linnum ? loc.charnum : 0);
   auto debugLoc = llvm::DebugLoc::get(loc.linnum, charnum, GetCurrentScope());
-  DBuilder.insertDeclare(var, divar, diexpr, debugLoc, IR->scopebb());
+  DBuilder.insertDeclare(storage, divar, diexpr, debugLoc, IR->scopebb());
+}
+
+// Sets the (current) value for a debuginfo variable.
+void ldc::DIBuilder::SetValue(const Loc &loc, llvm::Value *value,
+                              ldc::DILocalVariable divar,
+                              ldc::DIExpression diexpr) {
+  unsigned charnum = (loc.linnum ? loc.charnum : 0);
+  auto debugLoc = llvm::DebugLoc::get(loc.linnum, charnum, GetCurrentScope());
+  DBuilder.insertDbgValueIntrinsic(value,
+#if LDC_LLVM_VER < 600
+                                   0,
+#endif
+                                   divar, diexpr, debugLoc, IR->scopebb());
 }
 
 ldc::DIFile ldc::DIBuilder::CreateFile(Loc &loc) {
@@ -1020,7 +1034,7 @@ void ldc::DIBuilder::EmitValue(llvm::Value *val, VarDeclaration *vd) {
 
 void ldc::DIBuilder::EmitLocalVariable(llvm::Value *ll, VarDeclaration *vd,
                                        Type *type, bool isThisPtr,
-                                       bool forceAsLocal,
+                                       bool forceAsLocal, bool isRefRVal,
                                        llvm::ArrayRef<int64_t> addr) {
   if (!mustEmitFullDebugInfo())
     return;
@@ -1040,12 +1054,28 @@ void ldc::DIBuilder::EmitLocalVariable(llvm::Value *ll, VarDeclaration *vd,
   if (static_cast<llvm::MDNode *>(TD) == nullptr)
     return; // unsupported
 
-  if (vd->isRef() || vd->isOut()) {
+  const bool isRefOrOut = vd->isRef() || vd->isOut(); // incl. special-ref vars
+
+  // For MSVC x64 targets, declare params rewritten by ExplicitByvalRewrite as
+  // DI references, as if they were ref parameters.
+  const bool isPassedExplicitlyByval =
+      isTargetMSVCx64 && !isRefOrOut && isaArgument(ll) && addr.empty();
+
+  bool useDbgValueIntrinsic = false;
+  if (isRefOrOut || isPassedExplicitlyByval) {
+    // With the exception of special-ref loop variables, the reference/pointer
+    // itself is constant. So we don't have to attach the debug information to a
+    // memory location and can use llvm.dbg.value to set the constant pointer
+    // for the DI reference.
+    useDbgValueIntrinsic =
+        isPassedExplicitlyByval || (!isSpecialRefVar(vd) && isRefRVal);
 #if LDC_LLVM_VER >= 308
-    auto T = DtoType(type);
+    // Note: createReferenceType expects the size to be the size of a pointer,
-    TD = DBuilder.createReferenceType(llvm::dwarf::DW_TAG_reference_type, TD,
+    // not the size of the type the reference refers to.
-                                      getTypeAllocSize(T) * 8, // size (bits)
+    TD = DBuilder.createReferenceType(
-                                      DtoAlignment(type) * 8); // align (bits)
+        llvm::dwarf::DW_TAG_reference_type, TD,
+        gDataLayout->getPointerSizeInBits(), // size (bits)
+        DtoAlignment(type) * 8);             // align (bits)
 #else
     TD = DBuilder.createReferenceType(llvm::dwarf::DW_TAG_reference_type, TD);
 #endif
@@ -1122,10 +1152,15 @@ void ldc::DIBuilder::EmitLocalVariable(llvm::Value *ll, VarDeclaration *vd,
 #endif
   variableMap[vd] = debugVariable;
 
-  // declare
+  if (useDbgValueIntrinsic) {
-  Declare(vd->loc, ll, debugVariable, addr.empty()
+    SetValue(vd->loc, ll, debugVariable,
-                                          ? DBuilder.createExpression()
+             addr.empty() ? DBuilder.createExpression()
-                                          : DBuilder.createExpression(addr));
+                          : DBuilder.createExpression(addr));
+  } else {
+    Declare(vd->loc, ll, debugVariable,
+            addr.empty() ? DBuilder.createExpression()
+                         : DBuilder.createExpression(addr));
+  }
 }
 
 void ldc::DIBuilder::EmitGlobalVariable(llvm::GlobalVariable *llVar,

gen/dibuilder.h

@@ -115,15 +115,22 @@ public:
 
   /// \brief Emits all things necessary for making debug info for a local
   /// variable vd.
-  /// \param ll       LL lvalue of the variable.
+  /// \param ll       LL value which, in combination with `addr`, yields the
+  /// storage/lvalue of the variable. For special-ref loop variables, specify
+  /// the storage/lvalue of the reference/pointer.
   /// \param vd       Variable declaration to emit debug info for.
   /// \param type     Type of variable if different from vd->type
   /// \param isThisPtr Variable is hidden this pointer
   /// \param forceAsLocal Emit as local even if the variable is a parameter
-  /// \param addr     An array of complex address operations.
+  /// \param isRefRVal Only relevant for ref/out parameters: indicates whether
+  /// ll & addr specify the reference's rvalue, i.e., the lvalue of the original
+  /// variable, instead of the reference's lvalue.
+  /// \param addr     An array of complex address operations encoding a DWARF
+  /// expression.
   void
   EmitLocalVariable(llvm::Value *ll, VarDeclaration *vd, Type *type = nullptr,
                     bool isThisPtr = false, bool forceAsLocal = false,
+                    bool isRefRVal = false,
                     llvm::ArrayRef<int64_t> addr = llvm::ArrayRef<int64_t>());
 
   /// \brief Emits all things necessary for making debug info for a global
@@ -139,8 +146,10 @@ private:
   llvm::LLVMContext &getContext();
   Module *getDefinedModule(Dsymbol *s);
   DIScope GetCurrentScope();
-  void Declare(const Loc &loc, llvm::Value *var, ldc::DILocalVariable divar,
+  void Declare(const Loc &loc, llvm::Value *storage, ldc::DILocalVariable divar,
                ldc::DIExpression diexpr);
+  void SetValue(const Loc &loc, llvm::Value *value, ldc::DILocalVariable divar,
+                ldc::DIExpression diexpr);
   void AddFields(AggregateDeclaration *sd, ldc::DIFile file,
                  llvm::SmallVector<llvm::Metadata *, 16> &elems);
   DIFile CreateFile(Loc &loc);

gen/functions.cpp

@@ -461,7 +461,7 @@ void applyTargetMachineAttributes(llvm::Function &func,
 
   // Frame pointer elimination
   func.addFnAttr("no-frame-pointer-elim",
-                 opts::disableFpElim ? "true" : "false");
+                 opts::disableFPElim() ? "true" : "false");
 }
 
 } // anonymous namespace
@@ -773,8 +773,14 @@ void defineParameters(IrFuncTy &irFty, VarDeclarations &parameters) {
       ++llArgIdx;
     }
 
-    if (global.params.symdebug)
+    // The debuginfos for captured params are handled later by
-      gIR->DBuilder.EmitLocalVariable(irparam->value, vd, paramType);
+    // DtoCreateNestedContext().
+    if (global.params.symdebug && vd->nestedrefs.dim == 0) {
+      // Reference (ref/out) parameters have no storage themselves as they are
+      // constant pointers, so pass the reference rvalue to EmitLocalVariable().
+      gIR->DBuilder.EmitLocalVariable(irparam->value, vd, paramType, false,
+                                      false, /*isRefRVal=*/true);
+    }
   }
 }
 

gen/nested.cpp

@@ -130,10 +130,13 @@ DValue *DtoNestedVariable(Loc &loc, Type *astype, VarDeclaration *vd,
     Logger::cout() << "of type: " << *val->getType() << '\n';
   }
   const bool isRefOrOut = vd->isRef() || vd->isOut();
-  if (!isSpecialRefVar(vd) && (byref || isRefOrOut)) {
+  if (isSpecialRefVar(vd)) {
+    // Handled appropriately by makeVarDValue() and EmitLocalVariable(), pass
+    // storage of pointer (reference lvalue).
+  } else if (byref || isRefOrOut) {
     val = DtoAlignedLoad(val);
-    // ref/out variables get a reference-debuginfo-type in
+    // ref/out variables get a reference-debuginfo-type in EmitLocalVariable();
-    // DIBuilder::EmitLocalVariable()
+    // pass the GEP as reference lvalue in that case.
     if (!isRefOrOut)
       gIR->DBuilder.OpDeref(dwarfAddrOps);
     IF_LOG {
@@ -143,11 +146,13 @@ DValue *DtoNestedVariable(Loc &loc, Type *astype, VarDeclaration *vd,
   }
 
   if (!skipDIDeclaration && global.params.symdebug) {
+#if LDC_LLVM_VER < 500
     // Because we are passing a GEP instead of an alloca to
     // llvm.dbg.declare, we have to make the address dereference explicit.
     gIR->DBuilder.OpDeref(dwarfAddrOps);
-    gIR->DBuilder.EmitLocalVariable(gep, vd, nullptr, false, true,
+#endif
-                                    dwarfAddrOps);
+    gIR->DBuilder.EmitLocalVariable(gep, vd, nullptr, false,
+                                    /*forceAsLocal=*/true, false, dwarfAddrOps);
   }
 
   return makeVarDValue(astype, vd, val);
@@ -378,12 +383,19 @@ static void DtoCreateNestedContextType(FuncDeclaration *fd) {
     irLocal.nestedDepth = depth;
 
     LLType *t = nullptr;
-    if (vd->isRef() || vd->isOut())
+    if (vd->isRef() || vd->isOut()) {
       t = DtoType(vd->type->pointerTo());
-    else if (vd->isParameter() && (vd->storage_class & STClazy))
+    } else if (vd->isParameter() && (vd->storage_class & STClazy)) {
-      t = getIrParameter(vd)->value->getType()->getContainedType(0);
+      // The LL type is a delegate (LL struct).
-    else
+      // Depending on the used TargetABI, the LL parameter is either a struct or
+      // a pointer to a struct (`byval` attribute, ExplicitByvalRewrite).
+      t = getIrParameter(vd)->value->getType();
+      if (t->isPointerTy())
+        t = t->getPointerElementType();
+      assert(t->isStructTy());
+    } else {
       t = DtoMemType(vd->type);
+    }
 
     builder.addType(t, getTypeAllocSize(t));
 
@@ -481,6 +493,7 @@ void DtoCreateNestedContext(FuncGenState &funcGen) {
         if (vd->isRef() || vd->isOut()) {
           Logger::println("Captured by reference, copying pointer to nested frame");
           DtoAlignedStore(parm->value, gep);
+          // pass GEP as reference lvalue to EmitLocalVariable()
         } else {
           Logger::println("Copying to nested frame");
           // The parameter value is an alloca'd stack slot.
@@ -497,11 +510,14 @@ void DtoCreateNestedContext(FuncGenState &funcGen) {
       }
 
       if (global.params.symdebug) {
-        LLSmallVector<int64_t, 2> addr;
+        LLSmallVector<int64_t, 1> dwarfAddrOps;
+#if LDC_LLVM_VER < 500
         // Because we are passing a GEP instead of an alloca to
         // llvm.dbg.declare, we have to make the address dereference explicit.
-        gIR->DBuilder.OpDeref(addr);
+        gIR->DBuilder.OpDeref(dwarfAddrOps);
-        gIR->DBuilder.EmitLocalVariable(gep, vd, nullptr, false, false, addr);
+#endif
+        gIR->DBuilder.EmitLocalVariable(gep, vd, nullptr, false, false, false,
+                                        dwarfAddrOps);
       }
     }
   }

ir/irtypeaggr.cpp

@@ -203,22 +203,44 @@ IrTypeAggr::IrTypeAggr(AggregateDeclaration *ad)
       aggr(ad) {}
 
 bool IrTypeAggr::isPacked(AggregateDeclaration *ad) {
-  const auto aggregateSize = (ad->sizeok == SIZEOKdone ? ad->structsize : ~0u);
+  // If the aggregate's size is unknown, any field with type alignment > 1 will
+  // make it packed.
+  unsigned aggregateSize = ~0u;
+  unsigned aggregateAlignment = 1;
+  if (ad->sizeok == SIZEOKdone) {
+    aggregateSize = ad->structsize;
+
+    const auto naturalAlignment = ad->alignsize;
+    auto explicitAlignment = STRUCTALIGN_DEFAULT;
+    if (auto sd = ad->isStructDeclaration())
+      explicitAlignment = sd->alignment;
+
+    aggregateAlignment = explicitAlignment == STRUCTALIGN_DEFAULT
+                             ? naturalAlignment
+                             : explicitAlignment;
+  }
+
+  // Classes apparently aren't padded; their size may not match the alignment.
+  assert((ad->isClassDeclaration() ||
+          (aggregateSize & (aggregateAlignment - 1)) == 0) &&
+         "Size not a multiple of alignment?");
 
   // For unions, only a subset of the fields are actually used for the IR type -
-  // don't care.
+  // don't care (about a few potentially needlessly packed IR structs).
   for (const auto field : ad->fields) {
-    // The aggregate's size and the field offset need to be multiples of the
+    const auto naturalFieldTypeAlignment = field->type->alignsize();
-    // field's natural alignment, otherwise the aggregate type is unnaturally
+    const auto explicitFieldTypeAlignment = field->type->alignment();
-    // aligned, and LLVM would insert padding.
+    const auto fieldTypeAlignment =
-    const auto naturalFieldAlignment = field->type->alignsize();
+        explicitFieldTypeAlignment == STRUCTALIGN_DEFAULT
-    const auto mask = naturalFieldAlignment - 1;
+            ? naturalFieldTypeAlignment
+            : explicitFieldTypeAlignment;
 
-    // If the aggregate's size is unknown, any field with natural alignment > 1
+    // The aggregate size, aggregate alignment and the field offset need to be
-    // will make it packed.
+    // multiples of the field type's alignment, otherwise the aggregate type is
-    if ((aggregateSize & mask) != 0 || (field->offset & mask) != 0) {
+    // unnaturally aligned, and LLVM would insert padding.
+    const auto mask = fieldTypeAlignment - 1;
+    if ((aggregateSize | aggregateAlignment | field->offset) & mask)
       return true;
-    }
   }
 
   return false;

runtime/CMakeLists.txt

@@ -1,7 +1,11 @@
-project(runtime)
+project(runtime C)
 
 cmake_minimum_required(VERSION 2.8.9)
 
+set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} ${CMAKE_CURRENT_SOURCE_DIR})
+
+include(CheckIncludeFile)
+
 # Verify required variables if this CMake project is NOT embedded in the LDC CMake project.
 if(NOT LDC_EXE)
     if(NOT LDC_EXE_FULL)
@@ -37,14 +41,19 @@ set(D_FLAGS               -w                                  CACHE STRING "Runt
 set(D_FLAGS_DEBUG         -g;-link-debuglib                   CACHE STRING "Runtime D compiler flags (debug libraries), separated by ';'")
 set(D_FLAGS_RELEASE       -O3;-release                        CACHE STRING "Runtime D compiler flags (release libraries), separated by ';'")
 set(COMPILE_ALL_D_FILES_AT_ONCE ON                            CACHE BOOL   "Compile all D files for a lib in a single command line instead of separately")
+set(RT_ARCHIVE_WITH_LDC   AUTO                                CACHE STRING "Whether to archive the static runtime libs via LDC instead of CMake archiver")
 set(RT_CFLAGS             ""                                  CACHE STRING "Runtime extra C compiler flags, separated by ' '")
 set(LD_FLAGS              ""                                  CACHE STRING "Runtime extra C linker flags, separated by ' '")
 set(C_SYSTEM_LIBS         AUTO                                CACHE STRING "C system libraries for linking shared libraries and test runners, separated by ';'")
-set(TARGET_SYSTEM         AUTO                                CACHE STRING "Targeted platform for cross-compilation (e.g., 'Linux;UNIX', 'Darwin;APPLE;UNIX', 'Windows;MSVC')")
+set(TARGET_SYSTEM         AUTO                                CACHE STRING "Target OS/toolchain for cross-compilation (e.g., 'Linux;UNIX', 'Darwin;APPLE;UNIX', 'Windows;MSVC')")
+set(LDC_TARGET_PRESET     ""                                  CACHE STRING "Use a preset target configuration for cross-compilation, by passing format OS-arch (e.g., 'Linux-arm', 'Mac-x64', 'Windows-aarch64', 'Android-x86')")
 
 set(CMAKE_INSTALL_LIBDIR ${CMAKE_INSTALL_PREFIX}/lib${LIB_SUFFIX})
 
-# Auto-detect TARGET_SYSTEM
+# Use LDC_TARGET_PRESET to configure target runtime
+include(PresetRuntimeConfiguration)
+
+# Auto-detect TARGET_SYSTEM from host
 if("${TARGET_SYSTEM}" STREQUAL "AUTO")
     set(TARGET_SYSTEM ${CMAKE_SYSTEM_NAME})
     # Additionally add MSVC, APPLE and/or UNIX
@@ -75,6 +84,16 @@ else()
     set(MULTILIB_SUFFIX 64)
 endif()
 
+if(${RT_ARCHIVE_WITH_LDC} STREQUAL "AUTO")
+    # LLVM's MSVC archiver (v4.0.1) apparently has issues with objects with debuginfos,
+    # leading to MS linker warnings when linking against the static debug runtime libs.
+    if("${TARGET_SYSTEM}" MATCHES "MSVC")
+        set(RT_ARCHIVE_WITH_LDC OFF)
+    else()
+        set(RT_ARCHIVE_WITH_LDC ON)
+    endif()
+endif()
+
 set(SHARED_LIBS_SUPPORTED OFF)
 if("${TARGET_SYSTEM}" MATCHES "Linux|FreeBSD|APPLE"
    AND (NOT "${TARGET_SYSTEM}" MATCHES "Android"))
@@ -297,7 +316,7 @@ foreach(target ${LLVM_TARGETS_TO_BUILD})
 endforeach()
 
 # Always build zlib and other C parts of the runtime in release mode, regardless
-# of what the user chose for LDC itself.
+# of what the user chose for LDC itself. Also add other C_FLAGS here.
 # 1) Set up CMAKE_C_FLAGS_RELEASE
 if("${TARGET_SYSTEM}" MATCHES "MSVC")
     # Omit all references to the default C runtime libs.
@@ -309,6 +328,11 @@ if("${TARGET_SYSTEM}" MATCHES "MSVC")
     # warning C4996: zlib uses 'deprecated' POSIX names
     append("/wd4131 /wd4206 /wd4996" CMAKE_C_FLAGS_RELEASE)
 endif()
+CHECK_INCLUDE_FILE(unistd.h HAVE_UNISTD_H)
+if (HAVE_UNISTD_H)
+    # Needed for zlib
+    append("-DHAVE_UNISTD_H" CMAKE_C_FLAGS_RELEASE)
+endif()
 # 2) Set all other CMAKE_C_FLAGS variants to CMAKE_C_FLAGS_RELEASE
 set(variables
     CMAKE_C_FLAGS_DEBUG
@@ -421,13 +445,24 @@ macro(compile_phobos2 d_flags lib_suffix path_suffix all_at_once outlist_o outli
     )
 endmacro()
 
+# Define an optional 'D' CMake linker language for the static runtime libs,
+# using LDC as archiver. LDC handles (cross-)archiving internally via LLVM
+# and supports LTO objects.
+set(CMAKE_D_CREATE_STATIC_LIBRARY "${LDC_EXE_FULL} -lib -of=<TARGET> <OBJECTS>")
+foreach(f ${D_FLAGS})
+    append("\"${f}\"" CMAKE_D_CREATE_STATIC_LIBRARY)
+endforeach()
+
 # Sets the common properties for all library targets.
-function(set_common_library_properties target)
+function(set_common_library_properties target is_shared)
     set_target_properties(${target} PROPERTIES
         VERSION ${DMDFE_VERSION}
         SOVERSION ${DMDFE_PATCH_VERSION}
         LINKER_LANGUAGE C
     )
+    if(RT_ARCHIVE_WITH_LDC AND NOT is_shared)
+        set_target_properties(${target} PROPERTIES LINKER_LANGUAGE D)
+    endif()
 
     # ldc2 defaults to position-independent code on Linux to match the implicit
     # linker default on Ubuntu 16.10 and above. As we might be building on an
@@ -462,7 +497,7 @@ macro(build_runtime_libs druntime_o druntime_bc phobos2_o phobos2_bc c_flags ld_
         COMPILE_FLAGS               "${c_flags}"
         LINK_FLAGS                  "${ld_flags}"
     )
-    set_common_library_properties(druntime-ldc${target_suffix})
+    set_common_library_properties(druntime-ldc${target_suffix} "${is_shared}")
 
     # When building a shared library, we need to link in all the default
     # libraries otherwise implicitly added by LDC to make it loadable from
@@ -485,7 +520,7 @@ macro(build_runtime_libs druntime_o druntime_bc phobos2_o phobos2_bc c_flags ld_
             COMPILE_FLAGS               "${c_flags}"
             LINK_FLAGS                  "${ld_flags}"
         )
-        set_common_library_properties(phobos2-ldc${target_suffix})
+        set_common_library_properties(phobos2-ldc${target_suffix} "${is_shared}")
 
         if("${is_shared}" STREQUAL "ON")
             target_link_libraries(phobos2-ldc${target_suffix}
@@ -599,7 +634,7 @@ macro(build_all_runtime_variants d_flags c_flags ld_flags path_suffix outlist_ta
         # static profile-rt
         build_profile_runtime("${d_flags};${D_FLAGS};${D_FLAGS_RELEASE}" "${c_flags}" "${ld_flags}" "" "${path_suffix}" ${outlist_targets})
         get_target_suffix("" "${path_suffix}" target_suffix)
-        set_common_library_properties(ldc-profile-rt${target_suffix})
+        set_common_library_properties(ldc-profile-rt${target_suffix} OFF)
     endif()
 endmacro()
 

runtime/PresetRuntimeConfiguration.cmake

@@ -0,0 +1,97 @@
+# - Preset cross-compilation configurations for C/ASM and D compilation
+# and linking, for supported targets
+#
+# This module sets compiler flags for a few C and assembly files in
+# DRuntime and Phobos, and linker flags to link the standard library as
+# a shared library and build the test runners, for various
+# cross-compilation targets that the LDC developers have tried out.
+#
+# It is enabled by setting LDC_TARGET_PRESET to a supported platform,
+# after which the appropriate TARGET_SYSTEM is set and a target triple
+# is appended to D_FLAGS.
+#
+# You can pass in custom RT_CFLAGS and LD_FLAGS of your choosing, but
+# if they're left unconfigured, they will also be set to sensible
+# defaults.
+
+if(NOT LDC_TARGET_PRESET STREQUAL "")
+    if(LDC_TARGET_PRESET MATCHES "Android")
+        set(ANDROID_API "21")
+    endif()
+    # This initial RT_CFLAGS/LD_FLAGS configuration for Android is a
+    # convenience for natively compiling, because CMake cannot detect
+    # Android as a separate platform from Linux.
+    if(RT_CFLAGS STREQUAL "" AND LDC_TARGET_PRESET MATCHES "Android")
+        set(RT_CFLAGS_UNCONFIGURED True)
+        set(RT_CFLAGS "-ffunction-sections -funwind-tables -fstack-protector-strong -Wno-invalid-command-line-argument -Wno-unused-command-line-argument -no-canonical-prefixes -g -DNDEBUG -DANDROID  -D__ANDROID_API__=${ANDROID_API} -Wa,--noexecstack -Wformat -Werror=format-security -fpie")
+
+        if(LDC_TARGET_PRESET MATCHES "arm")
+            append("-target armv7-none-linux-androideabi${ANDROID_API} -march=armv7-a -mfloat-abi=softfp -mfpu=vfpv3-d16 -mthumb -Os" RT_CFLAGS)
+        elseif(LDC_TARGET_PRESET MATCHES "aarch64")
+            append("-target aarch64-none-linux-android -O2" RT_CFLAGS)
+        endif()
+    endif()
+
+    if(LD_FLAGS STREQUAL "" AND LDC_TARGET_PRESET MATCHES "Android")
+        set(LD_FLAGS_UNCONFIGURED True)
+        set(LD_FLAGS "-Wl,--gc-sections -Wl,-z,nocopyreloc -no-canonical-prefixes -Wl,--no-undefined -Wl,-z,noexecstack -Wl,-z,relro -Wl,-z,now -Wl,--warn-shared-textrel -Wl,--fatal-warnings -fpie -pie")
+
+        if(LDC_TARGET_PRESET MATCHES "arm")
+            append("-target armv7-none-linux-androideabi${ANDROID_API} -Wl,--fix-cortex-a8" LD_FLAGS)
+        elseif(LDC_TARGET_PRESET MATCHES "aarch64")
+            append("-target aarch64-none-linux-android" LD_FLAGS)
+        endif()
+    endif()
+
+    if(LDC_TARGET_PRESET MATCHES "Windows")
+        set(TARGET_SYSTEM "Windows;MSVC")
+        if(LDC_TARGET_PRESET MATCHES "x64")
+            # stub example, fill in with the rest
+            list(APPEND D_FLAGS "-mtriple=x86_64-pc-windows-msvc")
+        endif()
+    elseif(LDC_TARGET_PRESET MATCHES "Android")
+        set(TARGET_SYSTEM "Android;Linux;UNIX")
+
+        # Check if we're using the NDK by looking for the toolchains
+        # directory in CC
+        if(CMAKE_C_COMPILER MATCHES "toolchains")
+            # Extract the NDK path and platform from CC
+            string(REGEX REPLACE ".toolchains.+" "" NDK_PATH ${CMAKE_C_COMPILER})
+            string(REGEX REPLACE ".+/prebuilt/([^/]+)/.+" "\\1" NDK_HOST_PLATFORM ${CMAKE_C_COMPILER})
+
+            if(LDC_TARGET_PRESET MATCHES "arm")
+                set(TARGET_ARCH "arm")
+                set(LLVM_TARGET_TRIPLE "armv7-none-linux-android")
+                set(TOOLCHAIN_TARGET_TRIPLE "arm-linux-androideabi")
+            elseif(LDC_TARGET_PRESET MATCHES "aarch64")
+                set(TARGET_ARCH "arm64")
+                set(LLVM_TARGET_TRIPLE "aarch64-none-linux-android")
+                set(TOOLCHAIN_TARGET_TRIPLE "aarch64-linux-android")
+            else()
+                message(FATAL_ERROR "Android platform ${LDC_TARGET_PRESET} is not supported.")
+            endif()
+            list(APPEND D_FLAGS "-mtriple=${LLVM_TARGET_TRIPLE}")
+            set(TOOLCHAIN_VERSION "4.9")
+
+            if(RT_CFLAGS_UNCONFIGURED)
+                append("-gcc-toolchain ${NDK_PATH}/toolchains/${TOOLCHAIN_TARGET_TRIPLE}-${TOOLCHAIN_VERSION}/prebuilt/${NDK_HOST_PLATFORM} --sysroot ${NDK_PATH}/sysroot -isystem ${NDK_PATH}/sysroot/usr/include/${TOOLCHAIN_TARGET_TRIPLE}" RT_CFLAGS)
+
+                if(LDC_TARGET_PRESET MATCHES "arm")
+                    append("-fno-integrated-as" RT_CFLAGS)
+                endif()
+            endif()
+
+            if(LD_FLAGS_UNCONFIGURED)
+                set(LD_BFD "-fuse-ld=bfd")
+                # work around Windows bug, android-ndk/ndk#75
+                if(NDK_HOST_PLATFORM MATCHES "windows")
+                    set(LD_BFD "${LD_BFD}.exe")
+                endif()
+
+                append("--sysroot=${NDK_PATH}/platforms/android-${ANDROID_API}/arch-${TARGET_ARCH} -gcc-toolchain ${NDK_PATH}/toolchains/${TOOLCHAIN_TARGET_TRIPLE}-${TOOLCHAIN_VERSION}/prebuilt/${NDK_HOST_PLATFORM} ${LD_BFD}" LD_FLAGS)
+            endif()
+        endif()
+    else()
+        message(FATAL_ERROR "LDC_TARGET_PRESET ${LDC_TARGET_PRESET} is not supported yet, pull requests to add common flags are welcome.")
+    endif()
+endif()

runtime/druntime

@@ -1 +1 @@
-Subproject commit 1b4a0579ac100106014cdcfcc323cf3f8023e166
+Subproject commit 4d3b8c17936089f542ddaee00db9abc34e8e260a

runtime/ldc-build-runtime.d.in

@@ -12,6 +12,8 @@ struct Config {
     string ldcSourceDir;
     bool ninja;
     bool buildTestrunners;
+    string targetPreset;
+    string[] targetSystem;
     string[] dFlags;
     string[] cFlags;
     string[] linkerFlags;
@@ -117,10 +119,10 @@ void prepareLdcSource() {
     removeVersionSuffix("git-");
     removeVersionSuffix("-dirty");
 
-    const localArchiveFile = "ldc-src.tar.gz";
+    import std.format : format;
+    const localArchiveFile = "ldc-src.zip";
     if (!localArchiveFile.exists) {
-        import std.format : format;
+        const url = "https://github.com/ldc-developers/ldc/releases/download/v%1$s/ldc-%1$s-src.zip".format(ldcVersion);
-        const url = "https://github.com/ldc-developers/ldc/releases/download/v%1$s/ldc-%1$s-src.tar.gz".format(ldcVersion);
         writefln("Downloading LDC source archive: %s", url);
         import std.net.curl : download;
         download(url, localArchiveFile);
@@ -133,28 +135,36 @@ void prepareLdcSource() {
         }
     }
 
-    if (!ldcSrc.exists)
+    extractZipArchive(localArchiveFile, ".");
-        mkdir(ldcSrc);
+    rename("ldc-%1$s-src".format(ldcVersion), ldcSrc);
-
-    import std.array : split;
-    exec("tar -xzf ldc-src.tar.gz --strip 1 -C ldc-src".split);
 }
 
 void runCMake() {
-    import std.array : byPair, join;
+    import std.array : empty, byPair, join;
+    import std.regex : matchFirst;
 
     const wd = WorkingDirScope(config.buildDir);
 
+    if(config.dFlags.empty)
+        config.dFlags ~= "-w";
+
+    if(config.targetSystem.empty)
+        config.targetSystem ~= "AUTO";
+
     string[] args = [
         "cmake",
         "-DLDC_EXE_FULL=" ~ config.ldcExecutable,
         "-DD_VERSION=@D_VERSION@",
         "-DDMDFE_MINOR_VERSION=@DMDFE_MINOR_VERSION@",
         "-DDMDFE_PATCH_VERSION=@DMDFE_PATCH_VERSION@",
+        "-DLDC_TARGET_PRESET=" ~ config.targetPreset,
+        "-DTARGET_SYSTEM=" ~ config.targetSystem.join(";"),
         "-DD_FLAGS=" ~ config.dFlags.join(";"),
         "-DRT_CFLAGS=" ~ config.cFlags.join(" "),
         "-DLD_FLAGS=" ~ config.linkerFlags.join(" "),
     ];
+    if(config.targetPreset.matchFirst("^Android"))
+        args ~= ["-DCMAKE_SYSTEM_NAME=Linux", "-DCMAKE_C_COMPILER_WORKS=True"];
     foreach (pair; config.cmakeVars.byPair)
         args ~= "-D" ~ pair[0] ~ '=' ~ pair[1];
     if (config.ninja)
@@ -200,6 +210,23 @@ void exec(string[] command ...) {
     }
 }
 
+void extractZipArchive(string archivePath, string destination) {
+    import std.string : endsWith;
+    import std.zip;
+
+    auto archive = new ZipArchive(std.file.read(archivePath));
+    foreach (name, am; archive.directory) {
+        const destPath = buildNormalizedPath(destination, name);
+
+        const isDir = name.endsWith("/");
+        const destDir = isDir ? destPath : destPath.dirName;
+        mkdirRecurse(destDir);
+
+        if (!isDir)
+            std.file.write(destPath, archive.expand(am));
+    }
+}
+
 void parseCommandLine(string[] args) {
     import std.getopt : arraySep, getopt, defaultGetoptPrinter;
 
@@ -212,6 +239,8 @@ void parseCommandLine(string[] args) {
         "ldcSrcDir",   "Path to LDC source directory (if not specified: downloads & extracts source archive into '<buildDir>/ldc-src')", &config.ldcSourceDir,
         "ninja",       "Use Ninja as CMake build system", &config.ninja,
         "testrunners", "Build the testrunner executables too", &config.buildTestrunners,
+        "targetPreset","Target configuration preset by LDC devs, e.g. Android-arm", &config.targetPreset,
+        "targetSystem","Target OS/toolchain (separated by ';'), e.g. Windows;MSVC", &config.targetSystem,
         "dFlags",      "LDC flags for the D modules (separated by ';')", &config.dFlags,
         "cFlags",      "C/ASM compiler flags for the handful of C/ASM files (separated by ';')", &config.cFlags,
         "linkerFlags", "C linker flags for shared libraries and testrunner executables (separated by ';')", &config.linkerFlags,
@@ -237,7 +266,7 @@ void parseCommandLine(string[] args) {
             "  * CMake\n" ~
             "  * either Make or Ninja (recommended, enable with '--ninja')\n" ~
             "  * C toolchain (compiler and linker)\n" ~
-            "  * tar (can be worked around via '--ldcSrcDir=path/to/src')\n" ~
+            "--targetPreset currently supports Android-arm or Android-aarch64.\n" ~
             "All arguments are optional.\n" ~
             "CMake variables (see runtime/CMakeLists.txt in LDC source) can be specified via arguments like 'VAR=value'.\n",
             helpInformation.options

tests/codegen/dcompute_cl_addrspaces.d

@@ -1,6 +1,6 @@
 // REQUIRES: target_SPIRV
-// RUN: %ldc -c -mdcompute-targets=ocl-220 -m64 -output-ll -output-o %s && FileCheck %s --check-prefix=LL < kernels_ocl220_64.ll \
+// RUN: %ldc -c -mdcompute-targets=ocl-220 -m64 -mdcompute-file-prefix=addrspace -output-ll -output-o %s && FileCheck %s --check-prefix=LL < addrspace_ocl220_64.ll \
-// RUN: && %llvm-spirv -to-text kernels_ocl220_64.spv && FileCheck %s --check-prefix=SPT < kernels_ocl220_64.spt
+// RUN: && %llvm-spirv -to-text addrspace_ocl220_64.spv && FileCheck %s --check-prefix=SPT < addrspace_ocl220_64.spt
 @compute(CompileFor.deviceOnly) module dcompute_cl_addrspaces;
 import ldc.dcompute;
 

tests/codegen/dcompute_cu_addrspaces.d

@@ -1,6 +1,6 @@
 // REQUIRES: atleast_llvm309
 // REQUIRES: target_NVPTX
-// RUN: %ldc -c -mdcompute-targets=cuda-350 -m64 -output-ll -output-o %s && FileCheck %s --check-prefix=LL < kernels_cuda350_64.ll && FileCheck %s --check-prefix=PTX < kernels_cuda350_64.ptx
+// RUN: %ldc -c -mdcompute-targets=cuda-350 -m64 -mdcompute-file-prefix=addrspace -output-ll -output-o %s && FileCheck %s --check-prefix=LL < addrspace_cuda350_64.ll && FileCheck %s --check-prefix=PTX < addrspace_cuda350_64.ptx
 @compute(CompileFor.deviceOnly) module dcompute_cu_addrspaces;
 import ldc.dcompute;
 

tests/codegen/dcompute_host_and_device.d

@@ -0,0 +1,16 @@
+// Check that we can generate code for both the host and device in one compiler invocation
+// REQUIRES: atleast_llvm309
+// REQUIRES: target_NVPTX
+// RUN: %ldc -mdcompute-targets=cuda-350 -mdcompute-file-prefix=host_and_device -Iinputs %s %S/inputs/kernel.d
+
+import inputs.kernel : foo;
+
+int tlGlobal;
+__gshared int gGlobal;
+
+void main(string[] args)
+{
+    tlGlobal = 0;
+    gGlobal  = 0;
+    string s = foo.mangleof;
+}

tests/codegen/gh2346.d

@@ -0,0 +1,69 @@
+// RUN: %ldc -output-ll -of=%t.ll %s && FileCheck %s < %t.ll
+
+// Make sure the LL struct is packed due to an unnatural overall alignment of 1.
+// CHECK-DAG: %gh2346.UnalignedUInt = type <{ i32 }>
+struct UnalignedUInt {
+    align(1) uint a;
+}
+static assert(UnalignedUInt.alignof == 1);
+static assert(UnalignedUInt.sizeof == 4);
+
+// Then there's no need to pack naturally aligned containers.
+// CHECK-DAG: %gh2346.Container = type { i8, %gh2346.UnalignedUInt }
+struct Container {
+    ubyte one;
+    UnalignedUInt two;
+}
+static assert(Container.alignof == 1);
+static assert(Container.sizeof == 5);
+static assert(Container.two.offsetof == 1);
+
+// CHECK-DAG: %gh2346.UnalignedUInt2 = type <{ i32 }>
+struct UnalignedUInt2 {
+    align(2) uint a;
+}
+static assert(UnalignedUInt2.alignof == 2);
+static assert(UnalignedUInt2.sizeof == 4);
+
+// CHECK-DAG: %gh2346.Container2 = type { i8, [1 x i8], %gh2346.UnalignedUInt2 }
+struct Container2 {
+    ubyte one;
+    UnalignedUInt2 two;
+}
+static assert(Container2.alignof == 2);
+static assert(Container2.sizeof == 6);
+static assert(Container2.two.offsetof == 2);
+
+// CHECK-DAG: %gh2346.PackedContainer2 = type <{ i8, %gh2346.UnalignedUInt2 }>
+struct PackedContainer2 {
+    ubyte one;
+    align(1) UnalignedUInt2 two;
+}
+static assert(PackedContainer2.alignof == 1);
+static assert(PackedContainer2.sizeof == 5);
+static assert(PackedContainer2.two.offsetof == 1);
+
+// CHECK-DAG: %gh2346.WeirdContainer = type { i8, [1 x i8], %gh2346.UnalignedUInt, [2 x i8] }
+align(4) struct WeirdContainer {
+    ubyte one;
+    align(2) UnalignedUInt two;
+}
+static assert(WeirdContainer.alignof == 4);
+static assert(WeirdContainer.sizeof == 8);
+static assert(WeirdContainer.two.offsetof == 2);
+
+// CHECK-DAG: %gh2346.ExplicitlyUnalignedUInt2 = type <{ i32 }>
+align(2) struct ExplicitlyUnalignedUInt2 {
+    uint a;
+}
+static assert(ExplicitlyUnalignedUInt2.alignof == 2);
+static assert(ExplicitlyUnalignedUInt2.sizeof == 4);
+
+// CHECK-DAG: %gh2346.AnotherContainer = type { i8, [1 x i8], %gh2346.ExplicitlyUnalignedUInt2 }
+struct AnotherContainer {
+    ubyte one;
+    ExplicitlyUnalignedUInt2 two;
+}
+static assert(AnotherContainer.alignof == 2);
+static assert(AnotherContainer.sizeof == 6);
+static assert(AnotherContainer.two.offsetof == 2);

tests/codegen/inputs/kernel.d

@@ -0,0 +1,5 @@
+@compute(CompileFor.deviceOnly)
+module inputs.kernel;
+
+import ldc.dcompute;
+@kernel void foo() {}

tests/codegen/nested_lazy_gh2302.d

@@ -0,0 +1,18 @@
+// RUN: %ldc -output-ll -of=%t.ll %s && FileCheck %s < %t.ll
+// RUN: %ldc -run %s
+
+void foo()
+{
+    auto impl(T)(lazy T field)
+    {
+        // Make sure `field` is a closure variable with delegate type (LL struct).
+        // CHECK: %nest.impl = type { { i8*, i32 (i8*)* } }
+        auto ff() { return field; }
+        auto a = field;
+        return ff() + a;
+    }
+    auto r = impl(123);
+    assert(r == 246);
+}
+
+void main() { foo(); }

tests/compilable/dcompute.d

@@ -1,13 +1,13 @@
 // Tests that
-//  - we dont try to link with one file on the commandline that is @compute
+//  - we don't try to link with one file on the commandline that is @compute
-//  - truning on debugging doesn't ICE
+//  - turning on debugging doesn't ICE
-//  - dont analyse uninstansiated templates
+//  - don't analyse uninstantiated templates
 //  - typeid generated for hashing of struct (typeid(const(T))) is ignored and does not error
 
 // REQUIRES: target_NVPTX
-// RUN: %ldc -mdcompute-targets=cuda-350 -g
+// RUN: %ldc -mdcompute-targets=cuda-350 -g %s
 
-@compute(Compilefor.deviceOnly) module dcompute;
+@compute(CompileFor.deviceOnly) module dcompute;
 import ldc.dcompute;
 
 @kernel void foo()

tests/debuginfo/args_cdb.d

@@ -13,17 +13,21 @@ import core.simd;
 // enable case sensitive symbol lookup
 // CDB: .symopt-1
 
+struct Small { size_t val; }
+struct Large { size_t a, b, c, d, e, f, g, h; }
+
 int byValue(ubyte ub, ushort us, uint ui, ulong ul,
             float f, double d, real r,
             cdouble c, int delegate() dg, int function() fun,
             int[] slice, float[int] aa, ubyte[16] fa,
             float4 f4, double4 d4,
-            Interface ifc, TypeInfo_Class ti, typeof(null) np)
+            Small small, Large large,
+            TypeInfo_Class ti, typeof(null) np)
 {
-// CDB: bp `args_cdb.d:25`
+// CDB: bp `args_cdb.d:27`
 // CDB: g
-    // arguments implicitly passed by reference on x64 and not shown if unused
+    // arguments implicitly passed by reference aren't shown if unused
-    float cim = c.im + fa[7] + dg() + ifc.offset;
+    float cim = c.im + fa[7] + dg() + small.val + large.a;
     return 1;
 // CHECK: !args_cdb.byValue
 // CDB: dv /t
@@ -42,12 +46,13 @@ int byValue(ubyte ub, ushort us, uint ui, ulong ul,
 // CHECK: <function> * fun = {{0x[0-9a-f`]*}}
 // x86: struct int[] slice =
 // CHECK: unsigned char * aa = {{0x[0-9a-f`]*}}
-// "Internal implementation error for fa" with cdb on x64, ok in VS
+// x64: unsigned char (*)[16] fa
 // x86: unsigned char [16] fa
 // x86: float [4] f4 = float [4]
 // x86: double [4] d4 = double [4]
-// x64: Interface * ifc
+// CHECK: Small small
-// x86: Interface ifc
+// x64: Large * large
+// x86: Large large
 // CHECK: struct TypeInfo_Class * ti = {{0x[0-9a-f`]*}}
 // CHECK: void * np = {{0x[0`]*}}
 
@@ -64,31 +69,35 @@ int byValue(ubyte ub, ushort us, uint ui, ulong ul,
 
 // CDB: ?? dg
 // CHECK: int delegate()
-// CHECK-NEXT: context
+// CHECK-NEXT: context :
-// CHECK-NEXT: funcptr
+// CHECK-NEXT: funcptr :
 // CHECK-SAME: args_cdb.main.__lambda
 
 // CDB: ?? slice
 // CHECK: struct int[]
 // CHECK-NEXT: length : 2
-// CHECK-NEXT: ptr
+// CHECK-NEXT: ptr :
+// CHECK-SAME: 0n10
 
-// CDB: ?? fa[1]
+// CDB: ?? (*fa)[1]
-// "Internal implementation error for fa" with cdb on x64, ok in VS
+// x64: unsigned char 0x0e
-// no-x86: unsigned char 0x0e (displays 0xf6)
+// no-x86: would be fa[1], but displays garbage anyway
 
 // CDB: ?? f4[1]
 // CHECK: float 16
 
-// CDB: ?? d4[2]
+// CDB: ?? d4[1]
 // CHECK: double 17
 
-// CDB: ?? ifc
+// CDB: ?? small
-// CHECK: Interface
+// CHECK: Small
-// CHECK-NEXT: classinfo
+// x64-NEXT: val : 0x12
-// CHECK-NEXT: vtbl
+// no-x86-NEXT: val : 0x12 (displays garbage)
-// x64-NEXT: offset : 0x12
+
-// no-x86-NEXT: offset : 0x12 (displays 0)
+// CDB: ?? large
+// CHECK: Large
+// x64-NEXT: a : 0x13
+// no-x86-NEXT: a : 0x13 (displays garbage)
 
 // CDB: ?? ti
 // CHECK: TypeInfo_Class
@@ -100,9 +109,10 @@ int byPtr(ubyte* ub, ushort* us, uint* ui, ulong* ul,
           cdouble* c, int delegate()* dg, int function()* fun,
           int[]* slice, float[int]* aa, ubyte[16]* fa,
           float4* f4, double4* d4,
-          Interface* ifc, TypeInfo_Class* ti, typeof(null)* np)
+          Small* small, Large* large,
+          TypeInfo_Class* ti, typeof(null)* np)
 {
-// CDB: bp `args_cdb.d:106`
+// CDB: bp `args_cdb.d:115`
 // CDB: g
     return 3;
 // CHECK: !args_cdb.byPtr
@@ -127,8 +137,8 @@ int byPtr(ubyte* ub, ushort* us, uint* ui, ulong* ul,
 // CHECK-NEXT: im : 9
 // CDB: ?? *dg
 // CHECK: int delegate()
-// CHECK-NEXT: context
+// CHECK-NEXT: context :
-// CHECK-NEXT: funcptr
+// CHECK-NEXT: funcptr :
 // CHECK-SAME: args_cdb.main.__lambda
 // CDB: ?? *fun
 // CHECK: <function> *
@@ -145,13 +155,16 @@ int byPtr(ubyte* ub, ushort* us, uint* ui, ulong* ul,
 // CHECK: float 16
 // CDB: ?? (*d4)[2]
 // CHECK: double 17
-// CDB: ?? *ifc
+// CDB: ?? *small
-// CHECK: struct Interface
+// CHECK: struct Small
-// CHECK: offset : 0x12
+// CHECK-NEXT: val : 0x12
+// CDB: ?? *large
+// CHECK: struct Large
+// CHECK-NEXT: a : 0x13
+// CHECK-NEXT: b :
 // CDB: ?? *ti
 // CHECK: struct TypeInfo_Class
 // CHECK-NEXT: m_init : byte[]
-// shows bad member values
 // CDB: ?? *np
 // CHECK: void * {{0x[0`]*}}
 }
@@ -161,9 +174,10 @@ int byRef(ref ubyte ub, ref ushort us, ref uint ui, ref ulong ul,
           ref cdouble c, ref int delegate() dg, ref int function() fun,
           ref int[] slice, ref float[int] aa, ref ubyte[16] fa,
           ref float4 f4, ref double4 d4,
-          ref Interface ifc, ref TypeInfo_Class ti, ref typeof(null) np)
+          ref Small small, ref Large large,
+          ref TypeInfo_Class ti, ref typeof(null) np)
 {
-// CDB: bp `args_cdb.d:218`
+// CDB: bp `args_cdb.d:180`
 // CDB: g
 // CHECK: !args_cdb.byRef
 
@@ -189,7 +203,7 @@ int byRef(ref ubyte ub, ref ushort us, ref uint ui, ref ulong ul,
 // CHECK-NEXT: im : 9
 // CDB: ?? *dg
 // CHECK: int delegate()
-// CHECK-NEXT: context
+// CHECK-NEXT: context :
 // CHECK-NEXT: funcptr : {{0x[0-9a-f`]*}}
 // CHECK-SAME: args_cdb.main.__lambda
 // CDB: ?? *fun
@@ -206,16 +220,23 @@ int byRef(ref ubyte ub, ref ushort us, ref uint ui, ref ulong ul,
 // CHECK: float 16
 // CDB: ?? (*d4)[2]
 // CHECK: double 17
-// CDB: ?? *ifc
+// CDB: ?? *small
-// CHECK: struct Interface
+// CHECK: struct Small
-// CHECK: offset : 0x12
+// CHECK-NEXT: val : 0x12
+// CDB: ?? *large
+// CHECK: struct Large
+// CHECK-NEXT: a : 0x13
+// CHECK-NEXT: b :
 // CDB: ?? *ti
 // CHECK: struct TypeInfo_Class * {{0x[0-9a-f`]*}}
 // CHECK-NEXT: m_init : byte[]
 // CDB: ?? *np
-// CHECK: void * {{0x[0`]*}}
+// no-CHECK: void * {{0x[0`]*}} (not available)
+
+    // needs access to references to actually generate debug info
+    float cim = c.im + fa[7] + dg() + fun() + slice.length + aa.length + f4[0] + d4[1] +
+                small.val + large.a + ti.initializer.length;
 
-// needs access to references to actually generate debug info
     ub++;
     us++;
     ui++;
@@ -231,7 +252,8 @@ int byRef(ref ubyte ub, ref ushort us, ref uint ui, ref ulong ul,
     fa[0]++;
     f4.array[0] = f4.array[0] + 1;
     d4.array[0] = d4.array[0] + 1;
-    ifc = Interface(typeid(Object), null, 18);
+    small.val++;
+    large.a++;
     ti = typeid(TypeInfo);
     np = null;
     return 2;
@@ -254,13 +276,14 @@ int main()
     ubyte[16] fa; fa[] = 14;
     float4 f4 = 16;
     double4 d4 = 17;
-    Interface ifc = Interface(typeid(Object), null, 18);
+    Small small = Small(18);
+    Large large = Large(19);
     TypeInfo_Class ti = typeid(TypeInfo);
     typeof(null) np = null;
     
-    byValue(ub, us, ui, ul, f, d, r, c, dg, fun, slice, aa, fa, f4, d4, ifc, ti, np);
+    byValue(ub, us, ui, ul, f, d, r, c, dg, fun, slice, aa, fa, f4, d4, small, large, ti, np);
-    byPtr(&ub, &us, &ui, &ul, &f, &d, &r, &c, &dg, &fun, &slice, &aa, &fa, &f4, &d4, &ifc, &ti, &np);
+    byPtr(&ub, &us, &ui, &ul, &f, &d, &r, &c, &dg, &fun, &slice, &aa, &fa, &f4, &d4, &small, &large, &ti, &np);
-    byRef(ub, us, ui, ul, f, d, r, c, dg, fun, slice, aa, fa, f4, d4, ifc, ti, np);
+    byRef(ub, us, ui, ul, f, d, r, c, dg, fun, slice, aa, fa, f4, d4, small, large, ti, np);
 
     return 0;
 }

tests/debuginfo/nested_cdb.d

@@ -0,0 +1,67 @@
+// REQUIRES: atleast_llvm500
+// REQUIRES: Windows
+// REQUIRES: cdb
+// RUN: %ldc -g -of=%t.exe %s
+// RUN: sed -e "/^\\/\\/ CDB:/!d" -e "s,// CDB:,," %s \
+// RUN:    | %cdb -snul -lines -y . %t.exe >%t.out
+// RUN: FileCheck %s -check-prefix=CHECK -check-prefix=%arch < %t.out
+
+// CDB: ld /f nested_cdb*
+// enable case sensitive symbol lookup
+// CDB: .symopt-1
+
+void encloser(int arg0, ref int arg1)
+{
+    int enc_n = 123;
+// CDB: bp `nested_cdb.d:16`
+// CDB: g
+// CDB: dv /t
+// CHECK: int arg0 = 0n1
+// (cdb displays references as pointers)
+// CHECK-NEXT: int * arg1 = {{0x[0-9a-f`]*}}
+// CHECK-NEXT: int enc_n = 0n123
+// CDB: ?? *arg1
+// CHECK: int 0n2
+    enc_n += arg1;
+
+    void nested(int nes_i)
+    {
+        int blub = arg0 + arg1 + enc_n;
+// CDB: bp `nested_cdb.d:30`
+// CDB: g
+// CDB: dv /t
+// CHECK: int arg0 = 0n1
+// CHECK-NEXT: int * arg1 = {{0x[0-9a-f`]*}}
+// CHECK-NEXT: int enc_n = 0n125
+// CDB: ?? *arg1
+// CHECK: int 0n2
+        arg0 = arg1 = enc_n = nes_i;
+// CDB: bp `nested_cdb.d:39`
+// CDB: g
+// CDB: dv /t
+// CHECK: int arg0 = 0n456
+// CHECK-NEXT: int * arg1 = {{0x[0-9a-f`]*}}
+// CHECK-NEXT: int enc_n = 0n456
+// CDB: ?? *arg1
+// CHECK: int 0n456
+    }
+
+    nested(456);
+// CDB: bp `nested_cdb.d:50`
+// CDB: g
+// CDB: dv /t
+// CHECK: int arg0 = 0n456
+// CHECK-NEXT: int * arg1 = {{0x[0-9a-f`]*}}
+// CHECK-NEXT: int enc_n = 0n456
+// CDB: ?? *arg1
+// CHECK: int 0n456
+}
+
+void main()
+{
+    int arg1 = 2;
+    encloser(1, arg1);
+}
+
+// CDB: q
+// CHECK: quit

tests/debuginfo/nested_gdb.d

@@ -0,0 +1,61 @@
+// REQUIRES: gdb
+// RUN: %ldc %_gdb_dflags -g -of=%t %s
+// RUN: sed -e "/^\\/\\/ GDB:/!d" -e "s,// GDB:,," %s >%t.gdb
+// RUN: gdb %t --batch -x %t.gdb >%t.out 2>&1
+// RUN: FileCheck %s -check-prefix=CHECK < %t.out
+
+void encloser(int arg0, ref int arg1)
+{
+    int enc_n = 123;
+// GDB: b 10
+// GDB: r
+// GDB: p arg0
+// CHECK: $1 = 1
+// GDB: p arg1
+// CHECK: $2 = (int &) @{{0x[0-9a-f]*}}: 2
+// GDB: p enc_n
+// CHECK: $3 = 123
+    enc_n += arg1;
+
+    void nested(int nes_i)
+    {
+        int blub = arg0 + arg1 + enc_n;
+// GDB: b 23
+// GDB: c
+// GDB: p arg0
+// CHECK: $4 = 1
+// GDB: p arg1
+// CHECK: $5 = (int &) @{{0x[0-9a-f]*}}: 2
+// GDB: p enc_n
+// CHECK: $6 = 125
+        arg0 = arg1 = enc_n = nes_i;
+// GDB: b 32
+// GDB: c
+// GDB: p arg0
+// CHECK: $7 = 456
+// GDB: p arg1
+// CHECK: $8 = (int &) @{{0x[0-9a-f]*}}: 456
+// GDB: p enc_n
+// CHECK: $9 = 456
+    }
+
+    nested(456);
+// GDB: b 43
+// GDB: c
+// GDB: p arg0
+// no-CHECK: $10 = 456 (`<optimized out>` for LLVM < 5.0)
+// GDB: p arg1
+// CHECK: $11 = (int &) @{{0x[0-9a-f]*}}: 456
+// GDB: p enc_n
+// CHECK: $12 = 456
+}
+
+void main()
+{
+    int arg1 = 2;
+    encloser(1, arg1);
+}
+
+// GDB: c
+// GDB: q
+// CHECK: exited normally

tests/lit.site.cfg.in

@@ -81,7 +81,8 @@ if (platform.system() == 'Windows') and (config.default_target_bits == 32):
 if (platform.system() == 'Windows') and (config.default_target_bits == 64):
     config.available_features.add('Windows_x64')
 
-# Add "LTO" feature if linker support is available (LTO is supported from LLVM 3.9)
+# Add "LTO" feature if linker support and LTO plugin are available
+# (LTO is supported from LLVM 3.9)
 canDoLTO = False
 if (config.llvm_version >= 309):
     if (platform.system() == 'Darwin'):
@@ -89,14 +90,14 @@ if (config.llvm_version >= 309):
         p = subprocess.Popen(command, stdout=subprocess.PIPE,
                 stderr=subprocess.PIPE, universal_newlines=True)
         text = p.stderr.read()
-        if "LTO support" in text:
+        if ("LTO support" in text and os.path.isfile(config.ldc2_lib_dir + '/libLTO-ldc.dylib')):
             canDoLTO = True
     elif (platform.system() == 'Linux'):
         command = ['ld', '-plugin']
         p = subprocess.Popen(command, stdout=subprocess.PIPE,
                 stderr=subprocess.PIPE, universal_newlines=True)
         text = p.stderr.read()
-        if "plugin: missing argument" in text:
+        if ("plugin: missing argument" in text and os.path.isfile(config.ldc2_lib_dir + '/LLVMgold-ldc.so')):
             canDoLTO = True
 if canDoLTO:
     config.available_features.add('LTO')

utils/not.cpp

@@ -40,8 +40,13 @@ int main(int argc, const char **argv) {
   }
 
   std::string ErrMsg;
-  int Result = sys::ExecuteAndWait(*Program, argv, nullptr, nullptr, 0, 0,
+  int Result = sys::ExecuteAndWait(*Program, argv, nullptr,
-                                   &ErrMsg);
+#if LDC_LLVM_VER >= 600
+                                   {},
+#else
+                                   nullptr,
+#endif
+                                   0, 0, &ErrMsg);
 #ifdef _WIN32
   // Handle abort() in msvcrt -- It has exit code as 3.  abort(), aka
   // unreachable, should be recognized as a crash.  However, some binaries use