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