FLANG
CUDAIntrinsicCall.h
1//==-- Builder/CUDAIntrinsicCall.h - lowering of CUDA intrinsics ---*-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 FORTRAN_LOWER_CUDAINTRINSICCALL_H
10#define FORTRAN_LOWER_CUDAINTRINSICCALL_H
11
12#include "flang/Optimizer/Builder/IntrinsicCall.h"
13#include "mlir/Dialect/LLVMIR/NVVMDialect.h"
14
15namespace fir {
16
17struct CUDAIntrinsicLibrary : IntrinsicLibrary {
18
19 // Constructors.
20 explicit CUDAIntrinsicLibrary(fir::FirOpBuilder &builder, mlir::Location loc)
21 : IntrinsicLibrary(builder, loc) {}
22 CUDAIntrinsicLibrary() = delete;
23 CUDAIntrinsicLibrary(const CUDAIntrinsicLibrary &) = delete;
24
25 // CUDA intrinsic handlers.
26 mlir::Value genAtomicAdd(mlir::Type, llvm::ArrayRef<mlir::Value>);
27 fir::ExtendedValue genAtomicAddR2(mlir::Type,
29 template <int extent>
30 fir::ExtendedValue genAtomicAddVector(mlir::Type,
32 mlir::Value genAtomicAnd(mlir::Type, llvm::ArrayRef<mlir::Value>);
33 fir::ExtendedValue genAtomicCas(mlir::Type,
35 mlir::Value genAtomicDec(mlir::Type, llvm::ArrayRef<mlir::Value>);
36 fir::ExtendedValue genAtomicExch(mlir::Type,
38 mlir::Value genAtomicInc(mlir::Type, llvm::ArrayRef<mlir::Value>);
39 mlir::Value genAtomicMax(mlir::Type, llvm::ArrayRef<mlir::Value>);
40 mlir::Value genAtomicMin(mlir::Type, llvm::ArrayRef<mlir::Value>);
41 mlir::Value genAtomicOr(mlir::Type, llvm::ArrayRef<mlir::Value>);
42 mlir::Value genAtomicSub(mlir::Type, llvm::ArrayRef<mlir::Value>);
43 fir::ExtendedValue genAtomicXor(mlir::Type,
45 mlir::Value genBarrierArrive(mlir::Type, llvm::ArrayRef<mlir::Value>);
46 mlir::Value genBarrierArriveCnt(mlir::Type, llvm::ArrayRef<mlir::Value>);
47 void genBarrierInit(llvm::ArrayRef<fir::ExtendedValue>);
48 mlir::Value genBarrierTryWait(mlir::Type, llvm::ArrayRef<mlir::Value>);
49 mlir::Value genBarrierTryWaitSleep(mlir::Type, llvm::ArrayRef<mlir::Value>);
50 void genFenceProxyAsync(llvm::ArrayRef<fir::ExtendedValue>);
51 template <const char *fctName, int extent>
52 fir::ExtendedValue genLDXXFunc(mlir::Type,
54 mlir::Value genMatchAllSync(mlir::Type, llvm::ArrayRef<mlir::Value>);
55 mlir::Value genMatchAnySync(mlir::Type, llvm::ArrayRef<mlir::Value>);
56 template <typename OpTy>
57 mlir::Value genNVVMTime(mlir::Type, llvm::ArrayRef<mlir::Value>);
58 void genSyncThreads(llvm::ArrayRef<fir::ExtendedValue>);
59 mlir::Value genSyncThreadsAnd(mlir::Type, llvm::ArrayRef<mlir::Value>);
60 mlir::Value genSyncThreadsCount(mlir::Type, llvm::ArrayRef<mlir::Value>);
61 mlir::Value genSyncThreadsOr(mlir::Type, llvm::ArrayRef<mlir::Value>);
62 void genSyncWarp(llvm::ArrayRef<fir::ExtendedValue>);
63 mlir::Value genThisGrid(mlir::Type, llvm::ArrayRef<mlir::Value>);
64 mlir::Value genThisThreadBlock(mlir::Type, llvm::ArrayRef<mlir::Value>);
65 mlir::Value genThisWarp(mlir::Type, llvm::ArrayRef<mlir::Value>);
66 void genThreadFence(llvm::ArrayRef<fir::ExtendedValue>);
67 void genThreadFenceBlock(llvm::ArrayRef<fir::ExtendedValue>);
68 void genThreadFenceSystem(llvm::ArrayRef<fir::ExtendedValue>);
69 void genTMABulkCommitGroup(llvm::ArrayRef<fir::ExtendedValue>);
70 void genTMABulkG2S(llvm::ArrayRef<fir::ExtendedValue>);
71 void genTMABulkLoadC4(llvm::ArrayRef<fir::ExtendedValue>);
72 void genTMABulkLoadC8(llvm::ArrayRef<fir::ExtendedValue>);
73 void genTMABulkLoadI4(llvm::ArrayRef<fir::ExtendedValue>);
74 void genTMABulkLoadI8(llvm::ArrayRef<fir::ExtendedValue>);
75 void genTMABulkLoadR2(llvm::ArrayRef<fir::ExtendedValue>);
76 void genTMABulkLoadR4(llvm::ArrayRef<fir::ExtendedValue>);
77 void genTMABulkLoadR8(llvm::ArrayRef<fir::ExtendedValue>);
78 void genTMABulkS2G(llvm::ArrayRef<fir::ExtendedValue>);
79 void genTMABulkStoreC4(llvm::ArrayRef<fir::ExtendedValue>);
80 void genTMABulkStoreC8(llvm::ArrayRef<fir::ExtendedValue>);
81 void genTMABulkStoreI4(llvm::ArrayRef<fir::ExtendedValue>);
82 void genTMABulkStoreI8(llvm::ArrayRef<fir::ExtendedValue>);
83 void genTMABulkStoreR2(llvm::ArrayRef<fir::ExtendedValue>);
84 void genTMABulkStoreR4(llvm::ArrayRef<fir::ExtendedValue>);
85 void genTMABulkStoreR8(llvm::ArrayRef<fir::ExtendedValue>);
86 void genTMABulkWaitGroup(llvm::ArrayRef<fir::ExtendedValue>);
87 template <mlir::NVVM::VoteSyncKind kind>
88 mlir::Value genVoteSync(mlir::Type, llvm::ArrayRef<mlir::Value>);
89};
90
91const IntrinsicHandler *findCUDAIntrinsicHandler(llvm::StringRef name);
92
93} // namespace fir
94
95#endif // FORTRAN_LOWER_CUDAINTRINSICCALL_H
Definition BoxValue.h:478
Definition FIRBuilder.h:55
Definition FIRType.h:92
Definition AbstractConverter.h:34
Definition IntrinsicCall.h:564