clang  20.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 "SYCL.h"
13 #include "clang/Basic/Cuda.h"
14 #include "clang/Driver/Action.h"
15 #include "clang/Driver/Multilib.h"
16 #include "clang/Driver/Tool.h"
17 #include "clang/Driver/ToolChain.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 LibDevicePath;
36  std::string IncludePath;
37  llvm::StringMap<std::string> LibDeviceMap;
38 
39  // CUDA architectures for which we have raised an error in
40  // CheckCudaVersionSupportsArch.
41  mutable std::bitset<(int)OffloadArch::LAST> ArchsWithBadVersion;
42 
43 public:
44  CudaInstallationDetector(const Driver &D, const llvm::Triple &HostTriple,
45  const llvm::opt::ArgList &Args);
46 
47  void AddCudaIncludeArgs(const llvm::opt::ArgList &DriverArgs,
48  llvm::opt::ArgStringList &CC1Args) const;
49 
50  /// Emit an error if Version does not support the given Arch.
51  ///
52  /// If either Version or Arch is unknown, does not emit an error. Emits at
53  /// most one error per Arch.
55 
56  /// Check whether we detected a valid Cuda install.
57  bool isValid() const { return IsValid; }
58  /// Print information about the detected CUDA installation.
59  void print(raw_ostream &OS) const;
60 
61  /// Get the detected Cuda install's version.
62  CudaVersion version() const {
64  : Version;
65  }
66  /// Get the detected Cuda installation path.
67  StringRef getInstallPath() const { return InstallPath; }
68  /// Get the detected path to Cuda's bin directory.
69  StringRef getBinPath() const { return BinPath; }
70  /// Get the detected Cuda Include path.
71  StringRef getIncludePath() const { return IncludePath; }
72  /// Get the detected Cuda device library path.
73  StringRef getLibDevicePath() const { return LibDevicePath; }
74  /// Get libdevice file for given architecture
75  std::string getLibDeviceFile(StringRef Gpu) const {
76  return LibDeviceMap.lookup(Gpu);
77  }
79 };
80 
81 namespace tools {
82 namespace NVPTX {
83 
84 // Run ptxas, the NVPTX assembler.
85 class LLVM_LIBRARY_VISIBILITY Assembler final : public Tool {
86 public:
87  Assembler(const ToolChain &TC) : Tool("NVPTX::Assembler", "ptxas", TC) {}
88 
89  bool hasIntegratedCPP() const override { return false; }
90 
91  void ConstructJob(Compilation &C, const JobAction &JA,
92  const InputInfo &Output, const InputInfoList &Inputs,
93  const llvm::opt::ArgList &TCArgs,
94  const char *LinkingOutput) const override;
95 };
96 
97 // Runs fatbinary, which combines GPU object files ("cubin" files) and/or PTX
98 // assembly into a single output file.
99 class LLVM_LIBRARY_VISIBILITY FatBinary : public Tool {
100 public:
101  FatBinary(const ToolChain &TC) : Tool("NVPTX::Linker", "fatbinary", TC) {}
102 
103  bool hasIntegratedCPP() const override { return false; }
104 
105  void ConstructJob(Compilation &C, const JobAction &JA,
106  const InputInfo &Output, const InputInfoList &Inputs,
107  const llvm::opt::ArgList &TCArgs,
108  const char *LinkingOutput) const override;
109 };
110 
111 // Runs nvlink, which links GPU object files ("cubin" files) into a single file.
112 class LLVM_LIBRARY_VISIBILITY Linker final : public Tool {
113 public:
114  Linker(const ToolChain &TC) : Tool("NVPTX::Linker", "nvlink", TC) {}
115 
116  bool hasIntegratedCPP() const override { return false; }
117 
118  void ConstructJob(Compilation &C, const JobAction &JA,
119  const InputInfo &Output, const InputInfoList &Inputs,
120  const llvm::opt::ArgList &TCArgs,
121  const char *LinkingOutput) const override;
122 };
123 
124 class LLVM_LIBRARY_VISIBILITY OpenMPLinker : public Tool {
125  public:
127  : Tool("NVPTX::OpenMPLinker", "nvlink", TC) {}
128 
129  bool hasIntegratedCPP() const override { return false; }
130 
131  void ConstructJob(Compilation &C, const JobAction &JA,
132  const InputInfo &Output, const InputInfoList &Inputs,
133  const llvm::opt::ArgList &TCArgs,
134  const char *LinkingOutput) const override;
135 };
136 
137 class LLVM_LIBRARY_VISIBILITY SYCLLinker : public FatBinary {
138 public:
139  SYCLLinker(const ToolChain &TC) : FatBinary(TC) {}
140 
142  if (!SYCLToolChainLinker)
143  SYCLToolChainLinker.reset(new SYCL::Linker(getToolChain()));
144  return SYCLToolChainLinker.get();
145  }
146 private:
147  mutable std::unique_ptr<Tool> SYCLToolChainLinker;
148 };
149 
150 void getNVPTXTargetFeatures(const Driver &D, const llvm::Triple &Triple,
151  const llvm::opt::ArgList &Args,
152  std::vector<StringRef> &Features);
153 
154 } // end namespace NVPTX
155 } // end namespace tools
156 
157 namespace toolchains {
158 
159 class LLVM_LIBRARY_VISIBILITY NVPTXToolChain : public ToolChain {
160 public:
161  NVPTXToolChain(const Driver &D, const llvm::Triple &Triple,
162  const llvm::Triple &HostTriple, const llvm::opt::ArgList &Args,
163  bool Freestanding);
164 
165  NVPTXToolChain(const Driver &D, const llvm::Triple &Triple,
166  const llvm::opt::ArgList &Args);
167 
168  llvm::opt::DerivedArgList *
169  TranslateArgs(const llvm::opt::DerivedArgList &Args, StringRef BoundArch,
170  Action::OffloadKind DeviceOffloadKind) const override;
171 
172  void
173  addClangTargetOptions(const llvm::opt::ArgList &DriverArgs,
174  llvm::opt::ArgStringList &CC1Args,
175  Action::OffloadKind DeviceOffloadKind) const override;
176 
177  // Never try to use the integrated assembler with CUDA; always fork out to
178  // ptxas.
179  bool useIntegratedAs() const override { return false; }
180  bool isCrossCompiling() const override { return true; }
181  bool isPICDefault() const override { return false; }
182  bool isPIEDefault(const llvm::opt::ArgList &Args) const override {
183  return false;
184  }
185  bool HasNativeLLVMSupport() const override { return true; }
186  bool isPICDefaultForced() const override { return false; }
187  bool SupportsProfiling() const override { return false; }
188 
189  bool IsMathErrnoDefault() const override { return false; }
190 
191  bool supportsDebugInfoOption(const llvm::opt::Arg *A) const override;
192  void adjustDebugInfoKind(llvm::codegenoptions::DebugInfoKind &DebugInfoKind,
193  const llvm::opt::ArgList &Args) const override;
194 
195  // NVPTX supports only DWARF2.
196  unsigned GetDefaultDwarfVersion() const override { return 2; }
197  unsigned getMaxDwarfVersion() const override { return 2; }
198 
199  /// Uses nvptx-arch tool to get arch of the system GPU. Will return error
200  /// if unable to find one.
202  getSystemGPUArchs(const llvm::opt::ArgList &Args) const override;
203 
205 
206 protected:
207  Tool *buildAssembler() const override; // ptxas.
208  Tool *buildLinker() const override; // nvlink.
209 
210 private:
211  bool Freestanding = false;
212 };
213 
214 class LLVM_LIBRARY_VISIBILITY CudaToolChain : public NVPTXToolChain {
215 public:
216  CudaToolChain(const Driver &D, const llvm::Triple &Triple,
217  const ToolChain &HostTC, const llvm::opt::ArgList &Args,
218  const Action::OffloadKind OK);
219 
220  const llvm::Triple *getAuxTriple() const override {
221  return &HostTC.getTriple();
222  }
223 
224  bool HasNativeLLVMSupport() const override { return false; }
225 
226  std::string getInputFilename(const InputInfo &Input) const override;
227 
228  llvm::opt::DerivedArgList *
229  TranslateArgs(const llvm::opt::DerivedArgList &Args, StringRef BoundArch,
230  Action::OffloadKind DeviceOffloadKind) const override;
231  void
232  addClangTargetOptions(const llvm::opt::ArgList &DriverArgs,
233  llvm::opt::ArgStringList &CC1Args,
234  Action::OffloadKind DeviceOffloadKind) const override;
235 
236  llvm::DenormalMode getDefaultDenormalModeForType(
237  const llvm::opt::ArgList &DriverArgs, const JobAction &JA,
238  const llvm::fltSemantics *FPType = nullptr) const override;
239 
240  void AddCudaIncludeArgs(const llvm::opt::ArgList &DriverArgs,
241  llvm::opt::ArgStringList &CC1Args) const override;
242 
243  void addClangWarningOptions(llvm::opt::ArgStringList &CC1Args) const override;
244  CXXStdlibType GetCXXStdlibType(const llvm::opt::ArgList &Args) const override;
245  void
246  AddClangSystemIncludeArgs(const llvm::opt::ArgList &DriverArgs,
247  llvm::opt::ArgStringList &CC1Args) const override;
248  void AddClangCXXStdlibIncludeArgs(
249  const llvm::opt::ArgList &Args,
250  llvm::opt::ArgStringList &CC1Args) const override;
251  void AddIAMCUIncludeArgs(const llvm::opt::ArgList &DriverArgs,
252  llvm::opt::ArgStringList &CC1Args) const override;
253 
254  SanitizerMask getSupportedSanitizers() const override;
255 
256  VersionTuple
257  computeMSVCVersion(const Driver *D,
258  const llvm::opt::ArgList &Args) const override;
259 
260  Tool *SelectTool(const JobAction &JA) const override;
262 
263 protected:
264  Tool *buildAssembler() const override; // ptxas
265  Tool *buildLinker() const override; // fatbinary (ok, not really a linker)
266 
267 private:
268  const Action::OffloadKind OK;
269 };
270 
271 } // end namespace toolchains
272 } // end namespace driver
273 } // end namespace clang
274 
275 #endif // LLVM_CLANG_LIB_DRIVER_TOOLCHAINS_CUDA_H
const Decl * D
__device__ int
Compilation - A set of tasks to perform for a single driver invocation.
Definition: Compilation.h:45
A class to find a viable CUDA installation.
Definition: Cuda.h:28
void AddCudaIncludeArgs(const llvm::opt::ArgList &DriverArgs, llvm::opt::ArgStringList &CC1Args) const
Definition: Cuda.cpp:297
CudaInstallationDetector(const Driver &D, const llvm::Triple &HostTriple, const llvm::opt::ArgList &Args)
Definition: Cuda.cpp:135
StringRef getLibDevicePath() const
Get the detected Cuda device library path.
Definition: Cuda.h:73
CudaVersion version() const
Get the detected Cuda install's version.
Definition: Cuda.h:62
StringRef getInstallPath() const
Get the detected Cuda installation path.
Definition: Cuda.h:67
std::string getLibDeviceFile(StringRef Gpu) const
Get libdevice file for given architecture.
Definition: Cuda.h:75
StringRef getBinPath() const
Get the detected path to Cuda's bin directory.
Definition: Cuda.h:69
void CheckCudaVersionSupportsArch(OffloadArch Arch) const
Emit an error if Version does not support the given Arch.
Definition: Cuda.cpp:321
void print(raw_ostream &OS) const
Print information about the detected CUDA installation.
Definition: Cuda.cpp:338
StringRef getIncludePath() const
Get the detected Cuda Include path.
Definition: Cuda.h:71
bool isValid() const
Check whether we detected a valid Cuda install.
Definition: Cuda.h:57
Driver - Encapsulate logic for constructing compilation processes from a set of gcc-driver-like comma...
Definition: Driver.h:77
InputInfo - Wrapper for information about an input source.
Definition: InputInfo.h:22
ToolChain - Access to tools for a single platform.
Definition: ToolChain.h:92
const llvm::Triple & getTriple() const
Definition: ToolChain.h:271
Tool - Information on a specific compilation tool.
Definition: Tool.h:32
const llvm::Triple * getAuxTriple() const override
Get the toolchain's aux triple, if it has one.
Definition: Cuda.h:220
bool HasNativeLLVMSupport() const override
HasNativeLTOLinker - Check whether the linker and related tools have native LLVM support.
Definition: Cuda.h:224
unsigned getMaxDwarfVersion() const override
Definition: Cuda.h:197
CudaInstallationDetector CudaInstallation
Definition: Cuda.h:204
bool isPICDefault() const override
Test whether this toolchain defaults to PIC.
Definition: Cuda.h:181
bool isCrossCompiling() const override
Returns true if the toolchain is targeting a non-native architecture.
Definition: Cuda.h:180
bool IsMathErrnoDefault() const override
IsMathErrnoDefault - Does this tool chain use -fmath-errno by default.
Definition: Cuda.h:189
bool HasNativeLLVMSupport() const override
HasNativeLTOLinker - Check whether the linker and related tools have native LLVM support.
Definition: Cuda.h:185
NVPTXToolChain(const Driver &D, const llvm::Triple &Triple, const llvm::opt::ArgList &Args)
bool SupportsProfiling() const override
SupportsProfiling - Does this tool chain support -pg.
Definition: Cuda.h:187
unsigned GetDefaultDwarfVersion() const override
Definition: Cuda.h:196
bool useIntegratedAs() const override
Check if the toolchain should use the integrated assembler.
Definition: Cuda.h:179
bool isPIEDefault(const llvm::opt::ArgList &Args) const override
Test whether this toolchain defaults to PIE.
Definition: Cuda.h:182
NVPTXToolChain(const Driver &D, const llvm::Triple &Triple, const llvm::Triple &HostTriple, const llvm::opt::ArgList &Args, bool Freestanding)
bool isPICDefaultForced() const override
Tests whether this toolchain forces its default for PIC, PIE or non-PIC.
Definition: Cuda.h:186
bool hasIntegratedCPP() const override
Definition: Cuda.h:89
Assembler(const ToolChain &TC)
Definition: Cuda.h:87
FatBinary(const ToolChain &TC)
Definition: Cuda.h:101
bool hasIntegratedCPP() const override
Definition: Cuda.h:103
Linker(const ToolChain &TC)
Definition: Cuda.h:114
bool hasIntegratedCPP() const override
Definition: Cuda.h:116
bool hasIntegratedCPP() const override
Definition: Cuda.h:129
OpenMPLinker(const ToolChain &TC)
Definition: Cuda.h:126
SYCLLinker(const ToolChain &TC)
Definition: Cuda.h:139
Tool * GetSYCLToolChainLinker() const
Definition: Cuda.h:141
void getNVPTXTargetFeatures(const Driver &D, const llvm::Triple &Triple, const llvm::opt::ArgList &Args, std::vector< StringRef > &Features)
Definition: Cuda.cpp:741
The JSON file list parser is used to communicate input to InstallAPI.
OffloadArch
Definition: Cuda.h:55
CudaVersion
Definition: Cuda.h:20