1//===- NVPTX.cpp ----------------------------------------------------------===//
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#include "ABIInfoImpl.h"
10#include "TargetInfo.h"
11#include "clang/Basic/SyncScope.h"
12#include "llvm/ADT/STLExtras.h"
13#include "llvm/ADT/StringExtras.h"
14#include "llvm/IR/CallingConv.h"
15#include "llvm/IR/IntrinsicsNVPTX.h"
16#include "llvm/IR/LLVMContext.h"
17#include "llvm/Support/NVVMAttributes.h"
18
19using namespace clang;
20using namespace clang::CodeGen;
21
22//===----------------------------------------------------------------------===//
23// NVPTX ABI Implementation
24//===----------------------------------------------------------------------===//
25
26namespace {
27
28class NVPTXTargetCodeGenInfo;
29
30class NVPTXABIInfo : public ABIInfo {
31 NVPTXTargetCodeGenInfo &CGInfo;
32
33public:
34 NVPTXABIInfo(CodeGenTypes &CGT, NVPTXTargetCodeGenInfo &Info)
35 : ABIInfo(CGT), CGInfo(Info) {}
36
37 ABIArgInfo classifyReturnType(QualType RetTy) const;
38 ABIArgInfo classifyArgumentType(QualType Ty) const;
39
40 void computeInfo(CGFunctionInfo &FI) const override;
41 RValue EmitVAArg(CodeGenFunction &CGF, Address VAListAddr, QualType Ty,
42 AggValueSlot Slot) const override;
43 bool isUnsupportedType(QualType T) const;
44 ABIArgInfo coerceToIntArrayWithLimit(QualType Ty, unsigned MaxSize) const;
45};
46
47class NVPTXTargetCodeGenInfo : public TargetCodeGenInfo {
48public:
49 NVPTXTargetCodeGenInfo(CodeGenTypes &CGT)
50 : TargetCodeGenInfo(std::make_unique<NVPTXABIInfo>(args&: CGT, args&: *this)) {}
51
52 void setTargetAttributes(const Decl *D, llvm::GlobalValue *GV,
53 CodeGen::CodeGenModule &M) const override;
54 void setTargetAtomicMetadata(CodeGenFunction &CGF,
55 llvm::Instruction &AtomicInst,
56 const AtomicExpr *AE) const override;
57 bool shouldEmitStaticExternCAliases() const override;
58
59 StringRef getLLVMSyncScopeStr(const LangOptions &LangOpts, SyncScope Scope,
60 llvm::AtomicOrdering Ordering) const override;
61
62 llvm::Constant *getNullPointer(const CodeGen::CodeGenModule &CGM,
63 llvm::PointerType *T,
64 QualType QT) const override;
65
66 llvm::Type *getCUDADeviceBuiltinSurfaceDeviceType() const override {
67 // On the device side, surface reference is represented as an object handle
68 // in 64-bit integer.
69 return llvm::Type::getInt64Ty(C&: getABIInfo().getVMContext());
70 }
71
72 llvm::Type *getCUDADeviceBuiltinTextureDeviceType() const override {
73 // On the device side, texture reference is represented as an object handle
74 // in 64-bit integer.
75 return llvm::Type::getInt64Ty(C&: getABIInfo().getVMContext());
76 }
77
78 bool emitCUDADeviceBuiltinSurfaceDeviceCopy(CodeGenFunction &CGF, LValue Dst,
79 LValue Src) const override {
80 emitBuiltinSurfTexDeviceCopy(CGF, Dst, Src);
81 return true;
82 }
83
84 bool emitCUDADeviceBuiltinTextureDeviceCopy(CodeGenFunction &CGF, LValue Dst,
85 LValue Src) const override {
86 emitBuiltinSurfTexDeviceCopy(CGF, Dst, Src);
87 return true;
88 }
89
90 unsigned getDeviceKernelCallingConv() const override {
91 return llvm::CallingConv::PTX_Kernel;
92 }
93
94 void setCUDAKernelCallingConvention(const FunctionType *&FT) const override {
95 FT = getABIInfo().getContext().adjustFunctionType(
96 Fn: FT, EInfo: FT->getExtInfo().withCallingConv(cc: CC_DeviceKernel));
97 }
98
99 // Adds a NamedMDNode with GV, Name, and Operand as operands, and adds the
100 // resulting MDNode to the nvvm.annotations MDNode.
101 static void addNVVMMetadata(llvm::GlobalValue *GV, StringRef Name,
102 int Operand);
103
104private:
105 static void emitBuiltinSurfTexDeviceCopy(CodeGenFunction &CGF, LValue Dst,
106 LValue Src) {
107 llvm::Value *Handle = nullptr;
108 llvm::Constant *C =
109 llvm::dyn_cast<llvm::Constant>(Val: Src.getAddress().emitRawPointer(CGF));
110 // Lookup `addrspacecast` through the constant pointer if any.
111 if (auto *ASC = llvm::dyn_cast_or_null<llvm::AddrSpaceCastOperator>(Val: C))
112 C = llvm::cast<llvm::Constant>(Val: ASC->getPointerOperand());
113 if (auto *GV = llvm::dyn_cast_or_null<llvm::GlobalVariable>(Val: C)) {
114 // Load the handle from the specific global variable using
115 // `nvvm.texsurf.handle.internal` intrinsic.
116 Handle = CGF.EmitRuntimeCall(
117 callee: CGF.CGM.getIntrinsic(IID: llvm::Intrinsic::nvvm_texsurf_handle_internal,
118 Tys: {GV->getType()}),
119 args: {GV}, name: "texsurf_handle");
120 } else
121 Handle = CGF.EmitLoadOfScalar(lvalue: Src, Loc: SourceLocation());
122 CGF.EmitStoreOfScalar(value: Handle, lvalue: Dst);
123 }
124};
125
126/// Checks if the type is unsupported directly by the current target.
127bool NVPTXABIInfo::isUnsupportedType(QualType T) const {
128 ASTContext &Context = getContext();
129 if (!Context.getTargetInfo().hasFloat16Type() && T->isFloat16Type())
130 return true;
131 if (!Context.getTargetInfo().hasFloat128Type() &&
132 (T->isFloat128Type() ||
133 (T->isRealFloatingType() && Context.getTypeSize(T) == 128)))
134 return true;
135 if (const auto *EIT = T->getAs<BitIntType>())
136 return EIT->getNumBits() >
137 (Context.getTargetInfo().hasInt128Type() ? 128U : 64U);
138 if (!Context.getTargetInfo().hasInt128Type() && T->isIntegerType() &&
139 Context.getTypeSize(T) > 64U)
140 return true;
141 if (const auto *AT = T->getAsArrayTypeUnsafe())
142 return isUnsupportedType(T: AT->getElementType());
143 const auto *RD = T->getAsRecordDecl();
144 if (!RD)
145 return false;
146
147 // If this is a C++ record, check the bases first.
148 if (const CXXRecordDecl *CXXRD = dyn_cast<CXXRecordDecl>(Val: RD))
149 for (const CXXBaseSpecifier &I : CXXRD->bases())
150 if (isUnsupportedType(T: I.getType()))
151 return true;
152
153 for (const FieldDecl *I : RD->fields())
154 if (isUnsupportedType(T: I->getType()))
155 return true;
156 return false;
157}
158
159/// Coerce the given type into an array with maximum allowed size of elements.
160ABIArgInfo NVPTXABIInfo::coerceToIntArrayWithLimit(QualType Ty,
161 unsigned MaxSize) const {
162 // Alignment and Size are measured in bits.
163 const uint64_t Size = getContext().getTypeSize(T: Ty);
164 const uint64_t Alignment = getContext().getTypeAlign(T: Ty);
165 const unsigned Div = std::min<unsigned>(a: MaxSize, b: Alignment);
166 llvm::Type *IntType = llvm::Type::getIntNTy(C&: getVMContext(), N: Div);
167 const uint64_t NumElements = (Size + Div - 1) / Div;
168 return ABIArgInfo::getDirect(T: llvm::ArrayType::get(ElementType: IntType, NumElements));
169}
170
171ABIArgInfo NVPTXABIInfo::classifyReturnType(QualType RetTy) const {
172 if (RetTy->isVoidType())
173 return ABIArgInfo::getIgnore();
174
175 if (getContext().getLangOpts().OpenMP &&
176 getContext().getLangOpts().OpenMPIsTargetDevice &&
177 isUnsupportedType(T: RetTy))
178 return coerceToIntArrayWithLimit(Ty: RetTy, MaxSize: 64);
179
180 // note: this is different from default ABI
181 if (!RetTy->isScalarType())
182 return ABIArgInfo::getDirect();
183
184 // Treat an enum type as its underlying type.
185 if (const auto *ED = RetTy->getAsEnumDecl())
186 RetTy = ED->getIntegerType();
187
188 return (isPromotableIntegerTypeForABI(Ty: RetTy) ? ABIArgInfo::getExtend(Ty: RetTy)
189 : ABIArgInfo::getDirect());
190}
191
192ABIArgInfo NVPTXABIInfo::classifyArgumentType(QualType Ty) const {
193 // Treat an enum type as its underlying type.
194 if (const auto *ED = Ty->getAsEnumDecl())
195 Ty = ED->getIntegerType();
196
197 // Return aggregates type as indirect by value
198 if (isAggregateTypeForABI(T: Ty)) {
199 // Under CUDA device compilation, tex/surf builtin types are replaced with
200 // object types and passed directly.
201 if (getContext().getLangOpts().CUDAIsDevice) {
202 if (Ty->isCUDADeviceBuiltinSurfaceType())
203 return ABIArgInfo::getDirect(
204 T: CGInfo.getCUDADeviceBuiltinSurfaceDeviceType());
205 if (Ty->isCUDADeviceBuiltinTextureType())
206 return ABIArgInfo::getDirect(
207 T: CGInfo.getCUDADeviceBuiltinTextureDeviceType());
208 }
209 return getNaturalAlignIndirect(
210 Ty, /* AddrSpace */ getDataLayout().getAllocaAddrSpace(),
211 /* byval */ ByVal: true);
212 }
213
214 if (const auto *EIT = Ty->getAs<BitIntType>()) {
215 if ((EIT->getNumBits() > 128) ||
216 (!getContext().getTargetInfo().hasInt128Type() &&
217 EIT->getNumBits() > 64))
218 return getNaturalAlignIndirect(
219 Ty, /* AddrSpace */ getDataLayout().getAllocaAddrSpace(),
220 /* byval */ ByVal: true);
221 }
222
223 return (isPromotableIntegerTypeForABI(Ty) ? ABIArgInfo::getExtend(Ty)
224 : ABIArgInfo::getDirect());
225}
226
227void NVPTXABIInfo::computeInfo(CGFunctionInfo &FI) const {
228 if (!getCXXABI().classifyReturnType(FI))
229 FI.getReturnInfo() = classifyReturnType(RetTy: FI.getReturnType());
230
231 for (auto &&[ArgumentsCount, I] : llvm::enumerate(First: FI.arguments()))
232 I.info = ArgumentsCount < FI.getNumRequiredArgs()
233 ? classifyArgumentType(Ty: I.type)
234 : ABIArgInfo::getDirect();
235
236 // Always honor user-specified calling convention.
237 if (FI.getCallingConvention() != llvm::CallingConv::C)
238 return;
239
240 FI.setEffectiveCallingConvention(getRuntimeCC());
241}
242
243RValue NVPTXABIInfo::EmitVAArg(CodeGenFunction &CGF, Address VAListAddr,
244 QualType Ty, AggValueSlot Slot) const {
245 return emitVoidPtrVAArg(CGF, VAListAddr, ValueTy: Ty, /*IsIndirect=*/false,
246 ValueInfo: getContext().getTypeInfoInChars(T: Ty),
247 SlotSizeAndAlign: CharUnits::fromQuantity(Quantity: 1),
248 /*AllowHigherAlign=*/true, Slot);
249}
250
251void NVPTXTargetCodeGenInfo::setTargetAttributes(
252 const Decl *D, llvm::GlobalValue *GV, CodeGen::CodeGenModule &M) const {
253 if (GV->isDeclaration())
254 return;
255 const VarDecl *VD = dyn_cast_or_null<VarDecl>(Val: D);
256 if (VD) {
257 if (M.getLangOpts().CUDA) {
258 if (VD->getType()->isCUDADeviceBuiltinSurfaceType())
259 addNVVMMetadata(GV, Name: "surface", Operand: 1);
260 else if (VD->getType()->isCUDADeviceBuiltinTextureType())
261 addNVVMMetadata(GV, Name: "texture", Operand: 1);
262 return;
263 }
264 }
265
266 const FunctionDecl *FD = dyn_cast_or_null<FunctionDecl>(Val: D);
267 if (!FD)
268 return;
269
270 llvm::Function *F = cast<llvm::Function>(Val: GV);
271
272 // Perform special handling in OpenCL/CUDA mode
273 if (M.getLangOpts().OpenCL || M.getLangOpts().CUDA) {
274 // Use function attributes to check for kernel functions
275 // By default, all functions are device functions
276 if (FD->hasAttr<DeviceKernelAttr>() || FD->hasAttr<CUDAGlobalAttr>()) {
277 // OpenCL/CUDA kernel functions get kernel metadata
278 // And kernel functions are not subject to inlining
279 F->addFnAttr(Kind: llvm::Attribute::NoInline);
280 if (FD->hasAttr<CUDAGlobalAttr>()) {
281 for (auto IV : llvm::enumerate(First: FD->parameters()))
282 if (IV.value()->hasAttr<CUDAGridConstantAttr>())
283 F->addParamAttr(ArgNo: IV.index(),
284 Attr: llvm::Attribute::get(Context&: F->getContext(),
285 Kind: llvm::NVVMAttr::GridConstant));
286 }
287 if (CUDALaunchBoundsAttr *Attr = FD->getAttr<CUDALaunchBoundsAttr>())
288 M.handleCUDALaunchBoundsAttr(F, A: Attr);
289 }
290 }
291}
292
293void NVPTXTargetCodeGenInfo::addNVVMMetadata(llvm::GlobalValue *GV,
294 StringRef Name, int Operand) {
295 llvm::Module *M = GV->getParent();
296 llvm::LLVMContext &Ctx = M->getContext();
297
298 // Get "nvvm.annotations" metadata node
299 llvm::NamedMDNode *MD = M->getOrInsertNamedMetadata(Name: "nvvm.annotations");
300
301 SmallVector<llvm::Metadata *, 5> MDVals = {
302 llvm::ConstantAsMetadata::get(C: GV), llvm::MDString::get(Context&: Ctx, Str: Name),
303 llvm::ConstantAsMetadata::get(
304 C: llvm::ConstantInt::get(Ty: llvm::Type::getInt32Ty(C&: Ctx), V: Operand))};
305
306 // Append metadata to nvvm.annotations
307 MD->addOperand(M: llvm::MDNode::get(Context&: Ctx, MDs: MDVals));
308}
309
310bool NVPTXTargetCodeGenInfo::shouldEmitStaticExternCAliases() const {
311 return false;
312}
313
314void NVPTXTargetCodeGenInfo::setTargetAtomicMetadata(
315 CodeGenFunction &CGF, llvm::Instruction &AtomicInst,
316 const AtomicExpr *AE) const {
317 AtomicOptions AO = CGF.CGM.getAtomicOpts();
318 if (AO.getOption(Kind: clang::AtomicOptionKind::IgnoreDenormalMode))
319 addAtomicIgnoreDenormalModeMetadata(CGF, AtomicInst, /*AllowHalf=*/true);
320}
321
322StringRef NVPTXTargetCodeGenInfo::getLLVMSyncScopeStr(
323 const LangOptions &LangOpts, SyncScope Scope,
324 llvm::AtomicOrdering Ordering) const {
325 return *llvm::getAtomicScopeIRString(T: getABIInfo().getTarget().getTriple(),
326 S: getAtomicScope(S: Scope));
327}
328
329llvm::Constant *
330NVPTXTargetCodeGenInfo::getNullPointer(const CodeGen::CodeGenModule &CGM,
331 llvm::PointerType *PT,
332 QualType QT) const {
333 auto &Ctx = CGM.getContext();
334 if (PT->getAddressSpace() != Ctx.getTargetAddressSpace(AS: LangAS::opencl_local))
335 return llvm::ConstantPointerNull::get(T: PT);
336
337 auto NPT = llvm::PointerType::get(
338 C&: PT->getContext(), AddressSpace: Ctx.getTargetAddressSpace(AS: LangAS::opencl_generic));
339 return llvm::ConstantExpr::getAddrSpaceCast(
340 C: llvm::ConstantPointerNull::get(T: NPT), Ty: PT);
341}
342} // namespace
343
344void CodeGenModule::handleCUDALaunchBoundsAttr(llvm::Function *F,
345 const CUDALaunchBoundsAttr *Attr,
346 int32_t *MaxThreadsVal,
347 int32_t *MinBlocksVal,
348 int32_t *MaxClusterRankVal) {
349 llvm::APSInt MaxThreads(32);
350 MaxThreads = Attr->getMaxThreads()->EvaluateKnownConstInt(Ctx: getContext());
351 if (MaxThreads > 0) {
352 if (MaxThreadsVal)
353 *MaxThreadsVal = MaxThreads.getExtValue();
354 if (F)
355 F->addFnAttr(Kind: llvm::NVVMAttr::MaxNTID,
356 Val: llvm::utostr(X: MaxThreads.getExtValue()));
357 }
358
359 // min and max blocks is an optional argument for CUDALaunchBoundsAttr. If it
360 // was not specified in __launch_bounds__ or if the user specified a 0 value,
361 // we don't have to add a PTX directive.
362 if (Attr->getMinBlocks()) {
363 llvm::APSInt MinBlocks(32);
364 MinBlocks = Attr->getMinBlocks()->EvaluateKnownConstInt(Ctx: getContext());
365 if (MinBlocks > 0) {
366 if (MinBlocksVal)
367 *MinBlocksVal = MinBlocks.getExtValue();
368 if (F)
369 F->addFnAttr(Kind: llvm::NVVMAttr::MinCTASm,
370 Val: llvm::utostr(X: MinBlocks.getExtValue()));
371 }
372 }
373 if (Attr->getMaxBlocks()) {
374 llvm::APSInt MaxBlocks(32);
375 MaxBlocks = Attr->getMaxBlocks()->EvaluateKnownConstInt(Ctx: getContext());
376 if (MaxBlocks > 0) {
377 if (MaxClusterRankVal)
378 *MaxClusterRankVal = MaxBlocks.getExtValue();
379 if (F)
380 F->addFnAttr(Kind: llvm::NVVMAttr::MaxClusterRank,
381 Val: llvm::utostr(X: MaxBlocks.getExtValue()));
382 }
383 }
384}
385
386std::unique_ptr<TargetCodeGenInfo>
387CodeGen::createNVPTXTargetCodeGenInfo(CodeGenModule &CGM) {
388 return std::make_unique<NVPTXTargetCodeGenInfo>(args&: CGM.getTypes());
389}
390