2022-12-21 01:30:50 -08:00
|
|
|
#ifndef TRITON_CONVERSION_TRITONGPU_TO_LLVM_CONVERT_LAYOUT_OP_H
|
|
|
|
#define TRITON_CONVERSION_TRITONGPU_TO_LLVM_CONVERT_LAYOUT_OP_H
|
|
|
|
|
|
|
|
#include "TritonGPUToLLVMBase.h"
|
|
|
|
|
|
|
|
using namespace mlir;
|
|
|
|
using namespace mlir::triton;
|
|
|
|
|
|
|
|
using ::mlir::triton::gpu::DotOperandEncodingAttr;
|
|
|
|
|
|
|
|
bool isMmaToDotShortcut(MmaEncodingAttr &mmaLayout,
|
|
|
|
DotOperandEncodingAttr &dotOperandLayout);
|
|
|
|
|
|
|
|
void storeBlockedToShared(Value src, Value llSrc, ArrayRef<Value> srcStrides,
|
|
|
|
ArrayRef<Value> srcIndices, Value dst, Value smemBase,
|
|
|
|
Type elemPtrTy, Location loc,
|
|
|
|
ConversionPatternRewriter &rewriter);
|
|
|
|
|
|
|
|
void populateConvertLayoutOpToLLVMPatterns(
|
|
|
|
mlir::LLVMTypeConverter &typeConverter, RewritePatternSet &patterns,
|
|
|
|
int numWarps, AxisInfoAnalysis &axisInfoAnalysis,
|
2022-12-30 03:19:59 +08:00
|
|
|
const Allocation *allocation, Value smem,
|
|
|
|
ConvertTritonGPUOpToLLVMPatternBase::IndexCacheInfo &indexCacheInfo,
|
|
|
|
PatternBenefit benefit);
|
2022-12-21 01:30:50 -08:00
|
|
|
|
|
|
|
#endif
|