-
Notifications
You must be signed in to change notification settings - Fork 3.2k
Expand file tree
/
Copy pathTargetInfoBase.h
More file actions
39 lines (28 loc) · 1.17 KB
/
Copy pathTargetInfoBase.h
File metadata and controls
39 lines (28 loc) · 1.17 KB
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
#ifndef PROTONGPU_TO_LLVM_TARGETINFO_BASE_H
#define PROTONGPU_TO_LLVM_TARGETINFO_BASE_H
#include "mlir/IR/Attributes.h"
#include "triton/Conversion/MLIRTypes.h"
#include "triton/Conversion/TritonGPUToLLVM/TargetInfoBase.h"
#include "llvm/ADT/SmallVector.h"
namespace mlir::triton::proton::gpu {
class TargetInfoBase {
public:
explicit TargetInfoBase(const mlir::triton::TargetInfoBase &helper)
: helper(helper) {}
virtual const mlir::triton::TargetInfoBase &getTritonTargetInfo() const {
return helper;
}
// Return the global cycle counter value (i.e., synchronized across SMs) in
// nanoseconds, regardless of the clock frequency.
virtual Value globalTime(ConversionPatternRewriter &rewriter,
Location loc) const = 0;
virtual Value processorId(ConversionPatternRewriter &rewriter,
Location loc) const = 0;
virtual int getAddressSpace(Attribute addressSpace) const = 0;
virtual int getIndexPtrAddrSpace() const = 0;
virtual ~TargetInfoBase() = default;
protected:
const mlir::triton::TargetInfoBase &helper;
};
} // namespace mlir::triton::proton::gpu
#endif // PROTONGPU_TO_LLVM_TARGETINFO_BASE_H