clang  16.0.0git
Cuda.h
Go to the documentation of this file.
1 //===--- Cuda.h - Cuda ToolChain Implementations ----------------*- C++ -*-===//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 //===----------------------------------------------------------------------===//
8 
9 #ifndef LLVM_CLANG_LIB_DRIVER_TOOLCHAINS_CUDA_H
10 #define LLVM_CLANG_LIB_DRIVER_TOOLCHAINS_CUDA_H
11 
12 #include "clang/Basic/Cuda.h"
13 #include "clang/Driver/Action.h"
14 #include "clang/Driver/Multilib.h"
15 #include "clang/Driver/Tool.h"
16 #include "clang/Driver/ToolChain.h"
17 #include "llvm/ADT/Optional.h"
18 #include "llvm/Support/Compiler.h"
19 #include "llvm/Support/VersionTuple.h"
20 #include <bitset>
21 #include <set>
22 #include <vector>
23 
24 namespace clang {
25 namespace driver {
26 
27 /// A class to find a viable CUDA installation
29 private:
30  const Driver &D;
31  bool IsValid = false;
33  std::string InstallPath;
34  std::string BinPath;
35  std::string LibPath;
36  std::string LibDevicePath;
37  std::string IncludePath;
38  llvm::StringMap<std::string> LibDeviceMap;
39 
40  // CUDA architectures for which we have raised an error in
41  // CheckCudaVersionSupportsArch.
42  mutable std::bitset<(int)CudaArch::LAST> ArchsWithBadVersion;
43 
44 public:
45  CudaInstallationDetector(const Driver &D, const llvm::Triple &HostTriple,
46  const llvm::opt::ArgList &Args);
47 
48  void AddCudaIncludeArgs(const llvm::opt::ArgList &DriverArgs,
49  llvm::opt::ArgStringList &CC1Args) const;
50 
51  /// Emit an error if Version does not support the given Arch.
52  ///
53  /// If either Version or Arch is unknown, does not emit an error. Emits at
54  /// most one error per Arch.
55  void CheckCudaVersionSupportsArch(CudaArch Arch) const;
56 
57  /// Check whether we detected a valid Cuda install.
58  bool isValid() const { return IsValid; }
59  /// Print information about the detected CUDA installation.
60  void print(raw_ostream &OS) const;
61 
62  /// Get the detected Cuda install's version.
63  CudaVersion version() const {
65  : Version;
66  }
67  /// Get the detected Cuda installation path.
68  StringRef getInstallPath() const { return InstallPath; }
69  /// Get the detected path to Cuda's bin directory.
70  StringRef getBinPath() const { return BinPath; }
71  /// Get the detected Cuda Include path.
72  StringRef getIncludePath() const { return IncludePath; }
73  /// Get the detected Cuda library path.
74  StringRef getLibPath() const { return LibPath; }
75  /// Get the detected Cuda device library path.
76  StringRef getLibDevicePath() const { return LibDevicePath; }
77  /// Get libdevice file for given architecture
78  std::string getLibDeviceFile(StringRef Gpu) const {
79  return LibDeviceMap.lookup(Gpu);
80  }
82 };
83 
84 namespace tools {
85 namespace NVPTX {
86 
87 // Run ptxas, the NVPTX assembler.
88 class LLVM_LIBRARY_VISIBILITY Assembler : public Tool {
89  public:
90  Assembler(const ToolChain &TC) : Tool("NVPTX::Assembler", "ptxas", TC) {}
91 
92  bool hasIntegratedCPP() const override { return false; }
93 
94  void ConstructJob(Compilation &C, const JobAction &JA,
95  const InputInfo &Output, const InputInfoList &Inputs,
96  const llvm::opt::ArgList &TCArgs,
97  const char *LinkingOutput) const override;
98 };
99 
100 // Runs fatbinary, which combines GPU object files ("cubin" files) and/or PTX
101 // assembly into a single output file.
102 class LLVM_LIBRARY_VISIBILITY Linker : public Tool {
103  public:
104  Linker(const ToolChain &TC) : Tool("NVPTX::Linker", "fatbinary", TC) {}
105 
106  bool hasIntegratedCPP() const override { return false; }
107 
108  void ConstructJob(Compilation &C, const JobAction &JA,
109  const InputInfo &Output, const InputInfoList &Inputs,
110  const llvm::opt::ArgList &TCArgs,
111  const char *LinkingOutput) const override;
112 };
113 
114 void getNVPTXTargetFeatures(const Driver &D, const llvm::Triple &Triple,
115  const llvm::opt::ArgList &Args,
116  std::vector<StringRef> &Features);
117 
118 } // end namespace NVPTX
119 } // end namespace tools
120 
121 namespace toolchains {
122 
123 class LLVM_LIBRARY_VISIBILITY CudaToolChain : public ToolChain {
124 public:
125  CudaToolChain(const Driver &D, const llvm::Triple &Triple,
126  const ToolChain &HostTC, const llvm::opt::ArgList &Args);
127 
128  const llvm::Triple *getAuxTriple() const override {
129  return &HostTC.getTriple();
130  }
131 
132  std::string getInputFilename(const InputInfo &Input) const override;
133 
134  llvm::opt::DerivedArgList *
135  TranslateArgs(const llvm::opt::DerivedArgList &Args, StringRef BoundArch,
136  Action::OffloadKind DeviceOffloadKind) const override;
137  void addClangTargetOptions(const llvm::opt::ArgList &DriverArgs,
138  llvm::opt::ArgStringList &CC1Args,
139  Action::OffloadKind DeviceOffloadKind) const override;
140 
141  llvm::DenormalMode getDefaultDenormalModeForType(
142  const llvm::opt::ArgList &DriverArgs, const JobAction &JA,
143  const llvm::fltSemantics *FPType = nullptr) const override;
144 
145  // Never try to use the integrated assembler with CUDA; always fork out to
146  // ptxas.
147  bool useIntegratedAs() const override { return false; }
148  bool isCrossCompiling() const override { return true; }
149  bool isPICDefault() const override { return false; }
150  bool isPIEDefault(const llvm::opt::ArgList &Args) const override {
151  return false;
152  }
153  bool isPICDefaultForced() const override { return false; }
154  bool SupportsProfiling() const override { return false; }
155  bool supportsDebugInfoOption(const llvm::opt::Arg *A) const override;
156  void adjustDebugInfoKind(codegenoptions::DebugInfoKind &DebugInfoKind,
157  const llvm::opt::ArgList &Args) const override;
158  bool IsMathErrnoDefault() const override { return false; }
159 
160  void AddCudaIncludeArgs(const llvm::opt::ArgList &DriverArgs,
161  llvm::opt::ArgStringList &CC1Args) const override;
162 
163  void addClangWarningOptions(llvm::opt::ArgStringList &CC1Args) const override;
164  CXXStdlibType GetCXXStdlibType(const llvm::opt::ArgList &Args) const override;
165  void
166  AddClangSystemIncludeArgs(const llvm::opt::ArgList &DriverArgs,
167  llvm::opt::ArgStringList &CC1Args) const override;
168  void AddClangCXXStdlibIncludeArgs(
169  const llvm::opt::ArgList &Args,
170  llvm::opt::ArgStringList &CC1Args) const override;
171  void AddIAMCUIncludeArgs(const llvm::opt::ArgList &DriverArgs,
172  llvm::opt::ArgStringList &CC1Args) const override;
173 
174  SanitizerMask getSupportedSanitizers() const override;
175 
176  VersionTuple
177  computeMSVCVersion(const Driver *D,
178  const llvm::opt::ArgList &Args) const override;
179 
180  unsigned GetDefaultDwarfVersion() const override { return 2; }
181  // NVPTX supports only DWARF2.
182  unsigned getMaxDwarfVersion() const override { return 2; }
183 
186 
187 protected:
188  Tool *buildAssembler() const override; // ptxas
189  Tool *buildLinker() const override; // fatbinary (ok, not really a linker)
190 };
191 
192 } // end namespace toolchains
193 } // end namespace driver
194 } // end namespace clang
195 
196 #endif // LLVM_CLANG_LIB_DRIVER_TOOLCHAINS_CUDA_H
clang::driver::toolchains::CudaToolChain::getMaxDwarfVersion
unsigned getMaxDwarfVersion() const override
Definition: Cuda.h:182
clang::driver::tools::NVPTX::getNVPTXTargetFeatures
void getNVPTXTargetFeatures(const Driver &D, const llvm::Triple &Triple, const llvm::opt::ArgList &Args, std::vector< StringRef > &Features)
Definition: Cuda.cpp:561
clang::CudaVersion::NEW
@ NEW
string
string(SUBSTRING ${CMAKE_CURRENT_BINARY_DIR} 0 ${PATH_LIB_START} PATH_HEAD) string(SUBSTRING $
Definition: CMakeLists.txt:22
clang::driver::CudaInstallationDetector::getInstallPath
StringRef getInstallPath() const
Get the detected Cuda installation path.
Definition: Cuda.h:68
clang::driver::tools::NVPTX::Assembler
Definition: Cuda.h:88
llvm::SmallVector
Definition: LLVM.h:38
clang::driver::toolchains::CudaToolChain
Definition: Cuda.h:123
clang::driver::CudaInstallationDetector::isValid
bool isValid() const
Check whether we detected a valid Cuda install.
Definition: Cuda.h:58
int
__device__ int
Definition: __clang_hip_libdevice_declares.h:63
clang::driver::toolchains::CudaToolChain::getAuxTriple
const llvm::Triple * getAuxTriple() const override
Get the toolchain's aux triple, if it has one.
Definition: Cuda.h:128
clang::driver::tools::NVPTX::Assembler::hasIntegratedCPP
bool hasIntegratedCPP() const override
Definition: Cuda.h:92
clang::driver::toolchains::CudaToolChain::isCrossCompiling
bool isCrossCompiling() const override
Returns true if the toolchain is targeting a non-native architecture.
Definition: Cuda.h:148
clang::driver::InputInfo
InputInfo - Wrapper for information about an input source.
Definition: InputInfo.h:22
Action.h
clang::driver::toolchains::CudaToolChain::CudaInstallation
CudaInstallationDetector CudaInstallation
Definition: Cuda.h:185
clang::driver::Tool
Tool - Information on a specific compilation tool.
Definition: Tool.h:32
clang::driver::CudaInstallationDetector::getLibDevicePath
StringRef getLibDevicePath() const
Get the detected Cuda device library path.
Definition: Cuda.h:76
clang::driver::toolchains::CudaToolChain::HostTC
const ToolChain & HostTC
Definition: Cuda.h:184
clang::CudaVersion
CudaVersion
Definition: Cuda.h:20
clang::driver::toolchains::CudaToolChain::IsMathErrnoDefault
bool IsMathErrnoDefault() const override
IsMathErrnoDefault - Does this tool chain use -fmath-errno by default.
Definition: Cuda.h:158
clang::driver::CudaInstallationDetector::getLibPath
StringRef getLibPath() const
Get the detected Cuda library path.
Definition: Cuda.h:74
clang::driver::CudaInstallationDetector::getBinPath
StringRef getBinPath() const
Get the detected path to Cuda's bin directory.
Definition: Cuda.h:70
clang::driver::tools::NVPTX::Linker::hasIntegratedCPP
bool hasIntegratedCPP() const override
Definition: Cuda.h:106
Tool.h
clang::driver::ToolChain::getTriple
const llvm::Triple & getTriple() const
Definition: ToolChain.h:234
clang::driver::CudaInstallationDetector::getIncludePath
StringRef getIncludePath() const
Get the detected Cuda Include path.
Definition: Cuda.h:72
clang::driver::CudaInstallationDetector::getLibDeviceFile
std::string getLibDeviceFile(StringRef Gpu) const
Get libdevice file for given architecture.
Definition: Cuda.h:78
ToolChain.h
clang::CudaArch::LAST
@ LAST
clang::driver::tools::NVPTX::Assembler::Assembler
Assembler(const ToolChain &TC)
Definition: Cuda.h:90
clang::driver::toolchains::CudaToolChain::isPICDefaultForced
bool isPICDefaultForced() const override
Tests whether this toolchain forces its default for PIC, PIE or non-PIC.
Definition: Cuda.h:153
clang::driver::ToolChain
ToolChain - Access to tools for a single platform.
Definition: ToolChain.h:91
clang::driver::CudaInstallationDetector::WarnIfUnsupportedVersion
void WarnIfUnsupportedVersion()
Definition: Cuda.cpp:108
clang::CudaVersion::PARTIALLY_SUPPORTED
@ PARTIALLY_SUPPORTED
clang::driver::toolchains::CudaToolChain::isPICDefault
bool isPICDefault() const override
Test whether this toolchain defaults to PIC.
Definition: Cuda.h:149
clang::driver::tools::NVPTX::Linker
Definition: Cuda.h:102
clang::driver::toolchains::CudaToolChain::useIntegratedAs
bool useIntegratedAs() const override
Check if the toolchain should use the integrated assembler.
Definition: Cuda.h:147
clang::driver::tools::NVPTX::Linker::Linker
Linker(const ToolChain &TC)
Definition: Cuda.h:104
clang::driver::CudaInstallationDetector::AddCudaIncludeArgs
void AddCudaIncludeArgs(const llvm::opt::ArgList &DriverArgs, llvm::opt::ArgStringList &CC1Args) const
Definition: Cuda.cpp:290
clang::driver::CudaInstallationDetector
A class to find a viable CUDA installation.
Definition: Cuda.h:28
clang::driver::Compilation
Compilation - A set of tasks to perform for a single driver invocation.
Definition: Compilation.h:45
Cuda.h
clang::driver::CudaInstallationDetector::CudaInstallationDetector
CudaInstallationDetector(const Driver &D, const llvm::Triple &HostTriple, const llvm::opt::ArgList &Args)
Definition: Cuda.cpp:122
clang
Definition: CalledOnceCheck.h:17
clang::driver::Driver
Driver - Encapsulate logic for constructing compilation processes from a set of gcc-driver-like comma...
Definition: Driver.h:75
clang::driver::toolchains::CudaToolChain::isPIEDefault
bool isPIEDefault(const llvm::opt::ArgList &Args) const override
Test whether this toolchain defaults to PIE.
Definition: Cuda.h:150
clang::driver::CudaInstallationDetector::version
CudaVersion version() const
Get the detected Cuda install's version.
Definition: Cuda.h:63
clang::CudaArch
CudaArch
Definition: Cuda.h:49
Multilib.h
clang::codegenoptions::DebugInfoKind
DebugInfoKind
Definition: DebugInfoOptions.h:20
clang::driver::Action::OffloadKind
OffloadKind
Definition: Action.h:86
clang::driver::CudaInstallationDetector::print
void print(raw_ostream &OS) const
Print information about the detected CUDA installation.
Definition: Cuda.cpp:331
clang::driver::toolchains::CudaToolChain::GetDefaultDwarfVersion
unsigned GetDefaultDwarfVersion() const override
Definition: Cuda.h:180
clang::driver::JobAction
Definition: Action.h:398
clang::driver::CudaInstallationDetector::CheckCudaVersionSupportsArch
void CheckCudaVersionSupportsArch(CudaArch Arch) const
Emit an error if Version does not support the given Arch.
Definition: Cuda.cpp:314
clang::CudaVersion::UNKNOWN
@ UNKNOWN
clang::SanitizerMask
Definition: Sanitizers.h:30
clang::driver::toolchains::CudaToolChain::SupportsProfiling
bool SupportsProfiling() const override
SupportsProfiling - Does this tool chain support -pg.
Definition: Cuda.h:154