Google Git
Sign in
llvm / llvm-project / refs/heads/users/cdevadas/enhance-createFrom-function / . / mlir / lib / Conversion / AMDGPUToROCDL / AMDGPUToROCDL.cpp
blob: 6427807e944a1edd3dfb6cf9d17b472d54ec3c3e [file] [log] [blame] [edit]
//===- AMDGPUToROCDL.cpp - AMDGPU to ROCDL dialect conversion -------===//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//
#include "mlir/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.h"
#include "mlir/Conversion/GPUCommon/GPUCommonPass.h"
#include "mlir/Conversion/LLVMCommon/ConversionTarget.h"
#include "mlir/Conversion/LLVMCommon/Pattern.h"
#include "mlir/Conversion/LLVMCommon/TypeConverter.h"
#include "mlir/Dialect/AMDGPU/IR/AMDGPUDialect.h"
#include "mlir/Dialect/AMDGPU/Utils/Chipset.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
#include "mlir/Dialect/LLVMIR/LLVMTypes.h"
#include "mlir/Dialect/LLVMIR/ROCDLDialect.h"
#include "mlir/IR/Attributes.h"
#include "mlir/IR/BuiltinAttributes.h"
#include "mlir/IR/BuiltinTypes.h"
#include "mlir/IR/Matchers.h"
#include "mlir/IR/TypeUtilities.h"
#include "mlir/Pass/Pass.h"
#include "../LLVMCommon/MemRefDescriptor.h"
#include "llvm/ADT/STLExtras.h"
#include "llvm/ADT/TypeSwitch.h"
#include "llvm/Support/Casting.h"
#include "llvm/Support/ErrorHandling.h"
#include <optional>
namespace mlir {
#define GEN_PASS_DEF_CONVERTAMDGPUTOROCDLPASS
#include "mlir/Conversion/Passes.h.inc"
} // namespace mlir
using namespace mlir;
using namespace mlir::amdgpu;
// Define commonly used chipsets versions for convenience.
constexpr Chipset kGfx908 = Chipset(9, 0, 8);
constexpr Chipset kGfx90a = Chipset(9, 0, 0xa);
constexpr Chipset kGfx942 = Chipset(9, 4, 2);
constexpr Chipset kGfx950 = Chipset(9, 5, 0);
constexpr Chipset kGfx1250 = Chipset(12, 5, 0);
/// Convert an unsigned number `val` to i32.
static Value convertUnsignedToI32(ConversionPatternRewriter &rewriter,
Location loc, Value val) {
IntegerType i32 = rewriter.getI32Type();
// Force check that `val` is of int type.
auto valTy = cast<IntegerType>(val.getType());
if (i32 == valTy)
return val;
return valTy.getWidth() > 32
? Value(LLVM::TruncOp::create(rewriter, loc, i32, val))
: Value(LLVM::ZExtOp::create(rewriter, loc, i32, val));
}
static Value createI32Constant(ConversionPatternRewriter &rewriter,
Location loc, int32_t value) {
return LLVM::ConstantOp::create(rewriter, loc, rewriter.getI32Type(), value);
}
/// Convert an unsigned number `val` to i64.
static Value convertUnsignedToI64(ConversionPatternRewriter &rewriter,
Location loc, Value val) {
IntegerType i64 = rewriter.getI64Type();
// Force check that `val` is of int type.
auto valTy = cast<IntegerType>(val.getType());
if (i64 == valTy)
return val;
return valTy.getWidth() > 64
? Value(LLVM::TruncOp::create(rewriter, loc, i64, val))
: Value(LLVM::ZExtOp::create(rewriter, loc, i64, val));
}
static Value createI64Constant(ConversionPatternRewriter &rewriter,
Location loc, int64_t value) {
return LLVM::ConstantOp::create(rewriter, loc, rewriter.getI64Type(), value);
}
/// Returns the linear index used to access an element in the memref.
static Value getLinearIndexI32(ConversionPatternRewriter &rewriter,
Location loc, MemRefDescriptor &memRefDescriptor,
ValueRange indices, ArrayRef<int64_t> strides) {
IntegerType i32 = rewriter.getI32Type();
Value index;
for (auto [i, increment, stride] : llvm::enumerate(indices, strides)) {
if (stride != 1) { // Skip if stride is 1.
Value strideValue =
ShapedType::isDynamic(stride)
? convertUnsignedToI32(rewriter, loc,
memRefDescriptor.stride(rewriter, loc, i))
: LLVM::ConstantOp::create(rewriter, loc, i32, stride);
increment = LLVM::MulOp::create(rewriter, loc, increment, strideValue);
}
index = index ? LLVM::AddOp::create(rewriter, loc, index, increment)
: increment;
}
return index ? index : createI32Constant(rewriter, loc, 0);
}
/// Compute the contents of the `num_records` field for a given memref
/// descriptor - that is, the number of bytes that's one element past the
/// greatest possible valid index into the memref.
static Value getNumRecords(ConversionPatternRewriter &rewriter, Location loc,
MemRefType memrefType,
MemRefDescriptor &memrefDescriptor,
ArrayRef<int64_t> strides, int64_t elementByteWidth,
amdgpu::Chipset chipset, bool boundsCheck) {
if (chipset >= kGfx1250 && !boundsCheck) {
constexpr int64_t first45bits = (1ll << 45) - 1;
return createI64Constant(rewriter, loc, first45bits);
}
if (memrefType.hasStaticShape() &&
!llvm::any_of(strides, ShapedType::isDynamic)) {
int64_t size = memrefType.getRank() == 0 ? 1 : 0;
ArrayRef<int64_t> shape = memrefType.getShape();
for (uint32_t i = 0, e = memrefType.getRank(); i < e; ++i)
size = std::max(shape[i] * strides[i], size);
size = size * elementByteWidth;
return createI64Constant(rewriter, loc, size);
}
Value maxIndex;
for (uint32_t i = 0, e = memrefType.getRank(); i < e; ++i) {
Value size = memrefDescriptor.size(rewriter, loc, i);
Value stride = memrefDescriptor.stride(rewriter, loc, i);
Value maxThisDim = LLVM::MulOp::create(rewriter, loc, size, stride);
maxIndex = maxIndex
? LLVM::UMaxOp::create(rewriter, loc, maxIndex, maxThisDim)
: maxThisDim;
}
Value maxIndexI64 = convertUnsignedToI64(rewriter, loc, maxIndex);
Value byteWidthConst = createI64Constant(rewriter, loc, elementByteWidth);
return LLVM::MulOp::create(rewriter, loc, maxIndexI64, byteWidthConst);
}
static Value makeBufferRsrc(ConversionPatternRewriter &rewriter, Location loc,
Value basePointer, Value numRecords,
bool boundsCheck, amdgpu::Chipset chipset,
Value cacheSwizzleStride = nullptr,
unsigned addressSpace = 8) {
// The stride value is generally 0. However, on MI-300 and onward, you can
// enable a cache swizzling mode by setting bit 14 of the stride field
// and setting that stride to a cache stride.
Type i16 = rewriter.getI16Type();
Value stride;
if (chipset.majorVersion == 9 && chipset >= kGfx942 && cacheSwizzleStride) {
Value cacheStrideZext =
LLVM::ZExtOp::create(rewriter, loc, i16, cacheSwizzleStride);
Value swizzleBit = LLVM::ConstantOp::create(
rewriter, loc, i16, rewriter.getI16IntegerAttr(1 << 14));
stride = LLVM::OrOp::create(rewriter, loc, cacheStrideZext, swizzleBit,
/*isDisjoint=*/true);
} else {
stride = LLVM::ConstantOp::create(rewriter, loc, i16,
rewriter.getI16IntegerAttr(0));
}
uint32_t flags = 0;
if (chipset >= kGfx1250) {
// Flag word:
// bit 0: swizzle
// bit 1: 0 means (total_offset + payload > numRecords)
// 1 means ((total_offset + payload >) numRecords) || ((offset +
// payload) > stride) only applied when swizzle_enable = 0. keep at
// zero.
// whether oob is done depends on numRecords.
// bits 2-3: Type (must be 0)
} else {
// Get the number of elements.
// Flag word:
// bits 0-11: dst sel, ignored by these intrinsics
// bits 12-14: data format (ignored, must be nonzero, 7=float)
// bits 15-18: data format (ignored, must be nonzero, 4=32bit)
// bit 19: In nested heap (0 here)
// bit 20: Behavior on unmap (0 means "return 0 / ignore")
// bits 21-22: Index stride for swizzles (N/A)
// bit 23: Add thread ID (0)
// bit 24: Reserved to 1 (RDNA) or 0 (CDNA)
// bits 25-26: Reserved (0)
// bit 27: Buffer is non-volatile (CDNA only)
// bits 28-29: Out of bounds select (0 = structured, 1 = check index, 2 =
// none, 3 = either swizzles or testing against offset field) RDNA only
// bits 30-31: Type (must be 0)
flags |= (7 << 12) | (4 << 15);
if (chipset.majorVersion >= 10) {
flags |= (1 << 24);
uint32_t oob = boundsCheck ? 3 : 2;
flags |= (oob << 28);
}
}
Value flagsConst = createI32Constant(rewriter, loc, flags);
Type rsrcType =
LLVM::LLVMPointerType::get(rewriter.getContext(), addressSpace);
Value resource = rewriter.createOrFold<ROCDL::MakeBufferRsrcOp>(
loc, rsrcType, basePointer, stride, numRecords, flagsConst);
return resource;
}
namespace {
struct FatRawBufferCastLowering
: public ConvertOpToLLVMPattern<FatRawBufferCastOp> {
FatRawBufferCastLowering(const LLVMTypeConverter &converter, Chipset chipset)
: ConvertOpToLLVMPattern<FatRawBufferCastOp>(converter),
chipset(chipset) {}
Chipset chipset;
LogicalResult
matchAndRewrite(FatRawBufferCastOp op, FatRawBufferCastOpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
Location loc = op.getLoc();
Value memRef = adaptor.getSource();
Value unconvertedMemref = op.getSource();
MemRefType memrefType = cast<MemRefType>(unconvertedMemref.getType());
MemRefDescriptor descriptor(memRef);
DataLayout dataLayout = DataLayout::closest(op);
int64_t elementByteWidth =
dataLayout.getTypeSizeInBits(memrefType.getElementType()) / 8;
int64_t unusedOffset = 0;
SmallVector<int64_t, 5> strideVals;
if (failed(memrefType.getStridesAndOffset(strideVals, unusedOffset)))
return op.emitOpError("Can't lower non-stride-offset memrefs");
Value numRecords = adaptor.getValidBytes();
if (!numRecords)
numRecords =
getNumRecords(rewriter, loc, memrefType, descriptor, strideVals,
elementByteWidth, chipset, adaptor.getBoundsCheck());
Value basePointer =
adaptor.getResetOffset()
? descriptor.bufferPtr(rewriter, loc, *getTypeConverter(),
memrefType)
: descriptor.alignedPtr(rewriter, loc);
Value offset = adaptor.getResetOffset()
? LLVM::ConstantOp::create(rewriter, loc, getIndexType(),
rewriter.getIndexAttr(0))
: descriptor.offset(rewriter, loc);
bool hasSizes = memrefType.getRank() > 0;
// No need to unpack() and pack() all the individual sizes and strides,
// so we'll just extract the arrays.
Value sizes = hasSizes
? LLVM::ExtractValueOp::create(rewriter, loc, descriptor,
kSizePosInMemRefDescriptor)
: Value{};
Value strides =
hasSizes ? LLVM::ExtractValueOp::create(rewriter, loc, descriptor,
kStridePosInMemRefDescriptor)
: Value{};
Value fatPtr = makeBufferRsrc(
rewriter, loc, basePointer, numRecords, adaptor.getBoundsCheck(),
chipset, adaptor.getCacheSwizzleStride(), /*addressSpace=*/7);
Value result = MemRefDescriptor::poison(
rewriter, loc,
getTypeConverter()->convertType(op.getResult().getType()));
SmallVector<int64_t> pos{kAllocatedPtrPosInMemRefDescriptor};
result = LLVM::InsertValueOp::create(rewriter, loc, result, fatPtr, pos);
result = LLVM::InsertValueOp::create(rewriter, loc, result, fatPtr,
kAlignedPtrPosInMemRefDescriptor);
result = LLVM::InsertValueOp::create(rewriter, loc, result, offset,
kOffsetPosInMemRefDescriptor);
if (hasSizes) {
result = LLVM::InsertValueOp::create(rewriter, loc, result, sizes,
kSizePosInMemRefDescriptor);
result = LLVM::InsertValueOp::create(rewriter, loc, result, strides,
kStridePosInMemRefDescriptor);
}
rewriter.replaceOp(op, result);
return success();
}
};
/// Define lowering patterns for raw buffer ops
template <typename GpuOp, typename Intrinsic>
struct RawBufferOpLowering : public ConvertOpToLLVMPattern<GpuOp> {
RawBufferOpLowering(const LLVMTypeConverter &converter, Chipset chipset)
: ConvertOpToLLVMPattern<GpuOp>(converter), chipset(chipset) {}
Chipset chipset;
static constexpr uint32_t maxVectorOpWidth = 128;
LogicalResult
matchAndRewrite(GpuOp gpuOp, typename GpuOp::Adaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
Location loc = gpuOp.getLoc();
Value memref = adaptor.getMemref();
Value unconvertedMemref = gpuOp.getMemref();
MemRefType memrefType = cast<MemRefType>(unconvertedMemref.getType());
if (chipset.majorVersion < 9)
return gpuOp.emitOpError("raw buffer ops require GCN or higher");
Value storeData = adaptor.getODSOperands(0)[0];
if (storeData == memref) // no write component to this op
storeData = Value();
Type wantedDataType;
if (storeData)
wantedDataType = storeData.getType();
else
wantedDataType = gpuOp.getODSResults(0)[0].getType();
Value atomicCmpData = Value();
// Operand index 1 of a load is the indices, trying to read them can crash.
if (storeData) {
Value maybeCmpData = adaptor.getODSOperands(1)[0];
if (maybeCmpData != memref)
atomicCmpData = maybeCmpData;
}
Type llvmWantedDataType = this->typeConverter->convertType(wantedDataType);
Type i32 = rewriter.getI32Type();
// Get the type size in bytes.
DataLayout dataLayout = DataLayout::closest(gpuOp);
int64_t elementByteWidth =
dataLayout.getTypeSizeInBits(memrefType.getElementType()) / 8;
Value byteWidthConst = createI32Constant(rewriter, loc, elementByteWidth);
// If we want to load a vector<NxT> with total size <= 32
// bits, use a scalar load and bitcast it. Similarly, if bitsize(T) < 32
// and the total load size is >= 32, use a vector load of N / (bitsize(T) /
// 32) x i32 and bitcast. Also, the CAS intrinsic requires integer operands,
// so bitcast any floats to integers.
Type llvmBufferValType = llvmWantedDataType;
if (atomicCmpData) {
if (auto floatType = dyn_cast<FloatType>(wantedDataType))
llvmBufferValType = this->getTypeConverter()->convertType(
rewriter.getIntegerType(floatType.getWidth()));
}
if (auto dataVector = dyn_cast<VectorType>(wantedDataType)) {
uint32_t vecLen = dataVector.getNumElements();
uint32_t elemBits =
dataLayout.getTypeSizeInBits(dataVector.getElementType());
uint32_t totalBits = elemBits * vecLen;
bool usePackedFp16 =
isa_and_present<RawBufferAtomicFaddOp>(*gpuOp) && vecLen == 2;
if (totalBits > maxVectorOpWidth)
return gpuOp.emitOpError(
"Total width of loads or stores must be no more than " +
Twine(maxVectorOpWidth) + " bits, but we call for " +
Twine(totalBits) +
" bits. This should've been caught in validation");
if (!usePackedFp16 && elemBits < 32) {
if (totalBits > 32) {
if (totalBits % 32 != 0)
return gpuOp.emitOpError("Load or store of more than 32-bits that "
"doesn't fit into words. Can't happen\n");
llvmBufferValType = this->typeConverter->convertType(
VectorType::get(totalBits / 32, i32));
} else {
llvmBufferValType = this->typeConverter->convertType(
rewriter.getIntegerType(totalBits));
}
}
}
if (auto vecType = dyn_cast<VectorType>(llvmBufferValType)) {
// Buffer intrinsics doesn't support 1-element vectors, cast them to
// scalars.
if (vecType.getNumElements() == 1)
llvmBufferValType = vecType.getElementType();
}
SmallVector<Value, 6> args;
if (storeData) {
if (llvmBufferValType != llvmWantedDataType) {
Value castForStore = LLVM::BitcastOp::create(
rewriter, loc, llvmBufferValType, storeData);
args.push_back(castForStore);
} else {
args.push_back(storeData);
}
}
if (atomicCmpData) {
if (llvmBufferValType != llvmWantedDataType) {
Value castForCmp = LLVM::BitcastOp::create(
rewriter, loc, llvmBufferValType, atomicCmpData);
args.push_back(castForCmp);
} else {
args.push_back(atomicCmpData);
}
}
// Construct buffer descriptor from memref, attributes
int64_t offset = 0;
SmallVector<int64_t, 5> strides;
if (failed(memrefType.getStridesAndOffset(strides, offset)))
return gpuOp.emitOpError("Can't lower non-stride-offset memrefs");
MemRefDescriptor memrefDescriptor(memref);
Value ptr = memrefDescriptor.bufferPtr(
rewriter, loc, *this->getTypeConverter(), memrefType);
Value numRecords =
getNumRecords(rewriter, loc, memrefType, memrefDescriptor, strides,
elementByteWidth, chipset, adaptor.getBoundsCheck());
Value resource = makeBufferRsrc(rewriter, loc, ptr, numRecords,
adaptor.getBoundsCheck(), chipset);
args.push_back(resource);
// Indexing (voffset)
Value voffset = getLinearIndexI32(rewriter, loc, memrefDescriptor,
adaptor.getIndices(), strides);
if (std::optional<int32_t> indexOffset = adaptor.getIndexOffset();
indexOffset && *indexOffset > 0) {
Value extraOffsetConst = createI32Constant(rewriter, loc, *indexOffset);
voffset = voffset ? LLVM::AddOp::create(rewriter, loc, voffset,
extraOffsetConst)
: extraOffsetConst;
}
voffset = LLVM::MulOp::create(rewriter, loc, voffset, byteWidthConst);
args.push_back(voffset);
// SGPR offset.
Value sgprOffset = adaptor.getSgprOffset();
if (!sgprOffset)
sgprOffset = createI32Constant(rewriter, loc, 0);
sgprOffset = LLVM::MulOp::create(rewriter, loc, sgprOffset, byteWidthConst);
args.push_back(sgprOffset);
// bit 0: GLC = 0 (atomics drop value, less coherency)
// bits 1-2: SLC, DLC = 0 (similarly)
// bit 3: swizzled (0 for raw)
args.push_back(createI32Constant(rewriter, loc, 0));
llvm::SmallVector<Type, 1> resultTypes(gpuOp->getNumResults(),
llvmBufferValType);
Operation *lowered = Intrinsic::create(rewriter, loc, resultTypes, args,
ArrayRef<NamedAttribute>());
if (lowered->getNumResults() == 1) {
Value replacement = lowered->getResult(0);
if (llvmBufferValType != llvmWantedDataType) {
replacement = LLVM::BitcastOp::create(rewriter, loc, llvmWantedDataType,
replacement);
}
rewriter.replaceOp(gpuOp, replacement);
} else {
rewriter.eraseOp(gpuOp);
}
return success();
}
};
// TODO: AMDGPU backend already have all this bitpacking logic, we should move
// it to some common place.
/// Vmcnt, Expcnt and Lgkmcnt are decoded as follows:
/// Vmcnt = Waitcnt[3:0] (pre-gfx9)
/// Vmcnt = Waitcnt[15:14,3:0] (gfx9,10)
/// Vmcnt = Waitcnt[15:10] (gfx11)
/// Expcnt = Waitcnt[6:4] (pre-gfx11)
/// Expcnt = Waitcnt[2:0] (gfx11)
/// Lgkmcnt = Waitcnt[11:8] (pre-gfx10)
/// Lgkmcnt = Waitcnt[13:8] (gfx10)
/// Lgkmcnt = Waitcnt[9:4] (gfx11)
static FailureOr<unsigned> encodeWaitcnt(Chipset chipset, unsigned vmcnt,
unsigned expcnt, unsigned lgkmcnt) {
if (chipset.majorVersion < 9) {
vmcnt = std::min(15u, vmcnt);
expcnt = std::min(7u, expcnt);
lgkmcnt = std::min(15u, lgkmcnt);
return vmcnt | (expcnt << 4) | (lgkmcnt << 8);
}
if (chipset.majorVersion == 9) {
vmcnt = std::min(63u, vmcnt);
expcnt = std::min(7u, expcnt);
lgkmcnt = std::min(15u, lgkmcnt);
unsigned lowBits = vmcnt & 0xF;
unsigned highBits = (vmcnt >> 4) << 14;
unsigned otherCnts = (expcnt << 4) | (lgkmcnt << 8);
return lowBits | highBits | otherCnts;
}
if (chipset.majorVersion == 10) {
vmcnt = std::min(63u, vmcnt);
expcnt = std::min(7u, expcnt);
lgkmcnt = std::min(63u, lgkmcnt);
unsigned lowBits = vmcnt & 0xF;
unsigned highBits = (vmcnt >> 4) << 14;
unsigned otherCnts = (expcnt << 4) | (lgkmcnt << 8);
return lowBits | highBits | otherCnts;
}
if (chipset.majorVersion == 11) {
vmcnt = std::min(63u, vmcnt);
expcnt = std::min(7u, expcnt);
lgkmcnt = std::min(63u, lgkmcnt);
return (vmcnt << 10) | expcnt | (lgkmcnt << 4);
}
return failure();
}
struct MemoryCounterWaitOpLowering
: public ConvertOpToLLVMPattern<MemoryCounterWaitOp> {
MemoryCounterWaitOpLowering(const LLVMTypeConverter &converter,
Chipset chipset)
: ConvertOpToLLVMPattern<MemoryCounterWaitOp>(converter),
chipset(chipset) {}
Chipset chipset;
LogicalResult
matchAndRewrite(MemoryCounterWaitOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
if (chipset.majorVersion >= 12) {
Location loc = op.getLoc();
if (std::optional<int> ds = adaptor.getDs())
ROCDL::WaitDscntOp::create(rewriter, loc, *ds);
if (std::optional<int> load = adaptor.getLoad())
ROCDL::WaitLoadcntOp::create(rewriter, loc, *load);
if (std::optional<int> store = adaptor.getStore())
ROCDL::WaitStorecntOp::create(rewriter, loc, *store);
if (std::optional<int> exp = adaptor.getExp())
ROCDL::WaitExpcntOp::create(rewriter, loc, *exp);
if (std::optional<int> tensor = adaptor.getTensor())
ROCDL::WaitTensorcntOp::create(rewriter, loc, *tensor);
rewriter.eraseOp(op);
return success();
}
if (adaptor.getTensor())
return op.emitOpError("unsupported chipset");
auto getVal = [](Attribute attr) -> unsigned {
if (attr)
return cast<IntegerAttr>(attr).getInt();
// This value will be clamped to the maximum value for the chipset.
return 1024;
};
unsigned ds = getVal(adaptor.getDsAttr());
unsigned exp = getVal(adaptor.getExpAttr());
unsigned vmcnt = 1024;
Attribute load = adaptor.getLoadAttr();
Attribute store = adaptor.getStoreAttr();
if (load && store) {
vmcnt = getVal(load) + getVal(store);
} else if (load) {
vmcnt = getVal(load);
} else if (store) {
vmcnt = getVal(store);
}
FailureOr<unsigned> waitcnt = encodeWaitcnt(chipset, vmcnt, exp, ds);
if (failed(waitcnt))
return op.emitOpError("unsupported chipset");
rewriter.replaceOpWithNewOp<ROCDL::SWaitcntOp>(op, *waitcnt);
return success();
}
};
struct LDSBarrierOpLowering : public ConvertOpToLLVMPattern<LDSBarrierOp> {
LDSBarrierOpLowering(const LLVMTypeConverter &converter, Chipset chipset)
: ConvertOpToLLVMPattern<LDSBarrierOp>(converter), chipset(chipset) {}
Chipset chipset;
LogicalResult
matchAndRewrite(LDSBarrierOp op, LDSBarrierOp::Adaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
Location loc = op.getLoc();
// This ensures that waits on global memory aren't introduced on
// chips that don't have the BackOffBarrier feature enabled in LLVM.
bool requiresInlineAsm = chipset < kGfx90a;
Attribute mmra =
rewriter.getAttr<LLVM::MMRATagAttr>("amdgpu-synchronize-as", "local");
// Note: while there *is* a workgroup-one-as scope, this, when combined with
// the MMRA, will lead to the fence having no effect. This is because the
// codepaths for an atomic load or store will observe that a
// one-address-space atomic to LDS requires no synchronization because
// operations on LDS are totally ordered with respect to each other, and so
// will not emit the correct waitcnt operations that these fences are
// intended to produce. Therefore, we use a broader type of fence and rely
// on the MMRA to relax it to the semantics we want.
StringRef scope = "workgroup";
auto relFence = LLVM::FenceOp::create(rewriter, loc,
LLVM::AtomicOrdering::release, scope);
relFence->setDiscardableAttr(LLVM::LLVMDialect::getMmraAttrName(), mmra);
if (requiresInlineAsm) {
auto asmDialectAttr = LLVM::AsmDialectAttr::get(rewriter.getContext(),
LLVM::AsmDialect::AD_ATT);
const char *asmStr = ";;;WARNING: BREAKS DEBUG WATCHES\ns_barrier";
const char *constraints = "";
LLVM::InlineAsmOp::create(
rewriter, loc,
/*resultTypes=*/TypeRange(), /*operands=*/ValueRange(),
/*asm_string=*/asmStr, constraints, /*has_side_effects=*/true,
/*is_align_stack=*/false, LLVM::TailCallKind::None,
/*asm_dialect=*/asmDialectAttr,
/*operand_attrs=*/ArrayAttr());
} else if (chipset.majorVersion < 12) {
ROCDL::SBarrierOp::create(rewriter, loc);
} else {
ROCDL::BarrierSignalOp::create(rewriter, loc, -1);
ROCDL::BarrierWaitOp::create(rewriter, loc, -1);
}
auto acqFence = LLVM::FenceOp::create(rewriter, loc,
LLVM::AtomicOrdering::acquire, scope);
acqFence->setDiscardableAttr(LLVM::LLVMDialect::getMmraAttrName(), mmra);
rewriter.replaceOp(op, acqFence);
return success();
}
};
struct SchedBarrierOpLowering : public ConvertOpToLLVMPattern<SchedBarrierOp> {
SchedBarrierOpLowering(const LLVMTypeConverter &converter, Chipset chipset)
: ConvertOpToLLVMPattern<SchedBarrierOp>(converter), chipset(chipset) {}
Chipset chipset;
LogicalResult
matchAndRewrite(SchedBarrierOp op, SchedBarrierOp::Adaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
rewriter.replaceOpWithNewOp<ROCDL::SchedBarrier>(op,
(uint32_t)op.getOpts());
return success();
}
};
} // namespace
/// Pack small float vector operands (fp4/fp6/fp8/bf16) into the format
/// expected by scaled matrix multiply intrinsics (MFMA/WMMA).
///
/// Specifically:
/// 1. If the element type is bfloat16, bitcast it to i16 unless rocdl intrinsic
/// allows bf16. Newer MFMAs support bf16 types on operand, check
/// IntrinsicsAMDGPU.td file for reference.
/// 2. If instead we have a more than 64-bit quantity, use a <N / 4 x i32>
/// instead, which is what the f8f6f4 intrinsics use.
/// 3. If `input` is a vector of N <= 8 bytes, bitcast it to a (N * 8)-bit
/// integer.
///
/// Note that the type of `input` has already been LLVM type converted:
/// therefore 8-bit and smaller floats are represented as their corresponding
/// `iN` integers.
static Value packSmallFloatVectorOperand(ConversionPatternRewriter &rewriter,
Location loc, Value input,
bool allowBf16 = true) {
Type inputType = input.getType();
if (auto vectorType = dyn_cast<VectorType>(inputType)) {
if (vectorType.getElementType().isBF16() && !allowBf16)
return LLVM::BitcastOp::create(
rewriter, loc, vectorType.clone(rewriter.getI16Type()), input);
if (vectorType.getElementType().isInteger(8) &&
vectorType.getNumElements() <= 8)
return LLVM::BitcastOp::create(
rewriter, loc,
rewriter.getIntegerType(vectorType.getNumElements() * 8), input);
if (isa<IntegerType>(vectorType.getElementType()) &&
vectorType.getElementTypeBitWidth() <= 8) {
int64_t numWords = llvm::divideCeil(
vectorType.getNumElements() * vectorType.getElementTypeBitWidth(),
32);
return LLVM::BitcastOp::create(
rewriter, loc, VectorType::get(numWords, rewriter.getI32Type()),
input);
}
}
return input;
}
/// Converts sparse MFMA (smfmac) operands to the expected ROCDL types.
static Value convertSparseMFMAVectorOperand(ConversionPatternRewriter &rewriter,
Location loc, Value input,
bool allowBf16 = true) {
Type inputType = input.getType();
auto vectorType = cast<VectorType>(inputType);
// bf16 -> i16 when not allowed (pre-gfx950).
if (vectorType.getElementType().isBF16() && !allowBf16)
return LLVM::BitcastOp::create(
rewriter, loc, vectorType.clone(rewriter.getI16Type()), input);
// i8/fp8 vectors -> vector<Nxi32>.
if (isa<IntegerType>(vectorType.getElementType()) &&
vectorType.getElementTypeBitWidth() <= 8) {
int64_t numWords = llvm::divideCeil(
vectorType.getNumElements() * vectorType.getElementTypeBitWidth(), 32);
return LLVM::BitcastOp::create(
rewriter, loc, VectorType::get(numWords, rewriter.getI32Type()), input);
}
return input;
}
/// Converts the scaled MFMA/WMMA operands, `scalesA` and `scalesB`, from MLIR
/// AMDGPU dialect convention to ROCDL and LLVM AMDGPU intrinsics convention.
///
/// Specifically:
/// 1. If `input` is a i8 value, zero extend it to i32
/// 2. If `input` is a vector of length 4 or 8 and type i8, cast it to i32
///
/// Note that the type of `input` has already been LLVM type converted:
/// therefore 8-bit and smaller floats are represented as their corresponding
/// `iN` integers.
static Value castScaleOperand(ConversionPatternRewriter &rewriter, Location loc,
Value input) {
return TypeSwitch<Type, Value>(input.getType())
.Case([&](IntegerType) {
// Handle scalar i8: zero extend to i32.
return LLVM::ZExtOp::create(rewriter, loc, rewriter.getI32Type(),
input);
})
.Case([&](VectorType vectorType) {
// Handle vector<4xi8> -> i32 or vector<8xi8> -> i64.
int64_t numElements = vectorType.getNumElements();
assert((numElements == 4 || numElements == 8) &&
"scale operand must be a vector of length 4 or 8");
IntegerType outputType =
(numElements == 4) ? rewriter.getI32Type() : rewriter.getI64Type();
return LLVM::BitcastOp::create(rewriter, loc, outputType, input);
})
.DefaultUnreachable("unexpected input type for scale operand");
}
/// Maps f8 scale element types to WMMA scale format codes.
static std::optional<uint32_t> getWmmaScaleFormat(Type elemType) {
return TypeSwitch<Type, std::optional<uint32_t>>(elemType)
.Case([](Float8E8M0FNUType) { return 0; })
.Case([](Float8E4M3FNType) { return 2; })
.Default(std::nullopt);
}
/// Determines the ROCDL intrinsic name for scaled WMMA based on dimensions
/// and scale block size (16 or 32).
static std::optional<StringRef>
getScaledWmmaIntrinsicName(int64_t m, int64_t n, int64_t k, bool isScale16) {
if (m == 16 && n == 16 && k == 128)
return isScale16
? ROCDL::wmma_scale16_f32_16x16x128_f8f6f4::getOperationName()
: ROCDL::wmma_scale_f32_16x16x128_f8f6f4::getOperationName();
if (m == 32 && n == 16 && k == 128)
return isScale16 ? ROCDL::wmma_scale16_f32_32x16x128_f4::getOperationName()
: ROCDL::wmma_scale_f32_32x16x128_f4::getOperationName();
return std::nullopt;
}
/// Push an input operand. If it is a float type, nothing to do. If it is
/// an integer type, then we need to also push its signdness (1 for signed, 0
/// for unsigned) and we need to pack the input 16xi8 vector into a 4xi32
/// vector (or the 8xi8 vector into a 2xi32 one for gfx12+).
/// We also need to convert bfloat inputs to i16 to account for the bfloat
/// intrinsics having been defined before the AMD backend supported bfloat. We
/// similarly need to pack 8-bit float types into integers as if they were i8
/// (which they are for the backend's purposes).
static void wmmaPushInputOperand(
ConversionPatternRewriter &rewriter, Location loc,
const TypeConverter *typeConverter, bool isUnsigned, Value llvmInput,
Value mlirInput, SmallVectorImpl<Value> &operands,
SmallVectorImpl<NamedAttribute> &attrs, StringRef attrName) {
Type inputType = llvmInput.getType();
auto vectorType = dyn_cast<VectorType>(inputType);
if (!vectorType) {
operands.push_back(llvmInput);
return;
}
Type elemType = vectorType.getElementType();
if (elemType.getIntOrFloatBitWidth() > 8) {
operands.push_back(llvmInput);
return;
}
// We need to check the type of the input before conversion to properly test
// for int8. This is because, in LLVM, fp8 type is converted to int8, so the
// fp8/int8 information is lost during the conversion process.
auto mlirInputType = cast<VectorType>(mlirInput.getType());
bool isInputInteger = mlirInputType.getElementType().isInteger();
if (isInputInteger) {
// if element type is 8-bit signed or unsigned, ignore the isUnsigned flag
bool localIsUnsigned = isUnsigned;
if (elemType.isUnsignedInteger()) {
localIsUnsigned = true;
} else if (elemType.isSignedInteger()) {
localIsUnsigned = false;
}
attrs.push_back(
NamedAttribute(attrName, rewriter.getBoolAttr(!localIsUnsigned)));
}
int64_t numBits =
vectorType.getNumElements() * elemType.getIntOrFloatBitWidth();
Type i32 = rewriter.getI32Type();
Type intrinsicInType = numBits <= 32
? (Type)rewriter.getIntegerType(numBits)
: (Type)VectorType::get(numBits / 32, i32);
auto llvmIntrinsicInType = typeConverter->convertType(intrinsicInType);
Value castInput = rewriter.createOrFold<LLVM::BitcastOp>(
loc, llvmIntrinsicInType, llvmInput);
// The wave64-mode 16x16x16 intrinsics that take 4-bit integers only need
// (256 / 64) * 4 = 16 bits of input (on gfx12+) but take i32 arguments.
// Add in the zeros here.
if (numBits < 32)
castInput = LLVM::ZExtOp::create(rewriter, loc, i32, castInput);
operands.push_back(castInput);
}
/// Push the output operand. For many cases this is only pushing the output in
/// the operand list. But when we have f16 -> f16 or bf16 -> bf16 intrinsics,
/// since the same numbers of VGPRs is used, we need to decide if to store the
/// result in the upper 16 bits of the VGPRs or in the lower part. To store the
/// result in the lower 16 bits, set subwordOffset to 1, otherwise result will
/// be stored it in the upper part. The subwordOffset must not be set for gfx12,
/// as the instructions have been changed to return fewer registers instead.
static void wmmaPushOutputOperand(ConversionPatternRewriter &rewriter,
Location loc,
const TypeConverter *typeConverter,
Value output, int32_t subwordOffset,
bool clamp, SmallVectorImpl<Value> &operands,
SmallVectorImpl<NamedAttribute> &attrs) {
Type inputType = output.getType();
auto vectorType = dyn_cast<VectorType>(inputType);
Type elemType = vectorType.getElementType();
operands.push_back(output);
if (elemType.isF16() || elemType.isBF16() || elemType.isInteger(16)) {
attrs.push_back(
NamedAttribute("opsel", rewriter.getBoolAttr(subwordOffset)));
} else if (elemType.isInteger(32)) {
attrs.push_back(NamedAttribute("clamp", rewriter.getBoolAttr(clamp)));
}
}
/// Return true if `type` is the E5M2 variant of an 8-bit float that is
/// supported by the `_bf8` instructions on the given `chipset`.
static bool typeIsExpectedBf8ForChipset(Chipset chipset, Type type) {
return (chipset == kGfx942 && isa<Float8E5M2FNUZType>(type)) ||
(hasOcpFp8(chipset) && isa<Float8E5M2Type>(type));
}
/// Return true if `type` is the E4M3FN variant of an 8-bit float that is
/// supported by the `_fp8` instructions on the given `chipset`.
static bool typeIsExpectedFp8ForChipset(Chipset chipset, Type type) {
return (chipset == kGfx942 && isa<Float8E4M3FNUZType>(type)) ||
(hasOcpFp8(chipset) && isa<Float8E4M3FNType>(type));
}
/// Return the `rocdl` intrinsic corresponding to a MFMA operation `mfma`
/// if one exists. This includes checking to ensure the intrinsic is supported
/// on the architecture you are compiling for.
static std::optional<StringRef> mfmaOpToIntrinsic(MFMAOp mfma,
Chipset chipset) {
uint32_t m = mfma.getM(), n = mfma.getN(), k = mfma.getK(),
b = mfma.getBlocks();
Type sourceElem = getElementTypeOrSelf(mfma.getSourceA().getType());
Type destElem = getElementTypeOrSelf(mfma.getDestC().getType());
if (sourceElem.isF32() && destElem.isF32()) {
if (mfma.getReducePrecision() && chipset >= kGfx942) {
if (m == 32 && n == 32 && k == 4 && b == 1)
return ROCDL::mfma_f32_32x32x4_xf32::getOperationName();
if (m == 16 && n == 16 && k == 8 && b == 1)
return ROCDL::mfma_f32_16x16x8_xf32::getOperationName();
}
if (m == 32 && n == 32 && k == 1 && b == 2)
return ROCDL::mfma_f32_32x32x1f32::getOperationName();
if (m == 16 && n == 16 && k == 1 && b == 4)
return ROCDL::mfma_f32_16x16x1f32::getOperationName();
if (m == 4 && n == 4 && k == 1 && b == 16)
return ROCDL::mfma_f32_4x4x1f32::getOperationName();
if (m == 32 && n == 32 && k == 2 && b == 1)
return ROCDL::mfma_f32_32x32x2f32::getOperationName();
if (m == 16 && n == 16 && k == 4 && b == 1)
return ROCDL::mfma_f32_16x16x4f32::getOperationName();
}
if (sourceElem.isF16() && destElem.isF32()) {
if (chipset >= kGfx950) {
if (m == 32 && n == 32 && k == 16 && b == 1)
return ROCDL::mfma_f32_32x32x16_f16::getOperationName();
if (m == 16 && n == 16 && k == 32 && b == 1)
return ROCDL::mfma_f32_16x16x32_f16::getOperationName();
}
if (m == 32 && n == 32 && k == 4 && b == 2)
return ROCDL::mfma_f32_32x32x4f16::getOperationName();
if (m == 16 && n == 16 && k == 4 && b == 4)
return ROCDL::mfma_f32_16x16x4f16::getOperationName();
if (m == 4 && n == 4 && k == 4 && b == 16)
return ROCDL::mfma_f32_4x4x4f16::getOperationName();
if (m == 32 && n == 32 && k == 8 && b == 1)
return ROCDL::mfma_f32_32x32x8f16::getOperationName();
if (m == 16 && n == 16 && k == 16 && b == 1)
return ROCDL::mfma_f32_16x16x16f16::getOperationName();
}
if (sourceElem.isBF16() && destElem.isF32()) {
if (chipset >= kGfx950) {
if (m == 32 && n == 32 && k == 16 && b == 1)
return ROCDL::mfma_f32_32x32x16_bf16::getOperationName();
if (m == 16 && n == 16 && k == 32 && b == 1)
return ROCDL::mfma_f32_16x16x32_bf16::getOperationName();
}
if (chipset >= kGfx90a) {
if (m == 32 && n == 32 && k == 4 && b == 2)
return ROCDL::mfma_f32_32x32x4bf16_1k::getOperationName();
if (m == 16 && n == 16 && k == 4 && b == 4)
return ROCDL::mfma_f32_16x16x4bf16_1k::getOperationName();
if (m == 4 && n == 4 && k == 4 && b == 16)
return ROCDL::mfma_f32_4x4x4bf16_1k::getOperationName();
if (m == 32 && n == 32 && k == 8 && b == 1)
return ROCDL::mfma_f32_32x32x8bf16_1k::getOperationName();
if (m == 16 && n == 16 && k == 16 && b == 1)
return ROCDL::mfma_f32_16x16x16bf16_1k::getOperationName();
}
if (m == 32 && n == 32 && k == 2 && b == 2)
return ROCDL::mfma_f32_32x32x2bf16::getOperationName();
if (m == 16 && n == 16 && k == 2 && b == 4)
return ROCDL::mfma_f32_16x16x2bf16::getOperationName();
if (m == 4 && n == 4 && k == 2 && b == 16)
return ROCDL::mfma_f32_4x4x2bf16::getOperationName();
if (m == 32 && n == 32 && k == 4 && b == 1)
return ROCDL::mfma_f32_32x32x4bf16::getOperationName();
if (m == 16 && n == 16 && k == 8 && b == 1)
return ROCDL::mfma_f32_16x16x8bf16::getOperationName();
}
if (sourceElem.isInteger(8) && destElem.isInteger(32)) {
if (chipset >= kGfx950) {
if (m == 32 && n == 32 && k == 32 && b == 1)
return ROCDL::mfma_i32_32x32x32_i8::getOperationName();
if (m == 16 && n == 16 && k == 64 && b == 1)
return ROCDL::mfma_i32_16x16x64_i8::getOperationName();
}
if (m == 32 && n == 32 && k == 4 && b == 2)
return ROCDL::mfma_i32_32x32x4i8::getOperationName();
if (m == 16 && n == 16 && k == 4 && b == 4)
return ROCDL::mfma_i32_16x16x4i8::getOperationName();
if (m == 4 && n == 4 && k == 4 && b == 16)
return ROCDL::mfma_i32_4x4x4i8::getOperationName();
if (m == 32 && n == 32 && k == 8 && b == 1)
return ROCDL::mfma_i32_32x32x8i8::getOperationName();
if (m == 16 && n == 16 && k == 16 && b == 1)
return ROCDL::mfma_i32_16x16x16i8::getOperationName();
if (m == 32 && n == 32 && k == 16 && b == 1 && chipset >= kGfx942)
return ROCDL::mfma_i32_32x32x16_i8::getOperationName();
if (m == 16 && n == 16 && k == 32 && b == 1 && chipset >= kGfx942)
return ROCDL::mfma_i32_16x16x32_i8::getOperationName();
}
if (sourceElem.isF64() && destElem.isF64() && chipset >= kGfx90a) {
if (m == 16 && n == 16 && k == 4 && b == 1)
return ROCDL::mfma_f64_16x16x4f64::getOperationName();
if (m == 4 && n == 4 && k == 4 && b == 4)
return ROCDL::mfma_f64_4x4x4f64::getOperationName();
}
if (destElem.isF32() && typeIsExpectedBf8ForChipset(chipset, sourceElem)) {
// Known to be correct because there are no scalar f8 instructions and
// because a length mismatch will have been caught by the verifier.
Type sourceBElem =
cast<VectorType>(mfma.getSourceB().getType()).getElementType();
if (m == 16 && n == 16 && k == 32 && b == 1) {
if (typeIsExpectedBf8ForChipset(chipset, sourceBElem))
return ROCDL::mfma_f32_16x16x32_bf8_bf8::getOperationName();
if (typeIsExpectedFp8ForChipset(chipset, sourceBElem))
return ROCDL::mfma_f32_16x16x32_bf8_fp8::getOperationName();
}
if (m == 32 && n == 32 && k == 16 && b == 1) {
if (typeIsExpectedBf8ForChipset(chipset, sourceBElem))
return ROCDL::mfma_f32_32x32x16_bf8_bf8::getOperationName();
if (typeIsExpectedFp8ForChipset(chipset, sourceBElem))
return ROCDL::mfma_f32_32x32x16_bf8_fp8::getOperationName();
}
}
if (destElem.isF32() && typeIsExpectedFp8ForChipset(chipset, sourceElem)) {
Type sourceBElem =
cast<VectorType>(mfma.getSourceB().getType()).getElementType();
if (m == 16 && n == 16 && k == 32 && b == 1) {
if (typeIsExpectedBf8ForChipset(chipset, sourceBElem))
return ROCDL::mfma_f32_16x16x32_fp8_bf8::getOperationName();
if (typeIsExpectedFp8ForChipset(chipset, sourceBElem))
return ROCDL::mfma_f32_16x16x32_fp8_fp8::getOperationName();
}
if (m == 32 && n == 32 && k == 16 && b == 1) {
if (typeIsExpectedBf8ForChipset(chipset, sourceBElem))
return ROCDL::mfma_f32_32x32x16_fp8_bf8::getOperationName();
if (typeIsExpectedFp8ForChipset(chipset, sourceBElem))
return ROCDL::mfma_f32_32x32x16_fp8_fp8::getOperationName();
}
}
return std::nullopt;
}
static std::optional<uint32_t> smallFloatTypeToFormatCode(Type mlirElemType) {
return llvm::TypeSwitch<Type, std::optional<uint32_t>>(mlirElemType)
.Case([](Float8E4M3FNType) { return 0u; })
.Case([](Float8E5M2Type) { return 1u; })
.Case([](Float6E2M3FNType) { return 2u; })
.Case([](Float6E3M2FNType) { return 3u; })
.Case([](Float4E2M1FNType) { return 4u; })
.Default(std::nullopt);
}
/// If there is a scaled MFMA instruction for the input element types `aType`
/// and `bType`, output type `destType`, problem size M, N, K, and B (number of
/// blocks) on the given `chipset`, return a tuple consisting of the
/// OperationName of the intrinsic and the type codes that need to be passed to
/// that intrinsic. Note that this is also used to implement some un-scaled
/// MFMAs, since the compiler represents the ordinary instruction as a "scaled"
/// MFMA with a scale of 0.
static std::optional<std::tuple<StringRef, uint32_t, uint32_t>>
mfmaOpToScaledIntrinsic(Type aType, Type bType, Type destType, uint32_t m,
uint32_t n, uint32_t k, uint32_t b, Chipset chipset) {
aType = getElementTypeOrSelf(aType);
bType = getElementTypeOrSelf(bType);
destType = getElementTypeOrSelf(destType);
if (chipset < kGfx950)
return std::nullopt;
if (!isa<Float32Type>(destType))
return std::nullopt;
std::optional<uint32_t> aTypeCode = smallFloatTypeToFormatCode(aType);
std::optional<uint32_t> bTypeCode = smallFloatTypeToFormatCode(bType);
if (!aTypeCode || !bTypeCode)
return std::nullopt;
if (m == 32 && n == 32 && k == 64 && b == 1)
return std::tuple{ROCDL::mfma_scale_f32_32x32x64_f8f6f4::getOperationName(),
*aTypeCode, *bTypeCode};
if (m == 16 && n == 16 && k == 128 && b == 1)
return std::tuple{
ROCDL::mfma_scale_f32_16x16x128_f8f6f4::getOperationName(), *aTypeCode,
*bTypeCode};
return std::nullopt;
}
static std::optional<std::tuple<StringRef, uint32_t, uint32_t>>
mfmaOpToScaledIntrinsic(MFMAOp mfma, Chipset chipset) {
return mfmaOpToScaledIntrinsic(
mfma.getSourceA().getType(), mfma.getSourceB().getType(),
mfma.getDestC().getType(), mfma.getM(), mfma.getN(), mfma.getK(),
mfma.getBlocks(), chipset);
}
static std::optional<std::tuple<StringRef, uint32_t, uint32_t>>
mfmaOpToScaledIntrinsic(ScaledMFMAOp smfma, Chipset chipset) {
return mfmaOpToScaledIntrinsic(smfma.getSourceA().getType(),
smfma.getSourceB().getType(),
smfma.getDestC().getType(), smfma.getM(),
smfma.getN(), smfma.getK(), 1u, chipset);
}
/// Returns the `rocdl` intrinsic corresponding to a WMMA operation `wmma`
/// for RDNA3/4 architectures.
static std::optional<StringRef>
wmmaOpToIntrinsicRDNA(Type elemSourceType, Type elemBSourceType,
Type elemDestType, uint32_t k, bool isRDNA3) {
using fp8 = Float8E4M3FNType;
using bf8 = Float8E5M2Type;
// Handle k == 16 for RDNA3/4.
if (k == 16) {
// Common patterns for RDNA3 and RDNA4.
if (elemSourceType.isF16() && elemDestType.isF32())
return ROCDL::wmma_f32_16x16x16_f16::getOperationName();
if (elemSourceType.isBF16() && elemDestType.isF32())
return ROCDL::wmma_f32_16x16x16_bf16::getOperationName();
if (elemSourceType.isF16() && elemDestType.isF16())
return ROCDL::wmma_f16_16x16x16_f16::getOperationName();
if (elemSourceType.isBF16() && elemDestType.isBF16())
return ROCDL::wmma_bf16_16x16x16_bf16::getOperationName();
if (elemSourceType.isInteger(8) && elemDestType.isInteger(32))
return ROCDL::wmma_i32_16x16x16_iu8::getOperationName();
// RDNA3 specific patterns.
if (isRDNA3) {
if (elemSourceType.isInteger(4) && elemDestType.isInteger(32))
return ROCDL::wmma_i32_16x16x16_iu4::getOperationName();
return std::nullopt;
}
// RDNA4 specific patterns (fp8/bf8).
if (isa<fp8>(elemSourceType) && isa<fp8>(elemBSourceType) &&
elemDestType.isF32())
return ROCDL::wmma_f32_16x16x16_fp8_fp8::getOperationName();
if (isa<fp8>(elemSourceType) && isa<bf8>(elemBSourceType) &&
elemDestType.isF32())
return ROCDL::wmma_f32_16x16x16_fp8_bf8::getOperationName();
if (isa<bf8>(elemSourceType) && isa<bf8>(elemBSourceType) &&
elemDestType.isF32())
return ROCDL::wmma_f32_16x16x16_bf8_bf8::getOperationName();
if (isa<bf8>(elemSourceType) && isa<fp8>(elemBSourceType) &&
elemDestType.isF32())
return ROCDL::wmma_f32_16x16x16_bf8_fp8::getOperationName();
if (elemSourceType.isInteger(4) && elemDestType.isInteger(32))
return ROCDL::wmma_i32_16x16x16_iu4::getOperationName();
return std::nullopt;
}
// Handle k == 32 for RDNA4.
if (k == 32 && !isRDNA3) {
if (elemSourceType.isInteger(4) && elemDestType.isInteger(32))
return ROCDL::wmma_i32_16x16x32_iu4::getOperationName();
}
return std::nullopt;
}
/// Return the `rocdl` intrinsic corresponding to a WMMA operation `wmma`
/// for the gfx1250 architecture.
static std::optional<StringRef> wmmaOpToIntrinsicGfx1250(Type elemSourceType,
Type elemBSourceType,
Type elemDestType,
uint32_t k) {
using fp8 = Float8E4M3FNType;
using bf8 = Float8E5M2Type;
if (k == 4) {
if (elemSourceType.isF32() && elemDestType.isF32())
return ROCDL::wmma_f32_16x16x4_f32::getOperationName();
return std::nullopt;
}
if (k == 32) {
if (elemSourceType.isF16() && elemDestType.isF32())
return ROCDL::wmma_f32_16x16x32_f16::getOperationName();
if (elemSourceType.isBF16() && elemDestType.isF32())
return ROCDL::wmma_f32_16x16x32_bf16::getOperationName();
if (elemSourceType.isF16() && elemDestType.isF16())
return ROCDL::wmma_f16_16x16x32_f16::getOperationName();
if (elemSourceType.isBF16() && elemDestType.isBF16())
return ROCDL::wmma_bf16_16x16x32_bf16::getOperationName();
return std::nullopt;
}
if (k == 64) {
if (isa<fp8>(elemSourceType) && isa<fp8>(elemBSourceType)) {
if (elemDestType.isF32())
return ROCDL::wmma_f32_16x16x64_fp8_fp8::getOperationName();
if (elemDestType.isF16())
return ROCDL::wmma_f16_16x16x64_fp8_fp8::getOperationName();
}
if (isa<fp8>(elemSourceType) && isa<bf8>(elemBSourceType)) {
if (elemDestType.isF32())
return ROCDL::wmma_f32_16x16x64_fp8_bf8::getOperationName();
if (elemDestType.isF16())
return ROCDL::wmma_f16_16x16x64_fp8_bf8::getOperationName();
}
if (isa<bf8>(elemSourceType) && isa<bf8>(elemBSourceType)) {
if (elemDestType.isF32())
return ROCDL::wmma_f32_16x16x64_bf8_bf8::getOperationName();
if (elemDestType.isF16())
return ROCDL::wmma_f16_16x16x64_bf8_bf8::getOperationName();
}
if (isa<bf8>(elemSourceType) && isa<fp8>(elemBSourceType)) {
if (elemDestType.isF32())
return ROCDL::wmma_f32_16x16x64_bf8_fp8::getOperationName();
if (elemDestType.isF16())
return ROCDL::wmma_f16_16x16x64_bf8_fp8::getOperationName();
}
if (elemSourceType.isInteger(8) && elemDestType.isInteger(32))
return ROCDL::wmma_i32_16x16x64_iu8::getOperationName();
return std::nullopt;
}
if (k == 128) {
if (isa<fp8>(elemSourceType) && isa<fp8>(elemBSourceType)) {
if (elemDestType.isF32())
return ROCDL::wmma_f32_16x16x128_fp8_fp8::getOperationName();
if (elemDestType.isF16())
return ROCDL::wmma_f16_16x16x128_fp8_fp8::getOperationName();
}
if (isa<fp8>(elemSourceType) && isa<bf8>(elemBSourceType)) {
if (elemDestType.isF32())
return ROCDL::wmma_f32_16x16x128_fp8_bf8::getOperationName();
if (elemDestType.isF16())
return ROCDL::wmma_f16_16x16x128_fp8_bf8::getOperationName();
}
if (isa<bf8>(elemSourceType) && isa<bf8>(elemBSourceType)) {
if (elemDestType.isF32())
return ROCDL::wmma_f32_16x16x128_bf8_bf8::getOperationName();
if (elemDestType.isF16())
return ROCDL::wmma_f16_16x16x128_bf8_bf8::getOperationName();
}
if (isa<bf8>(elemSourceType) && isa<fp8>(elemBSourceType)) {
if (elemDestType.isF32())
return ROCDL::wmma_f32_16x16x128_bf8_fp8::getOperationName();
if (elemDestType.isF16())
return ROCDL::wmma_f16_16x16x128_bf8_fp8::getOperationName();
}
return std::nullopt;
}
return std::nullopt;
}
/// Returns the `rocdl` intrinsic corresponding to a SparseMFMA (smfmac)
/// operation if one exists. This includes checking to ensure the intrinsic is
/// supported on the architecture you are compiling for.
static std::optional<StringRef> smfmacOpToIntrinsic(SparseMFMAOp op,
Chipset chipset) {
bool isGfx950 = chipset >= kGfx950;
auto isFp8 = [&](Type t) { return typeIsExpectedFp8ForChipset(chipset, t); };
auto isBf8 = [&](Type t) { return typeIsExpectedBf8ForChipset(chipset, t); };
uint32_t m = op.getM(), n = op.getN(), k = op.getK();
Type sourceAElem = getElementTypeOrSelf(op.getSourceA().getType());
Type sourceBElem = getElementTypeOrSelf(op.getSourceB().getType());
Type destElem = getElementTypeOrSelf(op.getDestC().getType());
if (m == 16 && n == 16 && k == 32) {
if (sourceAElem.isF16() && sourceBElem.isF16() && destElem.isF32())
return ROCDL::smfmac_f32_16x16x32_f16::getOperationName();
if (sourceAElem.isBF16() && sourceBElem.isBF16() && destElem.isF32())
return ROCDL::smfmac_f32_16x16x32_bf16::getOperationName();
}
if (m == 16 && n == 16 && k == 64) {
if (isGfx950) {
if (sourceAElem.isF16() && sourceBElem.isF16() && destElem.isF32())
return ROCDL::smfmac_f32_16x16x64_f16::getOperationName();
if (sourceAElem.isBF16() && sourceBElem.isBF16() && destElem.isF32())
return ROCDL::smfmac_f32_16x16x64_bf16::getOperationName();
}
if (sourceAElem.isInteger(8) && sourceBElem.isInteger(8) &&
destElem.isInteger(32))
return ROCDL::smfmac_i32_16x16x64_i8::getOperationName();
if (isFp8(sourceAElem) && isFp8(sourceBElem) && destElem.isF32())
return ROCDL::smfmac_f32_16x16x64_fp8_fp8::getOperationName();
if (isFp8(sourceAElem) && isBf8(sourceBElem) && destElem.isF32())
return ROCDL::smfmac_f32_16x16x64_fp8_bf8::getOperationName();
if (isBf8(sourceAElem) && isFp8(sourceBElem) && destElem.isF32())
return ROCDL::smfmac_f32_16x16x64_bf8_fp8::getOperationName();
if (isBf8(sourceAElem) && isBf8(sourceBElem) && destElem.isF32())
return ROCDL::smfmac_f32_16x16x64_bf8_bf8::getOperationName();
}
if (m == 16 && n == 16 && k == 128 && isGfx950) {
if (sourceAElem.isInteger(8) && sourceBElem.isInteger(8) &&
destElem.isInteger(32))
return ROCDL::smfmac_i32_16x16x128_i8::getOperationName();
if (isFp8(sourceAElem) && isFp8(sourceBElem) && destElem.isF32())
return ROCDL::smfmac_f32_16x16x128_fp8_fp8::getOperationName();
if (isFp8(sourceAElem) && isBf8(sourceBElem) && destElem.isF32())
return ROCDL::smfmac_f32_16x16x128_fp8_bf8::getOperationName();
if (isBf8(sourceAElem) && isFp8(sourceBElem) && destElem.isF32())
return ROCDL::smfmac_f32_16x16x128_bf8_fp8::getOperationName();
if (isBf8(sourceAElem) && isBf8(sourceBElem) && destElem.isF32())
return ROCDL::smfmac_f32_16x16x128_bf8_bf8::getOperationName();
}
if (m == 32 && n == 32 && k == 16) {
if (sourceAElem.isF16() && sourceBElem.isF16() && destElem.isF32())
return ROCDL::smfmac_f32_32x32x16_f16::getOperationName();
if (sourceAElem.isBF16() && sourceBElem.isBF16() && destElem.isF32())
return ROCDL::smfmac_f32_32x32x16_bf16::getOperationName();
}
if (m == 32 && n == 32 && k == 32) {
if (isGfx950) {
if (sourceAElem.isF16() && sourceBElem.isF16() && destElem.isF32())
return ROCDL::smfmac_f32_32x32x32_f16::getOperationName();
if (sourceAElem.isBF16() && sourceBElem.isBF16() && destElem.isF32())
return ROCDL::smfmac_f32_32x32x32_bf16::getOperationName();
}
if (sourceAElem.isInteger(8) && sourceBElem.isInteger(8) &&
destElem.isInteger(32))
return ROCDL::smfmac_i32_32x32x32_i8::getOperationName();
if (isFp8(sourceAElem) && isFp8(sourceBElem) && destElem.isF32())
return ROCDL::smfmac_f32_32x32x32_fp8_fp8::getOperationName();
if (isFp8(sourceAElem) && isBf8(sourceBElem) && destElem.isF32())
return ROCDL::smfmac_f32_32x32x32_fp8_bf8::getOperationName();
if (isBf8(sourceAElem) && isFp8(sourceBElem) && destElem.isF32())
return ROCDL::smfmac_f32_32x32x32_bf8_fp8::getOperationName();
if (isBf8(sourceAElem) && isBf8(sourceBElem) && destElem.isF32())
return ROCDL::smfmac_f32_32x32x32_bf8_bf8::getOperationName();
}
if (m == 32 && n == 32 && k == 64 && isGfx950) {
if (sourceAElem.isInteger(8) && sourceBElem.isInteger(8) &&
destElem.isInteger(32))
return ROCDL::smfmac_i32_32x32x64_i8::getOperationName();
if (isFp8(sourceAElem) && isFp8(sourceBElem) && destElem.isF32())
return ROCDL::smfmac_f32_32x32x64_fp8_fp8::getOperationName();
if (isFp8(sourceAElem) && isBf8(sourceBElem) && destElem.isF32())
return ROCDL::smfmac_f32_32x32x64_fp8_bf8::getOperationName();
if (isBf8(sourceAElem) && isFp8(sourceBElem) && destElem.isF32())
return ROCDL::smfmac_f32_32x32x64_bf8_fp8::getOperationName();
if (isBf8(sourceAElem) && isBf8(sourceBElem) && destElem.isF32())
return ROCDL::smfmac_f32_32x32x64_bf8_bf8::getOperationName();
}
return std::nullopt;
}
/// Returns the `rocdl` intrinsic corresponding to a WMMA operation `wmma`
/// if one exists. This includes checking to ensure the intrinsic is supported
/// on the architecture you are compiling for.
static std::optional<StringRef> wmmaOpToIntrinsic(WMMAOp wmma,
Chipset chipset) {
auto sourceVectorType = cast<VectorType>(wmma.getSourceA().getType());
auto sourceBVectorType = cast<VectorType>(wmma.getSourceB().getType());
auto destVectorType = cast<VectorType>(wmma.getDestC().getType());
Type elemSourceType = sourceVectorType.getElementType();
Type elemBSourceType = sourceBVectorType.getElementType();
Type elemDestType = destVectorType.getElementType();
const uint32_t k = wmma.getK();
const bool isRDNA3 = chipset.majorVersion == 11;
const bool isRDNA4 = chipset.majorVersion == 12 && chipset.minorVersion == 0;
// Handle RDNA3 and RDNA4.
if (isRDNA3 || isRDNA4)
return wmmaOpToIntrinsicRDNA(elemSourceType, elemBSourceType, elemDestType,
k, isRDNA3);
// Handle gfx1250.
if (chipset == kGfx1250)
return wmmaOpToIntrinsicGfx1250(elemSourceType, elemBSourceType,
elemDestType, k);
return std::nullopt;
}
namespace {
struct MFMAOpLowering : public ConvertOpToLLVMPattern<MFMAOp> {
MFMAOpLowering(const LLVMTypeConverter &converter, Chipset chipset)
: ConvertOpToLLVMPattern<MFMAOp>(converter), chipset(chipset) {}
Chipset chipset;
LogicalResult
matchAndRewrite(MFMAOp op, MFMAOpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
Location loc = op.getLoc();
Type outType = typeConverter->convertType(op.getDestD().getType());
Type intrinsicOutType = outType;
if (auto outVecType = dyn_cast<VectorType>(outType))
if (outVecType.getElementType().isBF16())
intrinsicOutType = outVecType.clone(rewriter.getI16Type());
if (chipset.majorVersion != 9 || chipset < kGfx908)
return op->emitOpError("MFMA only supported on gfx908+");
uint32_t getBlgpField = static_cast<uint32_t>(op.getBlgp());
if (op.getNegateA() || op.getNegateB() || op.getNegateC()) {
if (chipset < kGfx942)
return op.emitOpError("negation unsupported on older than gfx942");
getBlgpField |=
op.getNegateA() | (op.getNegateB() << 1) | (op.getNegateC() << 2);
}
std::optional<StringRef> maybeIntrinsic = mfmaOpToIntrinsic(op, chipset);
std::optional<std::tuple<StringRef, uint32_t, uint32_t>>
maybeScaledIntrinsic = mfmaOpToScaledIntrinsic(op, chipset);
if (!maybeIntrinsic.has_value() && !maybeScaledIntrinsic.has_value())
return op.emitOpError("no intrinsic matching MFMA size on given chipset");
bool isScaled =
!maybeIntrinsic.has_value() && maybeScaledIntrinsic.has_value();
if (isScaled &&
(adaptor.getAbid() > 0 || getBlgpField > 0 || op.getCbsz() > 0)) {
return op.emitOpError(
"non-default abid, blgp, and cbsz aren't supported on MFMAs that can "
"be scaled as those fields are used for type information");
}
StringRef intrinsicName =
isScaled ? std::get<0>(*maybeScaledIntrinsic) : *maybeIntrinsic;
// Determine if we can use bf16 in the intrinsic. Newer MFMAs in gfx950+
// allows bf16 as the input. For reference check IntrinsicsAMDGPU.td file.
bool allowBf16 = [&]() {
if (chipset < kGfx950)
return false;
if (isScaled)
return true;
return intrinsicName.contains("16x16x32.bf16") ||
intrinsicName.contains("32x32x16.bf16");
}();
OperationState loweredOp(loc, intrinsicName);
loweredOp.addTypes(intrinsicOutType);
loweredOp.addOperands({packSmallFloatVectorOperand(
rewriter, loc, adaptor.getSourceA(), allowBf16),
packSmallFloatVectorOperand(
rewriter, loc, adaptor.getSourceB(), allowBf16),
adaptor.getDestC()});
if (isScaled) {
Value zero = createI32Constant(rewriter, loc, 0);
auto [_scaledName, aTypeCode, bTypeCode] = *maybeScaledIntrinsic;
loweredOp.addOperands({createI32Constant(rewriter, loc, aTypeCode),
createI32Constant(rewriter, loc, bTypeCode),
/*scale A byte=*/zero, /*scale A=*/zero,
/*scale B byte=*/zero, /*scale B=*/zero});
} else {
loweredOp.addOperands({createI32Constant(rewriter, loc, op.getCbsz()),
createI32Constant(rewriter, loc, op.getAbid()),
createI32Constant(rewriter, loc, getBlgpField)});
};
Value lowered = rewriter.create(loweredOp)->getResult(0);
if (outType != intrinsicOutType)
lowered = LLVM::BitcastOp::create(rewriter, loc, outType, lowered);
rewriter.replaceOp(op, lowered);
return success();
}
};
struct ScaledMFMAOpLowering : public ConvertOpToLLVMPattern<ScaledMFMAOp> {
ScaledMFMAOpLowering(const LLVMTypeConverter &converter, Chipset chipset)
: ConvertOpToLLVMPattern(converter), chipset(chipset) {}
Chipset chipset;
LogicalResult
matchAndRewrite(ScaledMFMAOp op, ScaledMFMAOpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
Location loc = op.getLoc();
Type intrinsicOutType = typeConverter->convertType(op.getDestD().getType());
if (chipset.majorVersion != 9 || chipset < kGfx950)
return op->emitOpError("scaled MFMA only supported on gfx908+");
std::optional<std::tuple<StringRef, uint32_t, uint32_t>>
maybeScaledIntrinsic = mfmaOpToScaledIntrinsic(op, chipset);
if (!maybeScaledIntrinsic.has_value())
return op.emitOpError(
"no intrinsic matching scaled MFMA size on given chipset");
auto [intrinsicName, aTypeCode, bTypeCode] = *maybeScaledIntrinsic;
OperationState loweredOp(loc, intrinsicName);
loweredOp.addTypes(intrinsicOutType);
loweredOp.addOperands(
{packSmallFloatVectorOperand(rewriter, loc, adaptor.getSourceA()),
packSmallFloatVectorOperand(rewriter, loc, adaptor.getSourceB()),
adaptor.getDestC()});
Value scalesIdxA =
createI32Constant(rewriter, loc, adaptor.getScalesIdxA());
Value scalesIdxB =
createI32Constant(rewriter, loc, adaptor.getScalesIdxB());
loweredOp.addOperands(
{createI32Constant(rewriter, loc, aTypeCode),
createI32Constant(rewriter, loc, bTypeCode),
/*scales idx A=*/scalesIdxA,
/*scales A*/
castScaleOperand(rewriter, loc, adaptor.getScalesA()),
/*scales idx B=*/scalesIdxB,
/*scales B*/
castScaleOperand(rewriter, loc, adaptor.getScalesB())});
Value lowered = rewriter.create(loweredOp)->getResult(0);
rewriter.replaceOp(op, lowered);
return success();
}
};
struct SparseMFMAOpLowering : public ConvertOpToLLVMPattern<SparseMFMAOp> {
SparseMFMAOpLowering(const LLVMTypeConverter &converter, Chipset chipset)
: ConvertOpToLLVMPattern<SparseMFMAOp>(converter), chipset(chipset) {}
Chipset chipset;
LogicalResult
matchAndRewrite(SparseMFMAOp op, SparseMFMAOpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
Location loc = op.getLoc();
auto outType =
typeConverter->convertType<VectorType>(op.getDestC().getType());
if (!outType)
return rewriter.notifyMatchFailure(op, "type conversion failed");
// smfmac is supported on gfx942 and gfx950.
if (chipset.majorVersion != 9 || chipset < kGfx942)
return op->emitOpError("sparse MFMA (smfmac) only supported on gfx942+");
bool isGfx950 = chipset >= kGfx950;
Value a = convertSparseMFMAVectorOperand(rewriter, loc,
adaptor.getSourceA(), isGfx950);
Value b = convertSparseMFMAVectorOperand(rewriter, loc,
adaptor.getSourceB(), isGfx950);
Value c = adaptor.getDestC();
std::optional<StringRef> maybeIntrinsic = smfmacOpToIntrinsic(op, chipset);
if (!maybeIntrinsic.has_value())
return op.emitOpError(
"no intrinsic matching sparse MFMA on the given chipset");
// Bitcast sparse indices from vector<4xi8> or vector<2xi16> to i32.
Value sparseIdx = LLVM::BitcastOp::create(
rewriter, loc, rewriter.getI32Type(), adaptor.getSparseIdx());
OperationState loweredOp(loc, maybeIntrinsic.value());
loweredOp.addTypes(outType);
loweredOp.addOperands({a, b, c, sparseIdx,
createI32Constant(rewriter, loc, op.getCbsz()),
createI32Constant(rewriter, loc, op.getAbid())});
Value lowered = rewriter.create(loweredOp)->getResult(0);
rewriter.replaceOp(op, lowered);
return success();
}
};
struct WMMAOpLowering : public ConvertOpToLLVMPattern<WMMAOp> {
WMMAOpLowering(const LLVMTypeConverter &converter, Chipset chipset)
: ConvertOpToLLVMPattern<WMMAOp>(converter), chipset(chipset) {}
Chipset chipset;
LogicalResult
matchAndRewrite(WMMAOp op, WMMAOpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
Location loc = op.getLoc();
auto outType =
typeConverter->convertType<VectorType>(op.getDestD().getType());
if (!outType)
return rewriter.notifyMatchFailure(op, "type conversion failed");
if (chipset.majorVersion != 11 && chipset.majorVersion != 12)
return op->emitOpError("WMMA only supported on gfx11 and gfx12");
bool isGFX1250 = chipset >= kGfx1250;
// The WMMA operations represent vectors of bf16s as vectors of i16s
// (except on gfx1250), so we need to bitcast bfloats to i16 and then
// bitcast them back.
auto aType = cast<VectorType>(adaptor.getSourceA().getType());
auto bType = cast<VectorType>(adaptor.getSourceB().getType());
auto destCType = cast<VectorType>(adaptor.getDestC().getType());
bool castAToI16 = aType.getElementType().isBF16() && !isGFX1250;
bool castBToI16 = bType.getElementType().isBF16() && !isGFX1250;
bool castDestCToI16 = destCType.getElementType().isBF16() && !isGFX1250;
bool castOutToI16 = outType.getElementType().isBF16() && !isGFX1250;
VectorType rawOutType = outType;
if (castOutToI16)
rawOutType = outType.clone(rewriter.getI16Type());
Value a = adaptor.getSourceA();
if (castAToI16)
a = LLVM::BitcastOp::create(rewriter, loc,
aType.clone(rewriter.getI16Type()), a);
Value b = adaptor.getSourceB();
if (castBToI16)
b = LLVM::BitcastOp::create(rewriter, loc,
bType.clone(rewriter.getI16Type()), b);
Value destC = adaptor.getDestC();
if (castDestCToI16)
destC = LLVM::BitcastOp::create(
rewriter, loc, destCType.clone(rewriter.getI16Type()), destC);
std::optional<StringRef> maybeIntrinsic = wmmaOpToIntrinsic(op, chipset);
if (!maybeIntrinsic.has_value())
return op.emitOpError("no intrinsic matching WMMA on the given chipset");
if (chipset.majorVersion >= 12 && op.getSubwordOffset() != 0)
return op.emitOpError("subwordOffset not supported on gfx12+");
SmallVector<Value, 4> operands;
SmallVector<NamedAttribute, 4> attrs;
wmmaPushInputOperand(rewriter, loc, typeConverter, op.getUnsignedA(), a,
op.getSourceA(), operands, attrs, "signA");
wmmaPushInputOperand(rewriter, loc, typeConverter, op.getUnsignedB(), b,
op.getSourceB(), operands, attrs, "signB");
wmmaPushOutputOperand(rewriter, loc, typeConverter, destC,
op.getSubwordOffset(), op.getClamp(), operands,
attrs);
OperationState loweredOp(loc, *maybeIntrinsic);
loweredOp.addTypes(rawOutType);
loweredOp.addOperands(operands);
loweredOp.addAttributes(attrs);
Operation *lowered = rewriter.create(loweredOp);
Operation *maybeCastBack = lowered;
if (rawOutType != outType)
maybeCastBack = LLVM::BitcastOp::create(rewriter, loc, outType,
lowered->getResult(0));
rewriter.replaceOp(op, maybeCastBack->getResults());
return success();
}
};
struct ScaledWMMAOpLowering : public ConvertOpToLLVMPattern<ScaledWMMAOp> {
ScaledWMMAOpLowering(const LLVMTypeConverter &converter, Chipset chipset)
: ConvertOpToLLVMPattern<ScaledWMMAOp>(converter), chipset(chipset) {}
Chipset chipset;
LogicalResult
matchAndRewrite(ScaledWMMAOp op, ScaledWMMAOpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
Location loc = op.getLoc();
auto outType =
typeConverter->convertType<VectorType>(op.getDestD().getType());
if (!outType)
return rewriter.notifyMatchFailure(op, "type conversion failed");
if (chipset < kGfx1250)
return op->emitOpError("WMMA scale only supported on gfx1250+");
int64_t m = op.getM();
int64_t n = op.getN();
int64_t k = op.getK();
Type aElemType = getElementTypeOrSelf(op.getSourceA().getType());
Type bElemType = getElementTypeOrSelf(op.getSourceB().getType());
std::optional<uint32_t> aFmtCode = smallFloatTypeToFormatCode(aElemType);
std::optional<uint32_t> bFmtCode = smallFloatTypeToFormatCode(bElemType);
if (!aFmtCode || !bFmtCode)
return op.emitOpError("unsupported element types for scaled_wmma");
// Get scale vector types and determine variant (scale vs scale16).
auto scaleAVecType = cast<VectorType>(op.getScaleA().getType());
auto scaleBVecType = cast<VectorType>(op.getScaleB().getType());
if (scaleAVecType.getNumElements() != scaleBVecType.getNumElements())
return op.emitOpError("scaleA and scaleB must have equal vector length");
// Extract scale format from element types.
Type scaleAElemType = scaleAVecType.getElementType();
Type scaleBElemType = scaleBVecType.getElementType();
std::optional<uint32_t> scaleAFmt = getWmmaScaleFormat(scaleAElemType);
std::optional<uint32_t> scaleBFmt = getWmmaScaleFormat(scaleBElemType);
if (!scaleAFmt || !scaleBFmt)
return op.emitOpError("unsupported scale element types");
// Determine which intrinsic to use based on dimensions.
bool isScale16 = (scaleAVecType.getNumElements() == 8);
std::optional<StringRef> intrinsicName =
getScaledWmmaIntrinsicName(m, n, k, isScale16);
if (!intrinsicName)
return op.emitOpError("unsupported scaled_wmma dimensions: ")
<< m << "x" << n << "x" << k;
SmallVector<NamedAttribute, 8> attrs;
// The f4 variant does not have fmtA and fmtB attributes.
bool is32x16 = (m == 32 && n == 16 && k == 128);
if (!is32x16) {
attrs.emplace_back("fmtA", rewriter.getI32IntegerAttr(*aFmtCode));
attrs.emplace_back("fmtB", rewriter.getI32IntegerAttr(*bFmtCode));
}
// modC uses default value of 0.
attrs.emplace_back("modC", rewriter.getI16IntegerAttr(0));
// Scale attributes. Convert user-facing firstScaleLane (0 or 16) to the
// half of the wave that is being selected (0 or 1).
attrs.emplace_back(
"scaleAType", rewriter.getI32IntegerAttr(op.getAFirstScaleLane() / 16));
attrs.emplace_back("fmtScaleA", rewriter.getI32IntegerAttr(*scaleAFmt));
attrs.emplace_back(
"scaleBType", rewriter.getI32IntegerAttr(op.getBFirstScaleLane() / 16));
attrs.emplace_back("fmtScaleB", rewriter.getI32IntegerAttr(*scaleBFmt));
// Reuse flags use default value of false.
attrs.emplace_back("reuseA", rewriter.getBoolAttr(false));
attrs.emplace_back("reuseB", rewriter.getBoolAttr(false));
// Convert typed float vectors to packed format.
Value sourceA =
packSmallFloatVectorOperand(rewriter, loc, adaptor.getSourceA());
Value sourceB =
packSmallFloatVectorOperand(rewriter, loc, adaptor.getSourceB());
// Pack scale vectors into i32/i64.
Value packedScaleA = castScaleOperand(rewriter, loc, adaptor.getScaleA());
Value packedScaleB = castScaleOperand(rewriter, loc, adaptor.getScaleB());
// Create the intrinsic call.
OperationState loweredOp(loc, *intrinsicName);
loweredOp.addTypes(outType);
loweredOp.addOperands(
{sourceA, sourceB, adaptor.getDestC(), packedScaleA, packedScaleB});
loweredOp.addAttributes(attrs);
Operation *lowered = rewriter.create(loweredOp);
rewriter.replaceOp(op, lowered->getResults());
return success();
}
};
struct TransposeLoadOpLowering
: public ConvertOpToLLVMPattern<TransposeLoadOp> {
TransposeLoadOpLowering(const LLVMTypeConverter &converter, Chipset chipset)
: ConvertOpToLLVMPattern<TransposeLoadOp>(converter), chipset(chipset) {}
Chipset chipset;
LogicalResult
matchAndRewrite(TransposeLoadOp op, TransposeLoadOpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
if (chipset != kGfx950)
return op.emitOpError("Non-gfx950 chipset not supported");
Location loc = op.getLoc();
auto srcMemRefType = cast<MemRefType>(op.getSrc().getType());
// Elements in subbyte memrefs are stored non-contiguously,
// reject if source is sub-byte memref. Use emulated memrefs instead.
size_t srcElementSize =
srcMemRefType.getElementType().getIntOrFloatBitWidth();
if (srcElementSize < 8)
return op.emitOpError("Expect source memref to have at least 8 bits "
"element size, got ")
<< srcElementSize;
auto resultType = cast<VectorType>(op.getResult().getType());
Value srcPtr =
getStridedElementPtr(rewriter, loc, srcMemRefType, adaptor.getSrc(),
(adaptor.getSrcIndices()));
size_t numElements = resultType.getNumElements();
size_t elementTypeSize =
resultType.getElementType().getIntOrFloatBitWidth();
// ROCDL transpose load intrinsics return vectors of 32-bit integers, if
// the element size is smaller than 16 bits.
Type rocdlResultType = VectorType::get((numElements * elementTypeSize) / 32,
rewriter.getIntegerType(32));
Type llvmResultType = typeConverter->convertType(resultType);
switch (elementTypeSize) {
case 4: {
assert(numElements == 16);
auto rocdlOp = ROCDL::ds_read_tr4_b64::create(rewriter, loc,
rocdlResultType, srcPtr);
rewriter.replaceOpWithNewOp<LLVM::BitcastOp>(op, llvmResultType, rocdlOp);
break;
}
case 6: {
assert(numElements == 16);
auto rocdlOp = ROCDL::ds_read_tr6_b96::create(rewriter, loc,
rocdlResultType, srcPtr);
rewriter.replaceOpWithNewOp<LLVM::BitcastOp>(op, llvmResultType, rocdlOp);
break;
}
case 8: {
assert(numElements == 8);
auto rocdlOp = ROCDL::ds_read_tr8_b64::create(rewriter, loc,
rocdlResultType, srcPtr);
rewriter.replaceOpWithNewOp<LLVM::BitcastOp>(op, llvmResultType, rocdlOp);
break;
}
case 16: {
assert(numElements == 4);
rewriter.replaceOpWithNewOp<ROCDL::ds_read_tr16_b64>(op, llvmResultType,
srcPtr);
break;
}
default:
return op.emitOpError("Unsupported element size for transpose load");
}
return success();
}
};
struct GatherToLDSOpLowering : public ConvertOpToLLVMPattern<GatherToLDSOp> {
GatherToLDSOpLowering(const LLVMTypeConverter &converter, Chipset chipset)
: ConvertOpToLLVMPattern<GatherToLDSOp>(converter), chipset(chipset) {}
Chipset chipset;
LogicalResult
matchAndRewrite(GatherToLDSOp op, GatherToLDSOpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
if (chipset.majorVersion < 9 || chipset.majorVersion > 10)
return op.emitOpError("pre-gfx9 and post-gfx10 not supported");
Location loc = op.getLoc();
auto srcMemRefType = cast<MemRefType>(op.getSrc().getType());
auto dstMemRefType = cast<MemRefType>(op.getDst().getType());
// TODO: instead of only transfering one element per thread, we could
// augment it to transfer multiple elements per thread by issuing multiple
// `global_load_lds` instructions.
Type transferType = op.getTransferType();
int loadWidth = [&]() -> int {
if (auto transferVectorType = dyn_cast<VectorType>(transferType)) {
return (transferVectorType.getNumElements() *
transferVectorType.getElementTypeBitWidth()) /
8;
}
return transferType.getIntOrFloatBitWidth() / 8;
}();
// Currently only 1, 2, 4, 12 and 16 byte loads are supported.
if (!llvm::is_contained({1, 2, 4, 12, 16}, loadWidth))
return op.emitOpError("chipset unsupported element size");
if (chipset != kGfx950 && llvm::is_contained({12, 16}, loadWidth))
return op.emitOpError("Gather to LDS instructions with 12-byte and "
"16-byte load widths are only supported on gfx950");
Value srcPtr =
getStridedElementPtr(rewriter, loc, srcMemRefType, adaptor.getSrc(),
(adaptor.getSrcIndices()));
Value dstPtr =
getStridedElementPtr(rewriter, loc, dstMemRefType, adaptor.getDst(),
(adaptor.getDstIndices()));
rewriter.replaceOpWithNewOp<ROCDL::LoadToLDSOp>(
op, srcPtr, dstPtr, rewriter.getI32IntegerAttr(loadWidth),
/*offset=*/rewriter.getI32IntegerAttr(0),
/*aux=*/rewriter.getI32IntegerAttr(0), ArrayAttr{}, ArrayAttr{},
ArrayAttr{});
return success();
}
};
namespace {
struct ExtPackedFp8OpLowering final
: public ConvertOpToLLVMPattern<ExtPackedFp8Op> {
ExtPackedFp8OpLowering(const LLVMTypeConverter &converter, Chipset chipset)
: ConvertOpToLLVMPattern<amdgpu::ExtPackedFp8Op>(converter),
chipset(chipset) {}
Chipset chipset;
LogicalResult
matchAndRewrite(ExtPackedFp8Op op, ExtPackedFp8OpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override;
};
struct ScaledExtPackedMatrixOpLowering final
: public ConvertOpToLLVMPattern<ScaledExtPackedMatrixOp> {
ScaledExtPackedMatrixOpLowering(const LLVMTypeConverter &converter,
Chipset chipset)
: ConvertOpToLLVMPattern<amdgpu::ScaledExtPackedMatrixOp>(converter),
chipset(chipset) {}
Chipset chipset;
LogicalResult
matchAndRewrite(ScaledExtPackedMatrixOp op,
ScaledExtPackedMatrixOpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override;
};
struct PackedTrunc2xFp8OpLowering final
: public ConvertOpToLLVMPattern<PackedTrunc2xFp8Op> {
PackedTrunc2xFp8OpLowering(const LLVMTypeConverter &converter,
Chipset chipset)
: ConvertOpToLLVMPattern<amdgpu::PackedTrunc2xFp8Op>(converter),
chipset(chipset) {}
Chipset chipset;
LogicalResult
matchAndRewrite(PackedTrunc2xFp8Op op, PackedTrunc2xFp8OpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override;
};
struct PackedStochRoundFp8OpLowering final
: public ConvertOpToLLVMPattern<PackedStochRoundFp8Op> {
PackedStochRoundFp8OpLowering(const LLVMTypeConverter &converter,
Chipset chipset)
: ConvertOpToLLVMPattern<amdgpu::PackedStochRoundFp8Op>(converter),
chipset(chipset) {}
Chipset chipset;
LogicalResult
matchAndRewrite(PackedStochRoundFp8Op op,
PackedStochRoundFp8OpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override;
};
struct ScaledExtPackedOpLowering final
: public ConvertOpToLLVMPattern<ScaledExtPackedOp> {
ScaledExtPackedOpLowering(const LLVMTypeConverter &converter, Chipset chipset)
: ConvertOpToLLVMPattern<amdgpu::ScaledExtPackedOp>(converter),
chipset(chipset) {}
Chipset chipset;
LogicalResult
matchAndRewrite(ScaledExtPackedOp op, ScaledExtPackedOpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override;
};
struct PackedScaledTruncOpLowering final
: public ConvertOpToLLVMPattern<PackedScaledTruncOp> {
PackedScaledTruncOpLowering(const LLVMTypeConverter &converter,
Chipset chipset)
: ConvertOpToLLVMPattern<amdgpu::PackedScaledTruncOp>(converter),
chipset(chipset) {}
Chipset chipset;
LogicalResult
matchAndRewrite(PackedScaledTruncOp op, PackedScaledTruncOpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override;
};
} // end namespace
LogicalResult ExtPackedFp8OpLowering::matchAndRewrite(
ExtPackedFp8Op op, ExtPackedFp8OpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const {
Location loc = op.getLoc();
if (!(chipset == kGfx942 || hasOcpFp8(chipset)))
return rewriter.notifyMatchFailure(
loc, "Fp8 conversion instructions are not available on target "
"architecture and their emulation is not implemented");
Type v4i8 =
getTypeConverter()->convertType(VectorType::get(4, rewriter.getI8Type()));
Type i32 = getTypeConverter()->convertType(rewriter.getI32Type());
Type f32 = getTypeConverter()->convertType(op.getResult().getType());
Value source = adaptor.getSource();
auto sourceVecType = dyn_cast<VectorType>(op.getSource().getType());
auto resultVecType = dyn_cast<VectorType>(op.getResult().getType());
Type sourceElemType = getElementTypeOrSelf(op.getSource());
// Extend to a v4i8
if (!sourceVecType || sourceVecType.getNumElements() < 4) {
Value longVec = LLVM::UndefOp::create(rewriter, loc, v4i8);
if (!sourceVecType) {
longVec = LLVM::InsertElementOp::create(
rewriter, loc, longVec, source, createI32Constant(rewriter, loc, 0));
} else {
for (int32_t i = 0, e = sourceVecType.getNumElements(); i < e; ++i) {
Value idx = createI32Constant(rewriter, loc, i);
Value elem = LLVM::ExtractElementOp::create(rewriter, loc, source, idx);
longVec =
LLVM::InsertElementOp::create(rewriter, loc, longVec, elem, idx);
}
}
source = longVec;
}
Value i32Source = LLVM::BitcastOp::create(rewriter, loc, i32, source);
if (resultVecType) {
if (typeIsExpectedBf8ForChipset(chipset, sourceElemType)) {
rewriter.replaceOpWithNewOp<ROCDL::CvtPkF32Bf8Op>(op, f32, i32Source,
op.getIndex());
} else if (typeIsExpectedFp8ForChipset(chipset, sourceElemType)) {
rewriter.replaceOpWithNewOp<ROCDL::CvtPkF32Fp8Op>(op, f32, i32Source,
op.getIndex());
}
} else {
if (typeIsExpectedBf8ForChipset(chipset, sourceElemType)) {
rewriter.replaceOpWithNewOp<ROCDL::CvtF32Bf8Op>(op, f32, i32Source,
op.getIndex());
} else if (typeIsExpectedFp8ForChipset(chipset, sourceElemType)) {
rewriter.replaceOpWithNewOp<ROCDL::CvtF32Fp8Op>(op, f32, i32Source,
op.getIndex());
}
}
return success();
}
int32_t getScaleSel(int32_t blockSize, unsigned bitWidth, int32_t scaleWaveHalf,
int32_t firstScaleByte) {
// When lowering amdgpu.scaled_ext_packed_matrix to rocdl.cvt.scale.pk*.f*.f*
// operations, the attributes blockSize, sourceType, scaleWaveHalf, and
// firstScaleByte are merged into a single attribute scaleSel. This is how
// those values are merged together. (Note: scaleWaveHalf isn't a high-level
// attribute but is derifed from firstScaleLane).
assert(llvm::is_contained({16, 32}, blockSize));
assert(llvm::is_contained({4u, 6u, 8u}, bitWidth));
const bool isFp8 = bitWidth == 8;
const bool isBlock16 = blockSize == 16;
if (!isFp8) {
int32_t bit0 = isBlock16;
assert(llvm::is_contained({0, 1, 2}, firstScaleByte));
int32_t bit1 = (firstScaleByte == 2) << 1;
assert(llvm::is_contained({0, 1}, scaleWaveHalf));
int32_t bit2 = scaleWaveHalf << 2;
return bit2 | bit1 | bit0;
}
int32_t bit0 = isBlock16;
// firstScaleByte is guaranteed to be defined by two bits.
assert(llvm::is_contained({0, 1, 2, 3}, firstScaleByte));
int32_t bits2and1 = firstScaleByte << 1;
assert(llvm::is_contained({0, 1}, scaleWaveHalf));
int32_t bit3 = scaleWaveHalf << 3;
int32_t bits = bit3 | bits2and1 | bit0;
// These are invalid cases.
assert(!llvm::is_contained(
{0b0011, 0b0101, 0b0111, 0b1000, 0b1001, 0b1011, 0b1111}, bits));
return bits;
}
static std::optional<StringRef>
scaledExtPacked816ToIntrinsic(Type srcElemType, Type destElemType) {
using fp4 = Float4E2M1FNType;
using fp8 = Float8E4M3FNType;
using bf8 = Float8E5M2Type;
using fp6 = Float6E2M3FNType;
using bf6 = Float6E3M2FNType;
if (isa<fp4>(srcElemType)) {
if (destElemType.isF16())
return ROCDL::CvtPkScalePk8F16Fp4Op::getOperationName();
if (destElemType.isBF16())
return ROCDL::CvtPkScalePk8Bf16Fp4Op::getOperationName();
if (destElemType.isF32())
return ROCDL::CvtPkScalePk8F32Fp4Op::getOperationName();
return std::nullopt;
}
if (isa<fp8>(srcElemType)) {
if (destElemType.isF16())
return ROCDL::CvtPkScalePk8F16Fp8Op::getOperationName();
if (destElemType.isBF16())
return ROCDL::CvtPkScalePk8Bf16Fp8Op::getOperationName();
if (destElemType.isF32())
return ROCDL::CvtPkScalePk8F32Fp8Op::getOperationName();
return std::nullopt;
}
if (isa<bf8>(srcElemType)) {
if (destElemType.isF16())
return ROCDL::CvtPkScalePk8F16Bf8Op::getOperationName();
if (destElemType.isBF16())
return ROCDL::CvtPkScalePk8Bf16Bf8Op::getOperationName();
if (destElemType.isF32())
return ROCDL::CvtPkScalePk8F32Bf8Op::getOperationName();
return std::nullopt;
}
if (isa<fp6>(srcElemType)) {
if (destElemType.isF16())
return ROCDL::CvtPkScalePk16F16Fp6Op::getOperationName();
if (destElemType.isBF16())
return ROCDL::CvtPkScalePk16Bf16Fp6Op::getOperationName();
if (destElemType.isF32())
return ROCDL::CvtPkScalePk16F32Fp6Op::getOperationName();
return std::nullopt;
}
if (isa<bf6>(srcElemType)) {
if (destElemType.isF16())
return ROCDL::CvtPkScalePk16F16Bf6Op::getOperationName();
if (destElemType.isBF16())
return ROCDL::CvtPkScalePk16Bf16Bf6Op::getOperationName();
if (destElemType.isF32())
return ROCDL::CvtPkScalePk16F32Bf6Op::getOperationName();
return std::nullopt;
}
llvm_unreachable("invalid combination of element types for packed conversion "
"instructions");
}
LogicalResult ScaledExtPackedMatrixOpLowering::matchAndRewrite(
ScaledExtPackedMatrixOp op, ScaledExtPackedMatrixOpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const {
using fp4 = Float4E2M1FNType;
using fp8 = Float8E4M3FNType;
using bf8 = Float8E5M2Type;
using fp6 = Float6E2M3FNType;
using bf6 = Float6E3M2FNType;
Location loc = op.getLoc();
if (chipset != kGfx1250) {
return rewriter.notifyMatchFailure(
loc,
"Scaled fp packed conversion instructions are not available on target "
"architecture and their emulation is not implemented");
}
// Convert user-facing firstScaleLane (0 or 16) to the half of the wave that
// is being selected.
int32_t scaleWaveHalf = op.getFirstScaleLane() / 16;
int32_t firstScaleByte = op.getFirstScaleByte();
int32_t blockSize = op.getBlockSize();
auto sourceType = cast<VectorType>(op.getSource().getType());
auto srcElemType = cast<FloatType>(sourceType.getElementType());
unsigned bitWidth = srcElemType.getWidth();
auto targetType = cast<VectorType>(op.getResult().getType());
auto destElemType = cast<FloatType>(targetType.getElementType());
IntegerType i32 = rewriter.getI32Type();
Value source = adaptor.getSource();
Type llvmResultType = typeConverter->convertType(op.getResult().getType());
Type packedType = nullptr;
if (isa<fp4>(srcElemType)) {
packedType = i32;
packedType = getTypeConverter()->convertType(packedType);
} else if (isa<fp8, bf8>(srcElemType)) {
packedType = VectorType::get(2, i32);
packedType = getTypeConverter()->convertType(packedType);
} else if (isa<fp6, bf6>(srcElemType)) {
packedType = VectorType::get(3, i32);
packedType = getTypeConverter()->convertType(packedType);
} else {
llvm_unreachable("invalid element type for packed scaled ext");
}
if (!packedType || !llvmResultType) {
return rewriter.notifyMatchFailure(op, "type conversion failed");
}
std::optional<StringRef> maybeIntrinsic =
scaledExtPacked816ToIntrinsic(srcElemType, destElemType);
if (!maybeIntrinsic.has_value())
return op.emitOpError(
"no intrinsic matching packed scaled conversion on the given chipset");
int32_t scaleSel =
getScaleSel(blockSize, bitWidth, scaleWaveHalf, firstScaleByte);
Value castedScale =
LLVM::BitcastOp::create(rewriter, loc, i32, adaptor.getScale());
Value castedSource =
LLVM::BitcastOp::create(rewriter, loc, packedType, source);
OperationState loweredOp(loc, *maybeIntrinsic);
loweredOp.addTypes({llvmResultType});
loweredOp.addOperands({castedSource, castedScale});
SmallVector<NamedAttribute, 1> attrs;
attrs.push_back(
NamedAttribute("scaleSel", rewriter.getI32IntegerAttr(scaleSel)));
loweredOp.addAttributes(attrs);
Operation *lowered = rewriter.create(loweredOp);
rewriter.replaceOp(op, lowered);
return success();
}
LogicalResult ScaledExtPackedOpLowering::matchAndRewrite(
ScaledExtPackedOp op, ScaledExtPackedOpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const {
Location loc = op.getLoc();
if (chipset != kGfx950)
return rewriter.notifyMatchFailure(
loc, "Scaled fp conversion instructions are not available on target "
"architecture and their emulation is not implemented");
Type i32 = getTypeConverter()->convertType(rewriter.getI32Type());
Value source = adaptor.getSource();
Value scale = adaptor.getScale();
VectorType sourceVecType = cast<VectorType>(op.getSource().getType());
Type sourceElemType = sourceVecType.getElementType();
VectorType destVecType = cast<VectorType>(op.getResult().getType());
Type destElemType = destVecType.getElementType();
VectorType packedVecType;
if (isa<Float8E5M2Type, Float8E4M3FNType>(sourceElemType)) {
VectorType v4i8 = VectorType::get(4, rewriter.getI8Type());
packedVecType = cast<VectorType>(getTypeConverter()->convertType(v4i8));
} else if (isa<Float4E2M1FNType>(sourceElemType)) {
VectorType v8i4 = VectorType::get(8, rewriter.getI4Type());
packedVecType = cast<VectorType>(getTypeConverter()->convertType(v8i4));
} else {
llvm_unreachable("invalid element type for scaled ext");
}
// Extend to a packedVectorType
if (sourceVecType.getNumElements() < packedVecType.getNumElements()) {
Value longVec = LLVM::ZeroOp::create(rewriter, loc, packedVecType);
if (!sourceVecType) {
longVec = LLVM::InsertElementOp::create(
rewriter, loc, longVec, source, createI32Constant(rewriter, loc, 0));
} else {
for (int32_t i = 0, e = sourceVecType.getNumElements(); i < e; ++i) {
Value idx = createI32Constant(rewriter, loc, i);
Value elem = LLVM::ExtractElementOp::create(rewriter, loc, source, idx);
longVec =
LLVM::InsertElementOp::create(rewriter, loc, longVec, elem, idx);
}
}
source = longVec;
}
Value i32Source = LLVM::BitcastOp::create(rewriter, loc, i32, source);
if (isa<Float8E5M2Type>(sourceElemType) && destElemType.isF32())
rewriter.replaceOpWithNewOp<ROCDL::CvtScaleF32PkF32Bf8Op>(
op, destVecType, i32Source, scale, op.getIndex());
else if (isa<Float8E5M2Type>(sourceElemType) && destElemType.isF16())
rewriter.replaceOpWithNewOp<ROCDL::CvtScaleF32PkF16Bf8Op>(
op, destVecType, i32Source, scale, op.getIndex());
else if (isa<Float8E5M2Type>(sourceElemType) && destElemType.isBF16())
rewriter.replaceOpWithNewOp<ROCDL::CvtScaleF32PkBf16Bf8Op>(
op, destVecType, i32Source, scale, op.getIndex());
else if (isa<Float8E4M3FNType>(sourceElemType) && destElemType.isF32())
rewriter.replaceOpWithNewOp<ROCDL::CvtScaleF32PkF32Fp8Op>(
op, destVecType, i32Source, scale, op.getIndex());
else if (isa<Float8E4M3FNType>(sourceElemType) && destElemType.isF16())
rewriter.replaceOpWithNewOp<ROCDL::CvtScaleF32PkF16Fp8Op>(
op, destVecType, i32Source, scale, op.getIndex());
else if (isa<Float8E4M3FNType>(sourceElemType) && destElemType.isBF16())
rewriter.replaceOpWithNewOp<ROCDL::CvtScaleF32PkBf16Fp8Op>(
op, destVecType, i32Source, scale, op.getIndex());
else if (isa<Float4E2M1FNType>(sourceElemType) && destElemType.isF32())
rewriter.replaceOpWithNewOp<ROCDL::CvtScaleF32PkF32Fp4Op>(
op, destVecType, i32Source, scale, op.getIndex());
else if (isa<Float4E2M1FNType>(sourceElemType) && destElemType.isF16())
rewriter.replaceOpWithNewOp<ROCDL::CvtScaleF32PkF16Fp4Op>(
op, destVecType, i32Source, scale, op.getIndex());
else if (isa<Float4E2M1FNType>(sourceElemType) && destElemType.isBF16())
rewriter.replaceOpWithNewOp<ROCDL::CvtScaleF32PkBf16Fp4Op>(
op, destVecType, i32Source, scale, op.getIndex());
else
return failure();
return success();
}
LogicalResult PackedScaledTruncOpLowering::matchAndRewrite(
PackedScaledTruncOp op, PackedScaledTruncOpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const {
Location loc = op.getLoc();
if (chipset != kGfx950)
return rewriter.notifyMatchFailure(
loc, "Scaled fp conversion instructions are not available on target "
"architecture and their emulation is not implemented");
Type v2i16 = getTypeConverter()->convertType(
VectorType::get(2, rewriter.getI16Type()));
Type i32 = getTypeConverter()->convertType(rewriter.getI32Type());
Type resultType = op.getResult().getType();
Type resultElemType = getElementTypeOrSelf(resultType);
VectorType sourceVecType = cast<VectorType>(op.getSource().getType());
Type sourceElemType = sourceVecType.getElementType();
Type intResultType = isa<Float4E2M1FNType>(resultElemType) ? i32 : v2i16;
Value source = adaptor.getSource();
Value scale = adaptor.getScale();
Value existing = adaptor.getExisting();
if (existing)
existing = LLVM::BitcastOp::create(rewriter, loc, intResultType, existing);
else
existing = LLVM::ZeroOp::create(rewriter, loc, intResultType);
if (sourceVecType.getNumElements() < 2) {
Value c0 = createI32Constant(rewriter, loc, 0);
Value elem0 = LLVM::ExtractElementOp::create(rewriter, loc, source, c0);
VectorType v2 = VectorType::get(2, sourceElemType);
source = LLVM::ZeroOp::create(rewriter, loc, v2);
source = LLVM::InsertElementOp::create(rewriter, loc, source, elem0, c0);
}
Value sourceA, sourceB;
if (sourceElemType.isF32()) {
Value c0 = createI32Constant(rewriter, loc, 0);
Value c1 = createI32Constant(rewriter, loc, 1);
sourceA = LLVM::ExtractElementOp::create(rewriter, loc, source, c0);
sourceB = LLVM::ExtractElementOp::create(rewriter, loc, source, c1);
}
Value result;
if (sourceElemType.isF32() && isa<Float8E5M2Type>(resultElemType))
result = ROCDL::CvtScaleF32PkBf8F32Op::create(rewriter, loc, intResultType,
existing, sourceA, sourceB,
scale, op.getIndex());
else if (sourceElemType.isF16() && isa<Float8E5M2Type>(resultElemType))
result = ROCDL::CvtScaleF32PkBf8F16Op::create(
rewriter, loc, intResultType, existing, source, scale, op.getIndex());
else if (sourceElemType.isBF16() && isa<Float8E5M2Type>(resultElemType))
result = ROCDL::CvtScaleF32PkBf8Bf16Op::create(
rewriter, loc, intResultType, existing, source, scale, op.getIndex());
else if (sourceElemType.isF32() && isa<Float8E4M3FNType>(resultElemType))
result = ROCDL::CvtScaleF32PkFp8F32Op::create(rewriter, loc, intResultType,
existing, sourceA, sourceB,
scale, op.getIndex());
else if (sourceElemType.isF16() && isa<Float8E4M3FNType>(resultElemType))
result = ROCDL::CvtScaleF32PkFp8F16Op::create(
rewriter, loc, intResultType, existing, source, scale, op.getIndex());
else if (sourceElemType.isBF16() && isa<Float8E4M3FNType>(resultElemType))
result = ROCDL::CvtScaleF32PkFp8Bf16Op::create(
rewriter, loc, intResultType, existing, source, scale, op.getIndex());
else if (sourceElemType.isF32() && isa<Float4E2M1FNType>(resultElemType))
result = ROCDL::CvtScaleF32PkFp4F32Op::create(rewriter, loc, intResultType,
existing, sourceA, sourceB,
scale, op.getIndex());
else if (sourceElemType.isF16() && isa<Float4E2M1FNType>(resultElemType))
result = ROCDL::CvtScaleF32PkFp4F16Op::create(
rewriter, loc, intResultType, existing, source, scale, op.getIndex());
else if (sourceElemType.isBF16() && isa<Float4E2M1FNType>(resultElemType))
result = ROCDL::CvtScaleF32PkFp4Bf16Op::create(
rewriter, loc, intResultType, existing, source, scale, op.getIndex());
else
return failure();
result = rewriter.replaceOpWithNewOp<LLVM::BitcastOp>(
op, getTypeConverter()->convertType(resultType), result);
return success();
}
LogicalResult PackedTrunc2xFp8OpLowering::matchAndRewrite(
PackedTrunc2xFp8Op op, PackedTrunc2xFp8OpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const {
Location loc = op.getLoc();
if (!(chipset == kGfx942 || hasOcpFp8(chipset)))
return rewriter.notifyMatchFailure(
loc, "Fp8 conversion instructions are not available on target "
"architecture and their emulation is not implemented");
Type i32 = getTypeConverter()->convertType(rewriter.getI32Type());
Type resultType = op.getResult().getType();
Type resultElemType = getElementTypeOrSelf(resultType);
Value sourceA = adaptor.getSourceA();
Value sourceB = adaptor.getSourceB();
if (!sourceB)
sourceB = LLVM::UndefOp::create(rewriter, loc, sourceA.getType());
Value existing = adaptor.getExisting();
if (existing)
existing = LLVM::BitcastOp::create(rewriter, loc, i32, existing);
else
existing = LLVM::UndefOp::create(rewriter, loc, i32);
Value result;
if (typeIsExpectedBf8ForChipset(chipset, resultElemType))
result = ROCDL::CvtPkBf8F32Op::create(rewriter, loc, i32, sourceA, sourceB,
existing, op.getWordIndex());
else if (typeIsExpectedFp8ForChipset(chipset, resultElemType))
result = ROCDL::CvtPkFp8F32Op::create(rewriter, loc, i32, sourceA, sourceB,
existing, op.getWordIndex());
result = rewriter.replaceOpWithNewOp<LLVM::BitcastOp>(
op, getTypeConverter()->convertType(resultType), result);
return success();
}
LogicalResult PackedStochRoundFp8OpLowering::matchAndRewrite(
PackedStochRoundFp8Op op, PackedStochRoundFp8OpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const {
Location loc = op.getLoc();
if (!(chipset == kGfx942 || hasOcpFp8(chipset)))
return rewriter.notifyMatchFailure(
loc, "Fp8 conversion instructions are not available on target "
"architecture and their emulation is not implemented");
Type i32 = getTypeConverter()->convertType(rewriter.getI32Type());
Type resultType = op.getResult().getType();
Type resultElemType = getElementTypeOrSelf(resultType);
Value source = adaptor.getSource();
Value stoch = adaptor.getStochiasticParam();
Value existing = adaptor.getExisting();
if (existing)
existing = LLVM::BitcastOp::create(rewriter, loc, i32, existing);
else
existing = LLVM::UndefOp::create(rewriter, loc, i32);
Value result;
if (typeIsExpectedBf8ForChipset(chipset, resultElemType))
result = ROCDL::CvtSrBf8F32Op::create(rewriter, loc, i32, source, stoch,
existing, op.getStoreIndex());
else if (typeIsExpectedFp8ForChipset(chipset, resultElemType))
result = ROCDL::CvtSrFp8F32Op::create(rewriter, loc, i32, source, stoch,
existing, op.getStoreIndex());
result = rewriter.replaceOpWithNewOp<LLVM::BitcastOp>(
op, getTypeConverter()->convertType(resultType), result);
return success();
}
// Implement the AMDGPU_DPPLowering class that will convert the amdgpu.dpp
// operation into the corresponding ROCDL instructions.
struct AMDGPUDPPLowering : public ConvertOpToLLVMPattern<DPPOp> {
AMDGPUDPPLowering(const LLVMTypeConverter &converter, Chipset chipset)
: ConvertOpToLLVMPattern<DPPOp>(converter), chipset(chipset) {}
Chipset chipset;
LogicalResult
matchAndRewrite(DPPOp DppOp, DPPOp::Adaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
// Convert the source operand to the corresponding LLVM type
Location loc = DppOp.getLoc();
Value src = adaptor.getSrc();
Value old = adaptor.getOld();
Type srcType = src.getType();
Type oldType = old.getType();
Type llvmType = nullptr;
if (srcType.getIntOrFloatBitWidth() < 32) {
llvmType = rewriter.getI32Type();
} else if (isa<FloatType>(srcType)) {
llvmType = (srcType.getIntOrFloatBitWidth() == 32)
? rewriter.getF32Type()
: rewriter.getF64Type();
} else if (isa<IntegerType>(srcType)) {
llvmType = (srcType.getIntOrFloatBitWidth() == 32)
? rewriter.getI32Type()
: rewriter.getI64Type();
}
auto llvmSrcIntType = typeConverter->convertType(
rewriter.getIntegerType(srcType.getIntOrFloatBitWidth()));
// If the source type is less of 32, use bitcast to convert it to i32.
auto convertOperand = [&](Value operand, Type operandType) {
if (operandType.getIntOrFloatBitWidth() <= 16) {
if (llvm::isa<FloatType>(operandType)) {
operand =
LLVM::BitcastOp::create(rewriter, loc, llvmSrcIntType, operand);
}
auto llvmVecType = typeConverter->convertType(mlir::VectorType::get(
32 / operandType.getIntOrFloatBitWidth(), llvmSrcIntType));
Value undefVec = LLVM::UndefOp::create(rewriter, loc, llvmVecType);
operand =
LLVM::InsertElementOp::create(rewriter, loc, undefVec, operand,
createI32Constant(rewriter, loc, 0));
operand = LLVM::BitcastOp::create(rewriter, loc, llvmType, operand);
}
return operand;
};
src = convertOperand(src, srcType);
old = convertOperand(old, oldType);
// This is taken from the following file llvm/lib/Target/AMDGPU/SIDefines.h
enum DppCtrl : unsigned {
ROW_SHL0 = 0x100,
ROW_SHR0 = 0x110,
ROW_ROR0 = 0x120,
WAVE_SHL1 = 0x130,
WAVE_ROL1 = 0x134,
WAVE_SHR1 = 0x138,
WAVE_ROR1 = 0x13C,
ROW_MIRROR = 0x140,
ROW_HALF_MIRROR = 0x141,
BCAST15 = 0x142,
BCAST31 = 0x143,
};
auto kind = DppOp.getKind();
auto permArgument = DppOp.getPermArgument();
uint32_t DppCtrl = 0;
switch (kind) {
case DPPPerm::quad_perm: {
auto quadPermAttr = cast<ArrayAttr>(*permArgument);
int32_t i = 0;
for (auto elem : quadPermAttr.getAsRange<IntegerAttr>()) {
uint32_t num = elem.getInt();
DppCtrl |= num << (i * 2);
i++;
}
break;
}
case DPPPerm::row_shl: {
auto intAttr = cast<IntegerAttr>(*permArgument);
DppCtrl = intAttr.getInt() + DppCtrl::ROW_SHL0;
break;
}
case DPPPerm::row_shr: {
auto intAttr = cast<IntegerAttr>(*permArgument);
DppCtrl = intAttr.getInt() + DppCtrl::ROW_SHR0;
break;
}
case DPPPerm::row_ror: {
auto intAttr = cast<IntegerAttr>(*permArgument);
DppCtrl = intAttr.getInt() + DppCtrl::ROW_ROR0;
break;
}
case DPPPerm::wave_shl:
DppCtrl = DppCtrl::WAVE_SHL1;
break;
case DPPPerm::wave_shr:
DppCtrl = DppCtrl::WAVE_SHR1;
break;
case DPPPerm::wave_rol:
DppCtrl = DppCtrl::WAVE_ROL1;
break;
case DPPPerm::wave_ror:
DppCtrl = DppCtrl::WAVE_ROR1;
break;
case DPPPerm::row_mirror:
DppCtrl = DppCtrl::ROW_MIRROR;
break;
case DPPPerm::row_half_mirror:
DppCtrl = DppCtrl::ROW_HALF_MIRROR;
break;
case DPPPerm::row_bcast_15:
DppCtrl = DppCtrl::BCAST15;
break;
case DPPPerm::row_bcast_31:
DppCtrl = DppCtrl::BCAST31;
break;
}
// Check for row_mask, bank_mask, bound_ctrl if they exist and create
// constants
auto rowMask = DppOp->getAttrOfType<IntegerAttr>("row_mask").getInt();
auto bankMask = DppOp->getAttrOfType<IntegerAttr>("bank_mask").getInt();
bool boundCtrl = DppOp->getAttrOfType<BoolAttr>("bound_ctrl").getValue();
// create a ROCDL_DPPMovOp instruction with the appropriate attributes
auto dppMovOp =
ROCDL::DPPUpdateOp::create(rewriter, loc, llvmType, old, src, DppCtrl,
rowMask, bankMask, boundCtrl);
Value result = dppMovOp.getRes();
if (srcType.getIntOrFloatBitWidth() < 32) {
result = LLVM::TruncOp::create(rewriter, loc, llvmSrcIntType, result);
if (!llvm::isa<IntegerType>(srcType)) {
result = LLVM::BitcastOp::create(rewriter, loc, srcType, result);
}
}
// We are replacing the AMDGPU_DPPOp instruction with the new
// ROCDL_DPPMovOp instruction
rewriter.replaceOp(DppOp, ValueRange(result));
return success();
}
};
struct AMDGPUSwizzleBitModeLowering
: public ConvertOpToLLVMPattern<SwizzleBitModeOp> {
using ConvertOpToLLVMPattern::ConvertOpToLLVMPattern;
LogicalResult
matchAndRewrite(SwizzleBitModeOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
Location loc = op.getLoc();
Type i32 = rewriter.getI32Type();
Value src = adaptor.getSrc();
SmallVector<Value> decomposed =
LLVM::decomposeValue(rewriter, loc, src, i32);
unsigned andMask = op.getAndMask();
unsigned orMask = op.getOrMask();
unsigned xorMask = op.getXorMask();
// bit 15 is 0 for the BitMode swizzle.
// https://gpuopen.com/learn/amd-gcn-assembly-cross-lane-operations/
unsigned mask = andMask | (orMask << 5) | (xorMask << 10);
Value maskValue = createI32Constant(rewriter, loc, mask);
SmallVector<Value> swizzled;
for (Value v : decomposed) {
Value res =
ROCDL::DsSwizzleOp::create(rewriter, loc, v.getType(), v, maskValue);
swizzled.emplace_back(res);
}
Value result = LLVM::composeValue(rewriter, loc, swizzled, src.getType());
rewriter.replaceOp(op, result);
return success();
}
};
struct AMDGPUPermlaneLowering : public ConvertOpToLLVMPattern<PermlaneSwapOp> {
using ConvertOpToLLVMPattern::ConvertOpToLLVMPattern;
AMDGPUPermlaneLowering(const LLVMTypeConverter &converter, Chipset chipset)
: ConvertOpToLLVMPattern<PermlaneSwapOp>(converter), chipset(chipset) {}
Chipset chipset;
LogicalResult
matchAndRewrite(PermlaneSwapOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
if (chipset < kGfx950)
return op->emitOpError("permlane_swap is only supported on gfx950+");
Location loc = op.getLoc();
Type i32 = rewriter.getI32Type();
Value src = adaptor.getSrc();
unsigned rowLength = op.getRowLength();
bool fi = op.getFetchInactive();
bool boundctrl = op.getBoundCtrl();
SmallVector<Value> decomposed =
LLVM::decomposeValue(rewriter, loc, src, i32);
SmallVector<Value> permuted;
for (Value v : decomposed) {
Value res;
Type i32pair = LLVM::LLVMStructType::getLiteral(
rewriter.getContext(), {v.getType(), v.getType()});
if (rowLength == 16)
res = ROCDL::Permlane16SwapOp::create(rewriter, loc, i32pair, v, v, fi,
boundctrl);
else if (rowLength == 32)
res = ROCDL::Permlane32SwapOp::create(rewriter, loc, i32pair, v, v, fi,
boundctrl);
else
llvm_unreachable("unsupported row length");
Value vdst0 = LLVM::ExtractValueOp::create(rewriter, loc, res, {0});
Value vdst1 = LLVM::ExtractValueOp::create(rewriter, loc, res, {1});
Value isEqual = LLVM::ICmpOp::create(rewriter, loc,
LLVM::ICmpPredicate::eq, vdst0, v);
// Per `permlane(16|32)` semantics: if the first extracted element equals
// 'v', the result is the second element; otherwise it is the first.
Value vdstNew =
LLVM::SelectOp::create(rewriter, loc, isEqual, vdst1, vdst0);
permuted.emplace_back(vdstNew);
}
Value result = LLVM::composeValue(rewriter, loc, permuted, src.getType());
rewriter.replaceOp(op, result);
return success();
}
};
static Value setValueAtOffset(ConversionPatternRewriter &rewriter, Location loc,
Value accumulator, Value value, int64_t shift) {
shift = shift % 32;
Value shiftAmount;
if (shift != 0) {
shiftAmount = createI32Constant(rewriter, loc, shift % 32);
value = LLVM::ShlOp::create(rewriter, loc, value, shiftAmount);
}
if (matchPattern(accumulator, mlir::m_Zero()))
return value;
constexpr bool isDisjoint = true;
return LLVM::OrOp::create(rewriter, loc, accumulator, value, isDisjoint);
}
template <typename BaseOp>
struct AMDGPUMakeDmaBaseLowering : public ConvertOpToLLVMPattern<BaseOp> {
using ConvertOpToLLVMPattern<BaseOp>::ConvertOpToLLVMPattern;
using Adaptor = typename ConvertOpToLLVMPattern<BaseOp>::OpAdaptor;
AMDGPUMakeDmaBaseLowering(const LLVMTypeConverter &converter, Chipset chipset)
: ConvertOpToLLVMPattern<BaseOp>(converter), chipset(chipset) {}
Chipset chipset;
LogicalResult
matchAndRewrite(BaseOp op, Adaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
if (chipset < kGfx1250)
return op->emitOpError("make_dma_base is only supported on gfx1250");
Location loc = op.getLoc();
constexpr int32_t constlen = 4;
Value consts[constlen];
for (int64_t i = 0; i < constlen; ++i)
consts[i] = createI32Constant(rewriter, loc, i);
constexpr int32_t sgprslen = constlen;
Value sgprs[sgprslen];
for (int64_t i = 0; i < sgprslen; ++i) {
sgprs[i] = consts[0];
}
sgprs[0] = consts[1];
if constexpr (BaseOp::isGather()) {
sgprs[0] = setValueAtOffset(rewriter, loc, sgprs[0], consts[1], 30);
auto type = cast<TDMGatherBaseType>(op.getResult().getType());
Type indexType = type.getIndexType();
unsigned indexSize = indexType.getIntOrFloatBitWidth();
assert(llvm::is_contained({16u, 32u}, indexSize) &&
"expected index_size to be 16 or 32");
unsigned idx = (indexSize / 16) - 1;
if (idx)
sgprs[0] = setValueAtOffset(rewriter, loc, sgprs[0], consts[1], 31);
}
ValueRange ldsIndices = adaptor.getLdsIndices();
Value lds = adaptor.getLds();
auto ldsMemRefType = cast<MemRefType>(op.getLds().getType());
Value ldsPtr = ConvertToLLVMPattern::getStridedElementPtr(
rewriter, loc, ldsMemRefType, lds, ldsIndices);
ValueRange globalIndices = adaptor.getGlobalIndices();
Value global = adaptor.getGlobal();
auto globalMemRefType = cast<MemRefType>(op.getGlobal().getType());
Value globalPtr = ConvertToLLVMPattern::getStridedElementPtr(
rewriter, loc, globalMemRefType, global, globalIndices);
Type i32 = rewriter.getI32Type();
Type i64 = rewriter.getI64Type();
sgprs[1] = LLVM::PtrToIntOp::create(rewriter, loc, i32, ldsPtr);
Value castForGlobalAddr =
LLVM::PtrToIntOp::create(rewriter, loc, i64, globalPtr);
sgprs[2] = LLVM::TruncOp::create(rewriter, loc, i32, castForGlobalAddr);
Value shift = LLVM::LShrOp::create(rewriter, loc, castForGlobalAddr,
createI64Constant(rewriter, loc, 32));
Value highHalf = LLVM::TruncOp::create(rewriter, loc, i32, shift);
Value mask = createI32Constant(rewriter, loc, (1ull << 25) - 1);
highHalf = LLVM::AndOp::create(rewriter, loc, highHalf, mask);
sgprs[3] = setValueAtOffset(rewriter, loc, highHalf, consts[2], 30);
Type v4i32 = this->typeConverter->convertType(VectorType::get(4, i32));
assert(v4i32 && "expected type conversion to succeed");
Value result = LLVM::PoisonOp::create(rewriter, loc, v4i32);
for (auto [sgpr, constant] : llvm::zip_equal(sgprs, consts))
result =
LLVM::InsertElementOp::create(rewriter, loc, result, sgpr, constant);
rewriter.replaceOp(op, result);
return success();
}
};
template <typename DescriptorOp>
struct AMDGPULowerDescriptor : public ConvertOpToLLVMPattern<DescriptorOp> {
using ConvertOpToLLVMPattern<DescriptorOp>::ConvertOpToLLVMPattern;
using OpAdaptor = typename ConvertOpToLLVMPattern<DescriptorOp>::OpAdaptor;
AMDGPULowerDescriptor(const LLVMTypeConverter &converter, Chipset chipset)
: ConvertOpToLLVMPattern<DescriptorOp>(converter), chipset(chipset) {}
Chipset chipset;
Value getDGroup0(OpAdaptor adaptor) const { return adaptor.getBase(); }
Value setWorkgroupMask(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter, Location loc,
Value sgpr0) const {
Value mask = op.getWorkgroupMask();
if (!mask)
return sgpr0;
Type i16 = rewriter.getI16Type();
mask = LLVM::BitcastOp::create(rewriter, loc, i16, mask);
Type i32 = rewriter.getI32Type();
Value extendedMask = LLVM::ZExtOp::create(rewriter, loc, i32, mask);
return setValueAtOffset(rewriter, loc, sgpr0, extendedMask, 0);
}
Value setDataSize(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter, Location loc,
Value sgpr0, ArrayRef<Value> consts) const {
unsigned elementTypeWidthInBits = op.getElementTypeWidth();
assert(llvm::is_contained({8u, 16u, 32u, 64u}, elementTypeWidthInBits) &&
"expected type width to be 8, 16, 32, or 64.");
int64_t idx = llvm::Log2_32(elementTypeWidthInBits / 8);
Value size = consts[idx];
return setValueAtOffset(rewriter, loc, sgpr0, size, 16);
}
Value setAtomicBarrier(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter, Location loc,
Value sgpr0, ArrayRef<Value> consts) const {
if (!adaptor.getAtomicBarrierAddress())
return sgpr0;
return setValueAtOffset(rewriter, loc, sgpr0, consts[1], 18);
}
Value setIterateEnable(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter, Location loc,
Value sgpr0, ArrayRef<Value> consts) const {
if (!adaptor.getGlobalIncrement())
return sgpr0;
// Value is ignored when in gather mode.
// TODO: emit error earlier?
return setValueAtOffset(rewriter, loc, sgpr0, consts[1], 19);
}
Value setPadEnable(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter, Location loc,
Value sgpr0, ArrayRef<Value> consts) const {
if (!op.getPadAmount())
return sgpr0;
return setValueAtOffset(rewriter, loc, sgpr0, consts[1], 20);
}
Value setEarlyTimeout(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter, Location loc,
Value sgpr0, ArrayRef<Value> consts) const {
if (!op.getWorkgroupMask())
return sgpr0;
return setValueAtOffset(rewriter, loc, sgpr0, consts[1], 21);
}
Value setPadInterval(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter, Location loc,
Value sgpr0, ArrayRef<Value> consts) const {
if (!op.getPadAmount())
return sgpr0;
// pre-condition: padInterval can be a power of two between 2 and 256.
// TODO: Validation if the value breaks the pre-condition.
// If the pre-condition fails, there is a possibility of
// affecting the higher bits. In a following PR implement
// RuntimeVerifiableOpInterface that instruments conditions that need to be
// checked at runtime.
IntegerType i32 = rewriter.getI32Type();
Value padInterval = adaptor.getPadInterval();
padInterval = LLVM::CountTrailingZerosOp::create(rewriter, loc, i32,
padInterval, false);
padInterval = LLVM::SubOp::create(rewriter, loc, padInterval, consts[1]);
// post-condition: padInterval can be a value between 0 and 7.
return setValueAtOffset(rewriter, loc, sgpr0, padInterval, 22);
}
Value setPadAmount(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter, Location loc,
Value sgpr0, ArrayRef<Value> consts) const {
if (!op.getPadAmount())
return sgpr0;
// pre-condition: padAmount is a value between 1-128.
// TODO: Validation if the value breaks the pre-condition.
// If the pre-condition fails, there is a possibility of
// affecting the higher bits. In a following PR implement
// RuntimeVerifiableOpInterface that instruments conditions that need to be
// checked at runtime.
Value padAmount = adaptor.getPadAmount();
padAmount = LLVM::SubOp::create(rewriter, loc, padAmount, consts[1]);
// post-condition: padAmount is a value between 0-127.
return setValueAtOffset(rewriter, loc, sgpr0, padAmount, 25);
}
Value setAtomicBarrierAddress(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter,
Location loc, Value sgpr1,
ArrayRef<Value> consts) const {
if (!adaptor.getAtomicBarrierAddress())
return sgpr1;
Value atomicBarrierAddress = adaptor.getAtomicBarrierAddress();
auto barrierAddressTy =
cast<MemRefType>(op.getAtomicBarrierAddress().getType());
ValueRange atomicBarrierIndices = adaptor.getAtomicBarrierIndices();
atomicBarrierAddress = ConvertToLLVMPattern::getStridedElementPtr(
rewriter, loc, barrierAddressTy, atomicBarrierAddress,
atomicBarrierIndices);
IntegerType i32 = rewriter.getI32Type();
// pre-condition: atomicBarrierAddress is aligned to 8 bytes which implies
// that the 3 LSBs are zero.
// TODO: Validation if the value breaks the pre-condition.
// In a following PR implement RuntimeVerifiableOpInterface
// that instruments conditions that need to be checked at runtime.
atomicBarrierAddress =
LLVM::PtrToIntOp::create(rewriter, loc, i32, atomicBarrierAddress);
atomicBarrierAddress =
LLVM::LShrOp::create(rewriter, loc, atomicBarrierAddress, consts[3]);
Value mask = createI32Constant(rewriter, loc, 0xFFFF);
atomicBarrierAddress =
LLVM::AndOp::create(rewriter, loc, atomicBarrierAddress, mask);
return setValueAtOffset(rewriter, loc, sgpr1, atomicBarrierAddress, 32);
}
std::pair<Value, Value> setTensorDimX(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter,
Location loc, Value sgpr1, Value sgpr2,
ArrayRef<Value> consts, uint64_t dimX,
uint32_t offset) const {
ArrayRef<int64_t> globalStaticSizes = adaptor.getGlobalStaticSizes();
ValueRange globalDynamicSizes = adaptor.getGlobalDynamicSizes();
SmallVector<OpFoldResult> mixedGlobalSizes =
getMixedValues(globalStaticSizes, globalDynamicSizes, rewriter);
if (mixedGlobalSizes.size() <= dimX)
return {sgpr1, sgpr2};
OpFoldResult tensorDimXOpFoldResult = *(mixedGlobalSizes.rbegin() + dimX);
// pre-condition: tensorDimX is less than 2^32-1
// TODO: Validation if the value breaks the pre-condition.
// In a following PR implement RuntimeVerifiableOpInterface that instruments
// conditions that need to be checked at runtime. This could also be fixed
// by saying that mixedGlobalSizes is a DynamicI32List.
Value tensorDimX;
if (auto attr = dyn_cast<Attribute>(tensorDimXOpFoldResult)) {
tensorDimX =
createI32Constant(rewriter, loc, cast<IntegerAttr>(attr).getInt());
} else {
IntegerType i32 = rewriter.getI32Type();
tensorDimX = cast<Value>(tensorDimXOpFoldResult);
tensorDimX = LLVM::TruncOp::create(rewriter, loc, i32, tensorDimX);
}
sgpr1 = setValueAtOffset(rewriter, loc, sgpr1, tensorDimX, offset);
Value c16 = createI32Constant(rewriter, loc, 16);
Value tensorDimXHigh = LLVM::LShrOp::create(rewriter, loc, tensorDimX, c16);
sgpr2 = setValueAtOffset(rewriter, loc, sgpr2, tensorDimXHigh, offset + 16);
return {sgpr1, sgpr2};
}
std::pair<Value, Value> setTensorDim0(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter,
Location loc, Value sgpr1, Value sgpr2,
ArrayRef<Value> consts) const {
return setTensorDimX(op, adaptor, rewriter, loc, sgpr1, sgpr2, consts, 0,
48);
}
std::pair<Value, Value> setTensorDim1(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter,
Location loc, Value sgpr2, Value sgpr3,
ArrayRef<Value> consts) const {
return setTensorDimX(op, adaptor, rewriter, loc, sgpr2, sgpr3, consts, 1,
80);
}
Value setTileDimX(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter, Location loc,
Value sgpr, ArrayRef<Value> consts, size_t dimX,
int64_t offset) const {
ArrayRef<int64_t> sharedStaticSizes = adaptor.getSharedStaticSizes();
ValueRange sharedDynamicSizes = adaptor.getSharedDynamicSizes();
SmallVector<OpFoldResult> mixedSharedSizes =
getMixedValues(sharedStaticSizes, sharedDynamicSizes, rewriter);
if (mixedSharedSizes.size() <= dimX)
return sgpr;
OpFoldResult tileDimXOpFoldResult = *(mixedSharedSizes.rbegin() + dimX);
// pre-condition: tileDimX is less than 2^16-1
// TODO: Validation if the value breaks the pre-condition.
// If the pre-condition fails, there is a possibility of
// affecting the higher bits. In a following PR implement
// RuntimeVerifiableOpInterface that instruments conditions that need to be
// checked at runtime. This could also be fixed by saying that
// mixedSharedSizes is a DynamicI16List.
Value tileDimX;
if (auto attr = dyn_cast<Attribute>(tileDimXOpFoldResult)) {
tileDimX =
createI32Constant(rewriter, loc, cast<IntegerAttr>(attr).getInt());
} else {
IntegerType i32 = rewriter.getI32Type();
tileDimX = cast<Value>(tileDimXOpFoldResult);
tileDimX = LLVM::TruncOp::create(rewriter, loc, i32, tileDimX);
}
return setValueAtOffset(rewriter, loc, sgpr, tileDimX, offset);
}
Value setTileDim0(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter, Location loc,
Value sgpr3, ArrayRef<Value> consts) const {
return setTileDimX(op, adaptor, rewriter, loc, sgpr3, consts, 0, 112);
}
Value setTileDim1(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter, Location loc,
Value sgpr4, ArrayRef<Value> consts) const {
return setTileDimX(op, adaptor, rewriter, loc, sgpr4, consts, 1, 128);
}
Value setValidIndices(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter, Location loc,
Value sgpr4, ArrayRef<Value> consts) const {
auto type = cast<VectorType>(op.getIndices().getType());
ArrayRef<int64_t> shape = type.getShape();
assert(shape.size() == 1 && "expected shape to be of rank 1.");
unsigned length = shape.back();
assert(0 < length && length <= 16 && "expected length to be at most 16.");
Value value = createI32Constant(rewriter, loc, length);
return setValueAtOffset(rewriter, loc, sgpr4, value, 128);
}
Value setTileDim1OrValidIndices(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter,
Location loc, Value sgpr4,
ArrayRef<Value> consts) const {
if constexpr (DescriptorOp::isGather())
return setValidIndices(op, adaptor, rewriter, loc, sgpr4, consts);
return setTileDim1(op, adaptor, rewriter, loc, sgpr4, consts);
}
Value setTileDim2(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter, Location loc,
Value sgpr4, ArrayRef<Value> consts) const {
// Value is ignored when in gather mode.
if constexpr (DescriptorOp::isGather())
return sgpr4;
return setTileDimX(op, adaptor, rewriter, loc, sgpr4, consts, 2, 144);
}
std::pair<Value, Value>
setTensorDimXStride(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter, Location loc,
Value sgprY, Value sgprZ, ArrayRef<Value> consts,
size_t dimX, int64_t offset) const {
ArrayRef<int64_t> globalStaticStrides = adaptor.getGlobalStaticStrides();
ValueRange globalDynamicStrides = adaptor.getGlobalDynamicStrides();
SmallVector<OpFoldResult> mixedGlobalStrides =
getMixedValues(globalStaticStrides, globalDynamicStrides, rewriter);
if (mixedGlobalStrides.size() <= (dimX + 1))
return {sgprY, sgprZ};
OpFoldResult tensorDimXStrideOpFoldResult =
*(mixedGlobalStrides.rbegin() + dimX + 1);
// pre-condition: tensorDimXStride is less than 2^48-1
// TODO: Validation if the value breaks the pre-condition.
// In a following PR implement RuntimeVerifiableOpInterface that instruments
// conditions that need to be checked at runtime.
Value tensorDimXStride;
if (auto attr = dyn_cast<Attribute>(tensorDimXStrideOpFoldResult))
tensorDimXStride =
createI64Constant(rewriter, loc, cast<IntegerAttr>(attr).getInt());
else
tensorDimXStride = cast<Value>(tensorDimXStrideOpFoldResult);
constexpr int64_t first48bits = (1ll << 48) - 1;
Value mask = createI64Constant(rewriter, loc, first48bits);
tensorDimXStride =
LLVM::AndOp::create(rewriter, loc, mask, tensorDimXStride);
IntegerType i32 = rewriter.getI32Type();
Value tensorDimXStrideLow =
LLVM::TruncOp::create(rewriter, loc, i32, tensorDimXStride);
sgprY = setValueAtOffset(rewriter, loc, sgprY, tensorDimXStrideLow, offset);
int64_t shift = (offset % 32) == 0 ? 32 : offset % 32;
Value shiftVal = createI64Constant(rewriter, loc, shift);
Value tensorDimXStrideHigh =
LLVM::LShrOp::create(rewriter, loc, tensorDimXStride, shiftVal);
tensorDimXStrideHigh =
LLVM::TruncOp::create(rewriter, loc, i32, tensorDimXStrideHigh);
sgprZ = setValueAtOffset(rewriter, loc, sgprZ, tensorDimXStrideHigh,
offset + shift);
return {sgprY, sgprZ};
}
std::pair<Value, Value>
setTensorDim0Stride(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter, Location loc,
Value sgpr5, Value sgpr6, ArrayRef<Value> consts) const {
return setTensorDimXStride(op, adaptor, rewriter, loc, sgpr5, sgpr6, consts,
0, 160);
}
std::pair<Value, Value>
setTensorDim1Stride(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter, Location loc,
Value sgpr5, Value sgpr6, ArrayRef<Value> consts) const {
// Value is ignored when in gather mode.
if constexpr (DescriptorOp::isGather())
return {sgpr5, sgpr6};
return setTensorDimXStride(op, adaptor, rewriter, loc, sgpr5, sgpr6, consts,
1, 208);
}
Value getDGroup1(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter, Location loc,
ArrayRef<Value> consts) const {
Value sgprs[8];
for (int64_t i = 0; i < 8; ++i) {
sgprs[i] = consts[0];
}
sgprs[0] = setWorkgroupMask(op, adaptor, rewriter, loc, sgprs[0]);
sgprs[0] = setDataSize(op, adaptor, rewriter, loc, sgprs[0], consts);
sgprs[0] = setAtomicBarrier(op, adaptor, rewriter, loc, sgprs[0], consts);
sgprs[0] = setIterateEnable(op, adaptor, rewriter, loc, sgprs[0], consts);
sgprs[0] = setPadEnable(op, adaptor, rewriter, loc, sgprs[0], consts);
sgprs[0] = setEarlyTimeout(op, adaptor, rewriter, loc, sgprs[0], consts);
sgprs[0] = setPadInterval(op, adaptor, rewriter, loc, sgprs[0], consts);
sgprs[0] = setPadAmount(op, adaptor, rewriter, loc, sgprs[0], consts);
sgprs[1] =
setAtomicBarrierAddress(op, adaptor, rewriter, loc, sgprs[1], consts);
std::tie(sgprs[1], sgprs[2]) =
setTensorDim0(op, adaptor, rewriter, loc, sgprs[1], sgprs[2], consts);
std::tie(sgprs[2], sgprs[3]) =
setTensorDim1(op, adaptor, rewriter, loc, sgprs[2], sgprs[3], consts);
sgprs[3] = setTileDim0(op, adaptor, rewriter, loc, sgprs[3], consts);
sgprs[4] =
setTileDim1OrValidIndices(op, adaptor, rewriter, loc, sgprs[4], consts);
sgprs[4] = setTileDim2(op, adaptor, rewriter, loc, sgprs[4], consts);
std::tie(sgprs[5], sgprs[6]) = setTensorDim0Stride(
op, adaptor, rewriter, loc, sgprs[5], sgprs[6], consts);
std::tie(sgprs[6], sgprs[7]) = setTensorDim1Stride(
op, adaptor, rewriter, loc, sgprs[6], sgprs[7], consts);
IntegerType i32 = rewriter.getI32Type();
Type v8i32 = this->typeConverter->convertType(VectorType::get(8, i32));
assert(v8i32 && "expected type conversion to succeed");
Value dgroup1 = LLVM::PoisonOp::create(rewriter, loc, v8i32);
for (auto [sgpr, constant] : llvm::zip_equal(sgprs, consts)) {
dgroup1 =
LLVM::InsertElementOp::create(rewriter, loc, dgroup1, sgpr, constant);
}
return dgroup1;
}
Value setTensorDimX(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter, Location loc,
Value sgpr0, ArrayRef<Value> consts, int64_t dimX,
int64_t offset) const {
ArrayRef<int64_t> globalStaticSizes = adaptor.getGlobalStaticSizes();
ValueRange globalDynamicSizes = adaptor.getGlobalDynamicSizes();
SmallVector<OpFoldResult> mixedGlobalSizes =
getMixedValues(globalStaticSizes, globalDynamicSizes, rewriter);
if (mixedGlobalSizes.size() <= static_cast<unsigned long>(dimX))
return sgpr0;
OpFoldResult tensorDimXOpFoldResult = *(mixedGlobalSizes.rbegin() + dimX);
Value tensorDimX;
if (auto attr = dyn_cast<Attribute>(tensorDimXOpFoldResult)) {
tensorDimX =
createI32Constant(rewriter, loc, cast<IntegerAttr>(attr).getInt());
} else {
IntegerType i32 = rewriter.getI32Type();
tensorDimX = cast<Value>(tensorDimXOpFoldResult);
tensorDimX = LLVM::TruncOp::create(rewriter, loc, i32, tensorDimX);
}
return setValueAtOffset(rewriter, loc, sgpr0, tensorDimX, offset);
}
Value setTensorDim2(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter, Location loc,
Value sgpr0, ArrayRef<Value> consts) const {
return setTensorDimX(op, adaptor, rewriter, loc, sgpr0, consts, 2, 0);
}
Value truncateAndSetValueAtOffset(ConversionPatternRewriter &rewriter,
Location loc, Value accumulator,
Value value, int64_t shift) const {
IntegerType i32 = rewriter.getI32Type();
value = LLVM::TruncOp::create(rewriter, loc, i32, value);
return setValueAtOffset(rewriter, loc, accumulator, value, shift);
}
Value setLDSAddrIncrement(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter, Location loc,
Value sgpr1, ArrayRef<Value> consts,
int64_t offset) const {
Value ldsAddrIncrement = adaptor.getLdsIncrement();
return setValueAtOffset(rewriter, loc, sgpr1, ldsAddrIncrement, offset);
}
std::pair<Value, Value>
setGlobalAddrIncrement(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter, Location loc,
Value sgpr2, Value sgpr3, ArrayRef<Value> consts,
int64_t offset) const {
Value globalAddrIncrement = adaptor.getGlobalIncrement();
sgpr2 = truncateAndSetValueAtOffset(rewriter, loc, sgpr2,
globalAddrIncrement, offset);
Value shift = createI64Constant(rewriter, loc, 32);
globalAddrIncrement =
LLVM::LShrOp::create(rewriter, loc, globalAddrIncrement, shift);
constexpr int64_t first16BitsHigh = (1ll << 16) - 1;
sgpr3 = truncateAndSetValueAtOffset(rewriter, loc, sgpr3,
globalAddrIncrement, offset + 32);
Value mask = createI32Constant(rewriter, loc, first16BitsHigh);
sgpr3 = LLVM::AndOp::create(rewriter, loc, sgpr3, mask);
return {sgpr2, sgpr3};
}
Value setTensorDim3OrLDSAddrIncrement(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter,
Location loc, Value sgpr1,
ArrayRef<Value> consts) const {
Value ldsIncrement = op.getLdsIncrement();
constexpr int64_t dim = 3;
constexpr int64_t offset = 32;
if (!ldsIncrement)
return setTensorDimX(op, adaptor, rewriter, loc, sgpr1, consts, dim,
offset);
return setLDSAddrIncrement(op, adaptor, rewriter, loc, sgpr1, consts,
offset);
}
std::pair<Value, Value> setTensorDim2StrideOrGlobalAddrIncrement(
DescriptorOp op, OpAdaptor adaptor, ConversionPatternRewriter &rewriter,
Location loc, Value sgpr2, Value sgpr3, ArrayRef<Value> consts) const {
Value globalIncrement = op.getGlobalIncrement();
constexpr int32_t dim = 2;
constexpr int32_t offset = 64;
if (!globalIncrement)
return setTensorDimXStride(op, adaptor, rewriter, loc, sgpr2, sgpr3,
consts, dim, offset);
return setGlobalAddrIncrement(op, adaptor, rewriter, loc, sgpr2, sgpr3,
consts, offset);
}
Value setIterateCount(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter, Location loc,
Value sgpr3, ArrayRef<Value> consts,
int32_t offset) const {
Value iterationCount = adaptor.getIterationCount();
IntegerType i32 = rewriter.getI32Type();
// pre-condition: iterationCount is in the inclusive interval [1, 256].
// TODO: validation if the value breaks the pre-condition.
// If the pre-condition fails, there is a possibility of
// affecting the higher bits. In a following PR implement
// RuntimeVerifiableOpInterface that instruments conditions that need to be
// checked at runtime.
iterationCount = LLVM::TruncOp::create(rewriter, loc, i32, iterationCount);
iterationCount =
LLVM::SubOp::create(rewriter, loc, iterationCount, consts[1]);
return setValueAtOffset(rewriter, loc, sgpr3, iterationCount, offset);
}
Value setTileDim3OrIterateCount(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter,
Location loc, Value sgpr3,
ArrayRef<Value> consts) const {
Value iterateCount = op.getIterationCount();
constexpr int32_t dim = 2;
constexpr int32_t offset = 112;
if (!iterateCount)
return setTileDimX(op, adaptor, rewriter, loc, sgpr3, consts, dim,
offset);
return setIterateCount(op, adaptor, rewriter, loc, sgpr3, consts, offset);
}
Value getDGroup2(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter, Location loc,
ArrayRef<Value> consts) const {
if constexpr (DescriptorOp::isGather())
return getDGroup2Gather(op, adaptor, rewriter, loc, consts);
return getDGroup2NonGather(op, adaptor, rewriter, loc, consts);
}
Value getDGroup2NonGather(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter, Location loc,
ArrayRef<Value> consts) const {
IntegerType i32 = rewriter.getI32Type();
Type v4i32 = this->typeConverter->convertType(VectorType::get(4, i32));
assert(v4i32 && "expected type conversion to succeed.");
bool onlyNeedsTwoDescriptors = !op.getLdsIncrement() && op.getRank() <= 2;
if (onlyNeedsTwoDescriptors)
return LLVM::ZeroOp::create(rewriter, loc, v4i32);
constexpr int64_t sgprlen = 4;
Value sgprs[sgprlen];
for (int i = 0; i < sgprlen; ++i)
sgprs[i] = consts[0];
sgprs[0] = setTensorDim2(op, adaptor, rewriter, loc, sgprs[0], consts);
sgprs[1] = setTensorDim3OrLDSAddrIncrement(op, adaptor, rewriter, loc,
sgprs[1], consts);
std::tie(sgprs[2], sgprs[3]) = setTensorDim2StrideOrGlobalAddrIncrement(
op, adaptor, rewriter, loc, sgprs[2], sgprs[3], consts);
sgprs[3] =
setTileDim3OrIterateCount(op, adaptor, rewriter, loc, sgprs[3], consts);
Value dgroup2 = LLVM::PoisonOp::create(rewriter, loc, v4i32);
for (auto [sgpr, constant] : llvm::zip(sgprs, consts))
dgroup2 =
LLVM::InsertElementOp::create(rewriter, loc, dgroup2, sgpr, constant);
return dgroup2;
}
Value getGatherIndices(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter, Location loc,
ArrayRef<Value> consts, bool firstHalf) const {
IntegerType i32 = rewriter.getI32Type();
Type v4i32 = this->typeConverter->convertType(VectorType::get(4, i32));
assert(v4i32 && "expected type conversion to succeed.");
Value indices = adaptor.getIndices();
auto vectorType = cast<VectorType>(indices.getType());
unsigned length = vectorType.getShape().back();
Type elementType = vectorType.getElementType();
unsigned maxLength = elementType == i32 ? 4 : 8;
int32_t offset = firstHalf ? 0 : maxLength;
unsigned discountedLength =
std::max(static_cast<int32_t>(length - offset), 0);
unsigned targetSize = std::min(maxLength, discountedLength);
SmallVector<Value> indicesVector;
for (unsigned i = offset; i < targetSize + offset; ++i) {
Value idx;
if (i < consts.size())
idx = consts[i];
else
idx = createI32Constant(rewriter, loc, i);
Value elem = LLVM::ExtractElementOp::create(rewriter, loc, indices, idx);
indicesVector.push_back(elem);
}
SmallVector<Value> indicesI32Vector;
if (elementType == i32) {
indicesI32Vector = indicesVector;
} else {
for (unsigned i = 0; i < targetSize; ++i) {
Value index = indicesVector[i];
indicesI32Vector.push_back(
LLVM::ZExtOp::create(rewriter, loc, i32, index));
}
if ((targetSize % 2) != 0)
// Add padding when not divisible by two.
indicesI32Vector.push_back(consts[0]);
}
SmallVector<Value> indicesToInsert;
if (elementType == i32) {
indicesToInsert = indicesI32Vector;
} else {
unsigned size = indicesI32Vector.size() / 2;
for (unsigned i = 0; i < size; ++i) {
Value first = indicesI32Vector[2 * i];
Value second = indicesI32Vector[2 * i + 1];
Value joined = setValueAtOffset(rewriter, loc, first, second, 16);
indicesToInsert.push_back(joined);
}
}
Value dgroup = LLVM::PoisonOp::create(rewriter, loc, v4i32);
for (auto [sgpr, constant] : llvm::zip_first(indicesToInsert, consts))
dgroup =
LLVM::InsertElementOp::create(rewriter, loc, dgroup, sgpr, constant);
return dgroup;
}
Value getDGroup2Gather(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter, Location loc,
ArrayRef<Value> consts) const {
return getGatherIndices(op, adaptor, rewriter, loc, consts, true);
}
std::pair<Value, Value>
setTensorDim3Stride(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter, Location loc,
Value sgpr0, Value sgpr1, ArrayRef<Value> consts) const {
constexpr int32_t dim = 3;
constexpr int32_t offset = 0;
return setTensorDimXStride(op, adaptor, rewriter, loc, sgpr0, sgpr1, consts,
dim, offset);
}
std::pair<Value, Value> setTensorDim4(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter,
Location loc, Value sgpr1, Value sgpr2,
ArrayRef<Value> consts) const {
constexpr int32_t dim = 4;
constexpr int32_t offset = 48;
return setTensorDimX(op, adaptor, rewriter, loc, sgpr1, sgpr2, consts, dim,
offset);
}
Value setTileDim4(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter, Location loc,
Value sgpr2, ArrayRef<Value> consts) const {
constexpr int32_t dim = 4;
constexpr int32_t offset = 80;
return setTileDimX(op, adaptor, rewriter, loc, sgpr2, consts, dim, offset);
}
Value getDGroup3(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter, Location loc,
ArrayRef<Value> consts) const {
if constexpr (DescriptorOp::isGather())
return getDGroup3Gather(op, adaptor, rewriter, loc, consts);
return getDGroup3NonGather(op, adaptor, rewriter, loc, consts);
}
Value getDGroup3NonGather(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter, Location loc,
ArrayRef<Value> consts) const {
IntegerType i32 = rewriter.getI32Type();
Type v4i32 = this->typeConverter->convertType(VectorType::get(4, i32));
assert(v4i32 && "expected type conversion to succeed.");
bool onlyNeedsTwoDescriptors = !op.getLdsIncrement() && op.getRank() <= 2;
if (onlyNeedsTwoDescriptors)
return LLVM::ZeroOp::create(rewriter, loc, v4i32);
constexpr int32_t sgprlen = 4;
Value sgprs[sgprlen];
for (int i = 0; i < sgprlen; ++i)
sgprs[i] = consts[0];
std::tie(sgprs[0], sgprs[1]) = setTensorDim3Stride(
op, adaptor, rewriter, loc, sgprs[0], sgprs[1], consts);
std::tie(sgprs[1], sgprs[2]) =
setTensorDim4(op, adaptor, rewriter, loc, sgprs[1], sgprs[2], consts);
sgprs[2] = setTileDim4(op, adaptor, rewriter, loc, sgprs[2], consts);
Value dgroup3 = LLVM::PoisonOp::create(rewriter, loc, v4i32);
for (auto [sgpr, constant] : llvm::zip(sgprs, consts))
dgroup3 =
LLVM::InsertElementOp::create(rewriter, loc, dgroup3, sgpr, constant);
return dgroup3;
}
Value getDGroup3Gather(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter, Location loc,
ArrayRef<Value> consts) const {
return getGatherIndices(op, adaptor, rewriter, loc, consts, false);
}
LogicalResult
matchAndRewrite(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
if (chipset < kGfx1250)
return op->emitOpError(
"make_dma_descriptor is only supported on gfx1250");
Location loc = op.getLoc();
SmallVector<Value> consts;
for (int64_t i = 0; i < 8; ++i)
consts.push_back(createI32Constant(rewriter, loc, i));
Value dgroup0 = this->getDGroup0(adaptor);
Value dgroup1 = this->getDGroup1(op, adaptor, rewriter, loc, consts);
Value dgroup2 = this->getDGroup2(op, adaptor, rewriter, loc, consts);
Value dgroup3 = this->getDGroup3(op, adaptor, rewriter, loc, consts);
SmallVector<Value> results = {dgroup0, dgroup1, dgroup2, dgroup3};
rewriter.replaceOpWithMultiple(op, {results});
return success();
}
};
template <typename SourceOp, typename TargetOp>
struct AMDGPUTensorLoadStoreOpLowering
: public ConvertOpToLLVMPattern<SourceOp> {
using ConvertOpToLLVMPattern<SourceOp>::ConvertOpToLLVMPattern;
using Adaptor = typename ConvertOpToLLVMPattern<SourceOp>::OneToNOpAdaptor;
AMDGPUTensorLoadStoreOpLowering(const LLVMTypeConverter &converter,
Chipset chipset)
: ConvertOpToLLVMPattern<SourceOp>(converter), chipset(chipset) {}
Chipset chipset;
LogicalResult
matchAndRewrite(SourceOp op, Adaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
if (chipset < kGfx1250)
return op->emitOpError("is only supported on gfx1250");
ValueRange desc = adaptor.getDesc();
rewriter.replaceOpWithNewOp<TargetOp>(op, desc[0], desc[1], desc[2],
desc[3], /*cachePolicy=*/0,
/*alias_scopes=*/nullptr,
/*noalias_scopes=*/nullptr,
/*tbaa=*/nullptr);
return success();
}
};
struct ConvertAMDGPUToROCDLPass
: public impl::ConvertAMDGPUToROCDLPassBase<ConvertAMDGPUToROCDLPass> {
using Base::Base;
void runOnOperation() override {
MLIRContext *ctx = &getContext();
FailureOr<Chipset> maybeChipset = Chipset::parse(chipset);
if (failed(maybeChipset)) {
emitError(UnknownLoc::get(ctx), "Invalid chipset name: " + chipset);
return signalPassFailure();
}
RewritePatternSet patterns(ctx);
LLVMTypeConverter converter(ctx);
populateAMDGPUToROCDLConversionPatterns(converter, patterns, *maybeChipset);
amdgpu::populateCommonGPUTypeAndAttributeConversions(converter);
LLVMConversionTarget target(getContext());
target.addIllegalDialect<::mlir::amdgpu::AMDGPUDialect>();
target.addLegalDialect<::mlir::LLVM::LLVMDialect>();
target.addLegalDialect<::mlir::ROCDL::ROCDLDialect>();
if (failed(applyPartialConversion(getOperation(), target,
std::move(patterns))))
signalPassFailure();
}
};
} // namespace
void mlir::amdgpu::populateCommonGPUTypeAndAttributeConversions(
TypeConverter &typeConverter) {
populateGpuMemorySpaceAttributeConversions(
typeConverter, [](gpu::AddressSpace space) {
switch (space) {
case gpu::AddressSpace::Global:
return ROCDL::ROCDLDialect::kGlobalMemoryAddressSpace;
case gpu::AddressSpace::Workgroup:
return ROCDL::ROCDLDialect::kSharedMemoryAddressSpace;
case gpu::AddressSpace::Private:
return ROCDL::ROCDLDialect::kPrivateMemoryAddressSpace;
}
llvm_unreachable("unknown address space enum value");
});
}
void mlir::populateAMDGPUTypeAndAttributeConversions(
TypeConverter &typeConverter) {
typeConverter.addTypeAttributeConversion(
[](BaseMemRefType type, amdgpu::AddressSpaceAttr as)
-> TypeConverter::AttributeConversionResult {
MLIRContext *ctx = as.getContext();
Type i64 = IntegerType::get(ctx, 64);
switch (as.getValue()) {
case amdgpu::AddressSpace::FatRawBuffer:
return IntegerAttr::get(i64, 7);
case amdgpu::AddressSpace::BufferRsrc:
return IntegerAttr::get(i64, 8);
case amdgpu::AddressSpace::FatStructuredBuffer:
return IntegerAttr::get(i64, 9);
}
return TypeConverter::AttributeConversionResult::abort();
});
typeConverter.addConversion([&](TDMBaseType type) -> Type {
Type i32 = IntegerType::get(type.getContext(), 32);
return typeConverter.convertType(VectorType::get(4, i32));
});
typeConverter.addConversion([&](TDMGatherBaseType type) -> Type {
Type i32 = IntegerType::get(type.getContext(), 32);
return typeConverter.convertType(VectorType::get(4, i32));
});
typeConverter.addConversion(
[&](TDMDescriptorType type,
SmallVectorImpl<Type> &result) -> std::optional<LogicalResult> {
Type i32 = IntegerType::get(type.getContext(), 32);
Type v4i32 = typeConverter.convertType(VectorType::get(4, i32));
Type v8i32 = typeConverter.convertType(VectorType::get(8, i32));
llvm::append_values(result, v4i32, v8i32, v4i32, v4i32);
return success();
});
auto addUnrealizedCast = [](OpBuilder &builder, TypeRange types,
ValueRange inputs,
Location loc) -> SmallVector<Value> {
// Only create unrealized_conversion_cast for TDMDescriptorType.
// All other types which are not expected, should be
// materialized by other target materialization functions.
if (inputs.size() != 1)
return {};
if (!isa<TDMDescriptorType>(inputs[0].getType()))
return {};
auto cast = UnrealizedConversionCastOp::create(builder, loc, types, inputs);
return cast.getResults();
};
typeConverter.addTargetMaterialization(addUnrealizedCast);
}
void mlir::populateAMDGPUToROCDLConversionPatterns(LLVMTypeConverter &converter,
RewritePatternSet &patterns,
Chipset chipset) {
populateAMDGPUTypeAndAttributeConversions(converter);
patterns
.add<FatRawBufferCastLowering,
RawBufferOpLowering<RawBufferLoadOp, ROCDL::RawPtrBufferLoadOp>,
RawBufferOpLowering<RawBufferStoreOp, ROCDL::RawPtrBufferStoreOp>,
RawBufferOpLowering<RawBufferAtomicFaddOp,
ROCDL::RawPtrBufferAtomicFaddOp>,
RawBufferOpLowering<RawBufferAtomicFmaxOp,
ROCDL::RawPtrBufferAtomicFmaxOp>,
RawBufferOpLowering<RawBufferAtomicSmaxOp,
ROCDL::RawPtrBufferAtomicSmaxOp>,
RawBufferOpLowering<RawBufferAtomicUminOp,
ROCDL::RawPtrBufferAtomicUminOp>,
RawBufferOpLowering<RawBufferAtomicCmpswapOp,
ROCDL::RawPtrBufferAtomicCmpSwap>,
AMDGPUDPPLowering, MemoryCounterWaitOpLowering, LDSBarrierOpLowering,
SchedBarrierOpLowering, MFMAOpLowering, ScaledMFMAOpLowering,
SparseMFMAOpLowering, WMMAOpLowering, ScaledWMMAOpLowering,
ExtPackedFp8OpLowering, ScaledExtPackedMatrixOpLowering,
ScaledExtPackedOpLowering, PackedScaledTruncOpLowering,
PackedTrunc2xFp8OpLowering, PackedStochRoundFp8OpLowering,
GatherToLDSOpLowering, TransposeLoadOpLowering,
AMDGPUPermlaneLowering, AMDGPUMakeDmaBaseLowering<MakeDmaBaseOp>,
AMDGPUMakeDmaBaseLowering<MakeGatherDmaBaseOp>,
AMDGPULowerDescriptor<MakeDmaDescriptorOp>,
AMDGPULowerDescriptor<MakeGatherDmaDescriptorOp>,
AMDGPUTensorLoadStoreOpLowering<TensorLoadToLDSOp,
ROCDL::TensorLoadToLDSOp>,
AMDGPUTensorLoadStoreOpLowering<TensorStoreFromLDSOp,
ROCDL::TensorStoreFromLDSOp>>(
converter, chipset);
patterns.add<AMDGPUSwizzleBitModeLowering>(converter);
}